maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA][HIP] Diagnose reference of host variable 

This patch diagnoses invalid references of global host variables in device,
global, or host device functions.

Differential Revision: https://reviews.llvm.org/D91281
This commit is contained in:
Yaxun (Sam) Liu 2020-11-11 08:39:00 -05:00
parent cd95338ee3
commit 5c8911d0ba
6 changed files with 186 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -8145,7 +8145,7 @@ def err_global_call_not_config : Error<
"call to global function %0 not configured">;
def err_ref_bad_target : Error<
"reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
"function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
"%select{function|variable}1 %2 in %select{__device__|__global__|__host__|__host__ __device__}3 function">;
def err_ref_bad_target_global_initializer : Error<
"reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
"function %1 in global initializer">;

3
clang/lib/Sema/SemaCUDA.cpp
View File

@ -743,7 +743,8 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
return true;
SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
<< IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
<< IdentifyCUDATarget(Callee) << /*function*/ 0 << Callee
<< IdentifyCUDATarget(Caller);
if (!Callee->getBuiltinID())
SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
diag::note_previous_decl, Caller, *this)

18
clang/lib/Sema/SemaExpr.cpp
View File

@ -354,6 +354,24 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
// CUDA/HIP: Diagnose invalid references of host global variables in device
// functions. Reference of device global variables in host functions is
// allowed through shadow variables therefore it is not diagnosed.
if (LangOpts.CUDAIsDevice) {
auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
auto Target = IdentifyCUDATarget(FD);
if (FD && Target != CFT_Host) {
const auto *VD = dyn_cast<VarDecl>(D);
if (VD && VD->hasGlobalStorage() && !VD->hasAttr<CUDADeviceAttr>() &&
!VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
!VD->getType()->isCUDADeviceBuiltinSurfaceType() &&
!VD->getType()->isCUDADeviceBuiltinTextureType() &&
!VD->isConstexpr() && !VD->getType().isConstQualified())
targetDiag(*Locs.begin(), diag::err_ref_bad_target)
<< /*host*/ 2 << /*variable*/ 1 << VD << Target;
}
}
if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
if (const auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);

3
clang/test/CodeGenCUDA/function-overload.cu
View File

@ -12,6 +12,9 @@
#include "Inputs/cuda.h"
// Check constructors/destructors for D/H functions
#ifdef __CUDA_ARCH__
__device__
#endif
int x;
struct s_cd_dh {
__host__ s_cd_dh() { x = 11; }

4
clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
View File

@ -124,7 +124,7 @@ __attribute__((device)) void test_shared64() {
val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
}
__UINT32_TYPE__ global_val32;
__attribute__((device)) __UINT32_TYPE__ global_val32;
__attribute__((device)) void test_global32() {
// CHECK-LABEL: test_global32
// CHECK: %0 = load i32, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
@ -138,7 +138,7 @@ __attribute__((device)) void test_global32() {
global_val32 = __builtin_amdgcn_atomic_dec32(&global_val32, global_val32, __ATOMIC_SEQ_CST, "workgroup");
}
__UINT64_TYPE__ global_val64;
__attribute__((device)) __UINT64_TYPE__ global_val64;
__attribute__((device)) void test_global64() {
// CHECK-LABEL: test_global64
// CHECK: %0 = load i64, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8

160
clang/test/SemaCUDA/device-use-host-var.cu Normal file
View File

@ -0,0 +1,160 @@
// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host %s
// host-no-diagnostics
#include "Inputs/cuda.h"
int global_host_var;
__device__ int global_dev_var;
__constant__ int global_constant_var;
__shared__ int global_shared_var;
constexpr int global_constexpr_var = 1;
const int global_const_var = 1;
template<typename F>
__global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}}
__device__ void dev_fun(int *out) {
int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var;
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
}
__global__ void global_fun(int *out) {
int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var;
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
}
__host__ __device__ void host_dev_fun(int *out) {
int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var;
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
}
inline __host__ __device__ void inline_host_dev_fun(int *out) {
int &ref_host_var = global_host_var;
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var;
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var;
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
}
void dev_lambda_capture_by_ref(int *out) {
int &ref_host_var = global_host_var;
kernel<<<1,1>>>([&]() {
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
// dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}}
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var; // dev-error {{capture host variable 'ref_host_var' by reference in device or host device lambda function}}
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
});
}
void dev_lambda_capture_by_copy(int *out) {
int &ref_host_var = global_host_var;
kernel<<<1,1>>>([=]() {
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
const int &ref_const_var = global_const_var;
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
*out = global_constexpr_var;
*out = global_const_var;
*out = ref_host_var;
*out = ref_dev_var;
*out = ref_constant_var;
*out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
});
}