diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -8145,7 +8145,7 @@ "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">; diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -743,7 +743,8 @@ 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) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -354,6 +354,24 @@ 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(CurContext); + auto Target = IdentifyCUDATarget(FD); + if (FD && Target != CFT_Host) { + const auto *VD = dyn_cast(D); + if (VD && VD->hasGlobalStorage() && !VD->hasAttr() && + !VD->hasAttr() && !VD->hasAttr() && + !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(D)) checkDeviceDecl(VD, Loc); diff --git a/clang/test/CodeGenCUDA/function-overload.cu b/clang/test/CodeGenCUDA/function-overload.cu --- a/clang/test/CodeGenCUDA/function-overload.cu +++ b/clang/test/CodeGenCUDA/function-overload.cu @@ -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; } diff --git a/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp b/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp --- a/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp +++ b/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp @@ -124,7 +124,7 @@ 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 @@ 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 diff --git a/clang/test/SemaCUDA/device-use-host-var.cu b/clang/test/SemaCUDA/device-use-host-var.cu new file mode 100644 --- /dev/null +++ b/clang/test/SemaCUDA/device-use-host-var.cu @@ -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 +__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; + }); +} +