[CUDA][HIP] Allow non-ODR use of host var in device

[tra]

clang/test/SemaCUDA/device-use-host-var.cu

55

I do not think it should be allowed. We end up instantiating the variable on device, even though the variable should be host-only.

Right now we allow it, but end up with an .extern .const which will make ptxas fail:
https://godbolt.org/z/sx9845

If we do allow it, we'll need to make sure that we only use the value, but do not allow instantiating the variable.

[yaxunl]

clang/test/SemaCUDA/device-use-host-var.cu

55

will only not allow this since it results in ODR-use of the host var.

[yaxunl]

Follow C++ about ODR-use of variables.

[yaxunl]

minor bug fix

[tra]

LGTM

I've added @rsmith to double check that we're handling it correctly.

[yaxunl]

ping

[tra]

I think we also may want to check that we allow sizeof(host_var) in the GPU code.

[yaxunl]

In D98193#2656326, @tra wrote:

I think we also may want to check that we allow sizeof(host_var) in the GPU code.

We have tests for that at line 94 of test/SemaCUDA/device-use-host-var.cu

[rsmith]

clang/lib/Sema/SemaExpr.cpp

17134

I suspect you want hasGlobalStorage rather than isFileVarDecl here (that is, preserve the condition from the deleted code), in order to disallow use of host-side local static variables from device-side functions:

__host__ void f() {
  static int n;
  struct X {
    __device__ void g() { ++n; }
  };
  // ...
}

[rsmith]

In principle this seems reasonable to me.

[yaxunl]

clang/lib/Sema/SemaExpr.cpp

17134

For function scope static variable, if the smallest enclosing function is device or device host function, the static variable without __device__ or __constant__ attribute is allowed and the variable is emitted at device side (https://godbolt.org/z/PY5d3WGas). In that case, a device function is allowed to access that static variable even if it does not have __device__ or __constant__ attribute.

I will make changes to handle the function scope static variable.

[yaxunl]

revised by Richard's comments. Check function-scope static var.

[yaxunl]

fix test failure on windows. need to specify triple since it affects name mangling.

[yaxunl]

ping

[tra]

Small test nit. LGTM otherwise.

clang/test/CodeGenCUDA/device-use-host-var.cu

23

Nit: You may want to add a CHECK-LABEL: <function name>. It does not make much of a difference in this file, but is useful in general to explicitly restrict the search scope.

[yaxunl]

clang/test/CodeGenCUDA/device-use-host-var.cu

23

will do

[yaxunl]

Rebase and fix HIP header bug exposed by this patch.

clang/lib/Headers/__clang_hip_math.h

 #endif
 
 #if defined (__cplusplus) && __cplusplus < 201103L
 // emulate static_assert on type sizes
 template<bool>
 struct __compare_result{};
 template<>
 struct __compare_result<true> {
-  static const bool valid;
+  static const __device__ bool valid;
 };
 
 __DEVICE__
 void __suppress_unused_warning(bool b){};
 template <unsigned int S, unsigned int T>
 __DEVICE__ void __static_assert_equal_size() {
   __suppress_unused_warning(__compare_result<S == T>::valid);
 }

clang/lib/Sema/SemaExpr.cpp

bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
 
   DiagnoseAvailabilityOfDecl(D, Locs, UnknownObjCClass, ObjCPropertyAccess,
                              AvoidPartialAvailabilityChecks, ClassReceiver);
 
   DiagnoseUnusedOfDecl(*this, D, Loc);
 
   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 (auto *VD = dyn_cast<ValueDecl>(D))
       checkDeviceDecl(VD, Loc);
 
     if (!Context.getTargetInfo().isTLSSupported())
       if (const auto *VD = dyn_cast<VarDecl>(D))
         if (VD->getTLSKind() != VarDecl::TLS_None)
           targetDiag(*Locs.begin(), diag::err_thread_unsupported);
MarkVarDeclODRUsed(VarDecl *Var, SourceLocation Loc, Sema &SemaRef,
   if (SemaRef.LangOpts.OpenMP)
     SemaRef.tryCaptureOpenMPLambdas(Var);
   SemaRef.tryCaptureVariable(Var, Loc, Sema::TryCapture_Implicit,
     /*EllipsisLoc*/ SourceLocation(),
     /*BuildAndDiagnose*/ true,
     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<FunctionDecl>(SemaRef.CurContext);
+    auto Target = SemaRef.IdentifyCUDATarget(FD);
+    auto IsEmittedOnDeviceSide = [](VarDecl *Var) {
+      if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
+          Var->hasAttr<CUDASharedAttr>() ||
+          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<FunctionDecl>(Var->getDeclContext())) {
+        return FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>();
+      }
+      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);
 }
 
 void Sema::MarkCaptureUsedInEnclosingContext(VarDecl *Capture,
                                              SourceLocation Loc,
                                              unsigned CapturingScopeIndex) {
   MarkVarDeclODRUsed(Capture, Loc, *this, &CapturingScopeIndex);
 }

clang/test/CodeGenCUDA/device-use-host-var.cu

+// 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;
+}

clang/test/Headers/hip-header.hip

 // REQUIRES: amdgpu-registered-target
 // 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__ | 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
 
 // CHECK-LABEL: amdgpu_kernel void @_Z4kernPff
 __global__ void kern(float *x, float y) {
   *x = sin(y);
 }
 
 // CHECK-LABEL: define{{.*}} i64 @_Z11test_size_tv
 // CHEC: ret i64 8
 __device__ size_t test_size_t() {
   return sizeof(size_t);
 }

clang/test/SemaCUDA/device-use-host-var.cu

 // 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;
+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<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}}
+  // 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) {
   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;
+  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;
   *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;
+  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;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_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;
+  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}}
   *out = global_dev_var;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_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;
+  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;
   *out = global_constant_var;
   *out = global_shared_var;
   *out = global_constexpr_var;
   *out = global_const_var;
 

 // Texture references are special. As far as C++ is concerned they are host
 // variables that are referenced from device code. However, they are handled
 // very differently by the compiler under the hood and such references are
 // allowed. Compiler should produce no warning here, but it should diagnose the
 // same case without the device_builtin_texture_type attribute.
 template <class, int = 1, int = 1>
 struct __attribute__((device_builtin_texture_type)) texture {
   static texture<int> ref;
-  __device__ int c() {
+  __device__ void c() {
     auto &x = ref;
   }
 };
 
 template <class, int = 1, int = 1>
 struct  not_a_texture {
   static not_a_texture<int> ref;
-  __device__ int c() {
+  __device__ void c() {
     auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
   }
 };
+
+template<>
+not_a_texture<int> not_a_texture<int>::ref;
+
+__device__ void test_not_a_texture() {
+  not_a_texture<int> inst;
+  inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int, 1, 1>::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;
+      }
+    };
+  }
+}