Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -38,7 +38,7 @@ struct __compare_result{}; template<> struct __compare_result { - static const bool valid; + static const __device__ bool valid; }; __DEVICE__ Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -355,24 +355,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); @@ -17143,6 +17125,31 @@ CaptureType, DeclRefType, FunctionScopeIndexToStopAt); + // Diagnose ODR-use 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.CUDA && SemaRef.LangOpts.CUDAIsDevice) { + auto *FD = dyn_cast_or_null(SemaRef.CurContext); + auto Target = SemaRef.IdentifyCUDATarget(FD); + auto IsEmittedOnDeviceSide = [](VarDecl *Var) { + if (Var->hasAttr() || Var->hasAttr() || + Var->hasAttr() || + Var->getType()->isCUDADeviceBuiltinSurfaceType() || + Var->getType()->isCUDADeviceBuiltinTextureType()) + return true; + // Function-scope static variable in device functions or kernels are + // emitted on device side. + if (auto *FD = dyn_cast(Var->getDeclContext())) { + return FD->hasAttr() || FD->hasAttr(); + } + return false; + }; + if (Var && Var->hasGlobalStorage() && !IsEmittedOnDeviceSide(Var)) { + SemaRef.targetDiag(Loc, diag::err_ref_bad_target) + << /*host*/ 2 << /*variable*/ 1 << Var << Target; + } + } + Var->markUsed(SemaRef.Context); } Index: clang/test/CodeGenCUDA/device-use-host-var.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/device-use-host-var.cu @@ -0,0 +1,40 @@ +// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa \ +// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s + +#include "Inputs/cuda.h" + +struct A { + int x; +}; + +constexpr int constexpr_var = 1; +constexpr A constexpr_struct{2}; +constexpr A constexpr_array[4] = {0, 0, 0, 3}; +constexpr char constexpr_str[] = "abcd"; +const int const_var = 4; + +// CHECK-DAG: @_ZL13constexpr_str.const = private unnamed_addr addrspace(4) constant [5 x i8] c"abcd\00" +// CHECK-DAG: @_ZL13constexpr_var = internal addrspace(4) constant i32 1 +// CHECK-DAG: @_ZL16constexpr_struct = internal addrspace(4) constant %struct.A { i32 2 } +// CHECK-DAG: @_ZL15constexpr_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 3 }] +// CHECK-NOT: external + +// CHECK-LABEL: define{{.*}}@_Z7dev_funPiPPKi +// CHECK: store i32 1 +// CHECK: store i32 2 +// CHECK: store i32 3 +// CHECK: store i32 4 +// CHECK: load i8, i8* getelementptr {{.*}} @_ZL13constexpr_str.const +// CHECK: store i32* {{.*}}@_ZL13constexpr_var +// CHECK: store i32* getelementptr {{.*}} @_ZL16constexpr_struct +// CHECK: store i32* getelementptr {{.*}} @_ZL15constexpr_array +__device__ void dev_fun(int *out, const int **out2) { + *out = constexpr_var; + *out = constexpr_struct.x; + *out = constexpr_array[3].x; + *out = const_var; + *out = constexpr_str[3]; + *out2 = &constexpr_var; + *out2 = &constexpr_struct.x; + *out2 = &constexpr_array[3].x; +} Index: clang/test/Headers/hip-header.hip =================================================================== --- clang/test/Headers/hip-header.hip +++ clang/test/Headers/hip-header.hip @@ -5,6 +5,13 @@ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -D__HIPCC_RTC__ | FileCheck %s +// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/Inputs/include \ +// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ +// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ +// RUN: -D__HIPCC_RTC__ -std=c++14 | FileCheck %s + // expected-no-diagnostics 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,37 +5,96 @@ #include "Inputs/cuda.h" -int global_host_var; +struct A { + int x; + static int host_var; +}; + +int A::host_var; + +namespace X { + int host_var; +} + +static int static_host_var; + __device__ int global_dev_var; __constant__ int global_constant_var; __shared__ int global_shared_var; -constexpr int global_constexpr_var = 1; + +int global_host_var; const int global_const_var = 1; +constexpr int global_constexpr_var = 1; + +int global_host_array[2] = {1, 2}; +const int global_const_array[2] = {1, 2}; +constexpr int global_constexpr_array[2] = {1, 2}; + +A global_host_struct_var{1}; +const A global_const_struct_var{1}; +constexpr A global_constexpr_struct_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}} + // Check access device variables are allowed. 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 = ref_dev_var; + *out = ref_constant_var; + *out = ref_shared_var; *out = global_dev_var; *out = global_constant_var; *out = global_shared_var; - *out = global_constexpr_var; + + // Check access of non-const host variables are not allowed. + *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} *out = global_const_var; + *out = global_constexpr_var; + global_host_var = 1; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} + // Check reference of non-constexpr host variables are not allowed. + int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} + const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}} + const int &ref_constexpr_var = global_constexpr_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; + + // Check access member of non-constexpr struct type host variable is not allowed. + *out = global_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}} + *out = global_const_struct_var.x; // dev-error {{reference to __host__ variable 'global_const_struct_var' in __device__ function}} + *out = global_constexpr_struct_var.x; + global_host_struct_var.x = 1; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}} + + // Check address taking of non-constexpr host variables is not allowed. + int *p = &global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}} + const int *cp = &global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}} + const int *cp2 = &global_constexpr_var; + + // Check access elements of non-constexpr host array is not allowed. + *out = global_host_array[1]; // dev-error {{reference to __host__ variable 'global_host_array' in __device__ function}} + *out = global_const_array[1]; // dev-error {{reference to __host__ variable 'global_const_array' in __device__ function}} + *out = global_constexpr_array[1]; + + // Check ODR-use of host variables in namespace is not allowed. + *out = X::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}} + + // Check ODR-use of static host varables in class or file scope is not allowed. + *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}} + *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}} + + // Check function-scope static variable is allowed. + static int static_var; + *out = static_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) { @@ -44,7 +103,7 @@ 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; + const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __global__ function}} *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}} *out = global_dev_var; @@ -67,7 +126,7 @@ 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; + const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}} *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} *out = global_dev_var; @@ -114,7 +173,7 @@ 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; + const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}} *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}} @@ -140,7 +199,7 @@ 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; + const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}} *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}} *out = global_dev_var; @@ -166,7 +225,7 @@ template struct __attribute__((device_builtin_texture_type)) texture { static texture ref; - __device__ int c() { + __device__ void c() { auto &x = ref; } }; @@ -174,7 +233,40 @@ template struct not_a_texture { static not_a_texture ref; - __device__ int c() { + __device__ void c() { auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}} } }; + +template<> +not_a_texture not_a_texture::ref; + +__device__ void test_not_a_texture() { + not_a_texture inst; + inst.c(); // dev-note {{in instantiation of member function 'not_a_texture::c' requested here}} +} + +// Test static variable in host function used by device function. +void test_static_var_host() { + for (int i = 0; i < 10; i++) { + static int x; + struct A { + __device__ int f() { + return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}} + } + }; + } +} + +// Test static variable in device function used by device function. +__device__ void test_static_var_device() { + for (int i = 0; i < 10; i++) { + static int x; + int y = x; + struct A { + __device__ int f() { + return x; + } + }; + } +}