Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -354,24 +354,6 @@ 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 (auto *VD = dyn_cast(D)) checkDeviceDecl(VD, Loc); @@ -18284,6 +18266,24 @@ } break; } + + // 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 (SemaRef.LangOpts.CUDAIsDevice) { + auto *FD = dyn_cast_or_null(SemaRef.CurContext); + auto Target = SemaRef.IdentifyCUDATarget(FD); + if (FD && Target != Sema::CFT_Host) { + if (Var && Var->hasGlobalStorage() && !Var->hasAttr() && + !Var->hasAttr() && + !Var->hasAttr() && + !Var->getType()->isCUDADeviceBuiltinSurfaceType() && + !Var->getType()->isCUDADeviceBuiltinTextureType() && + !Var->isConstexpr() && !Var->getType().isConstQualified()) + SemaRef.targetDiag(Loc, diag::err_ref_bad_target) + << /*host*/ 2 << /*variable*/ 1 << Var << Target; + } + } } /// Mark a variable referenced, and check whether it is odr-used Index: clang/test/SemaCUDA/device-use-host-var.cu =================================================================== --- clang/test/SemaCUDA/device-use-host-var.cu +++ clang/test/SemaCUDA/device-use-host-var.cu @@ -5,6 +5,11 @@ #include "Inputs/cuda.h" +struct A { + int x; + A() {} +}; + int global_host_var; __device__ int global_dev_var; __constant__ int global_constant_var; @@ -12,6 +17,9 @@ constexpr int global_constexpr_var = 1; const int global_const_var = 1; +A global_host_struct_var; +const A global_const_struct_var; + template __global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}} @@ -24,11 +32,13 @@ 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_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_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 = global_const_struct_var.x; *out = ref_host_var; *out = ref_dev_var; @@ -36,6 +46,12 @@ *out = ref_shared_var; *out = ref_constexpr_var; *out = ref_const_var; + + // Check non-ODR use of host varirables are allowed. + *out = sizeof(global_host_var); + *out = sizeof(global_host_struct_var.x); + decltype(global_host_var) var1; + decltype(global_host_struct_var.x) var2; } __global__ void global_fun(int *out) {