[CUDA][HIP] Diagnose reference of host variable

[yaxunl]

Summary

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

[yaxunl]

add comments about H->D not diagnosed

[tra]

LGTM.

clang/lib/Sema/SemaExpr.cpp

357

This could use a comment why we only check D->H references.

[yaxunl]

clang/lib/Sema/SemaExpr.cpp

357

Just added a comment:

Reference of device global variables in host functions is
allowed through shadow variables therefore it is not diagnosed.

[tra]

LGTM.

[tra]

It appears that we need to add special handling for texture/surface references. Nominally they are host-side objects, but they are accessed/used from device functions as far as Sema is concerned.

E.g. https://godbolt.org/z/z1YnE3
NVCC and older clang compile it, but the recent one fails because this patch does not let the change through.
I think isCUDADeviceBuiltinTextureType has problem handling texture refs within templates.

[tra]

In D91281#2441147, @tra wrote:

I think isCUDADeviceBuiltinTextureType has problem handling texture refs within templates.

Proposed fix: https://reviews.llvm.org/D92893

[yaxunl]

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.

[yaxunl]

In D91281#2607082, @yaxunl wrote:

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.

It seems we should only diagnose ODR-use (https://en.cppreference.com/w/cpp/language/definition) of host variable in device functions.

[yaxunl]

In D91281#2609766, @yaxunl wrote:

In D91281#2607082, @yaxunl wrote:

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.

It seems we should only diagnose ODR-use (https://en.cppreference.com/w/cpp/language/definition) of host variable in device functions.

Fix in https://reviews.llvm.org/D98193

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_config_scalar_return : Error<
   "CUDA special function '%0' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;
 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">;
 def err_capture_bad_target : Error<
   "capture host variable %0 by reference in device or host device lambda function">;
 def err_capture_bad_target_this_ptr : Error<
   "capture host side class data member by this pointer in device or host device lambda function">;
 def warn_kern_is_method : Extension<

clang/lib/Sema/SemaCUDA.cpp

bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   // Avoid emitting this error twice for the same location.  Using a hashtable
   // like this is unfortunate, but because we must continue parsing as normal
   // after encountering a deferred error, it's otherwise very tricky for us to
   // ensure that we only emit this deferred error once.
   if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
     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)
         << Callee;
   return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
          DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
 }
 

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 (const 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);

clang/test/CodeGenCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // Make sure we handle target overloads correctly.  Most of this is checked in
 // sema, but special functions like constructors and destructors are here.
 //
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
 
 #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; }
   __device__ s_cd_dh() { x = 12; }
 };
 
 struct s_cd_hd {
   __host__ __device__ s_cd_hd() { x = 31; }

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

__attribute__((device)) void test_shared64() {
   val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_SEQ_CST, "workgroup");
 
   // CHECK: %2 = load i64, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
   // CHECK-NEXT: %3 = call i64 @llvm.amdgcn.atomic.dec.i64.p0i64(i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), i64 %2, i32 7, i32 2, i1 false)
   // CHECK-NEXT: store i64 %3, i64* addrspacecast (i64 addrspace(3)* @_ZZ13test_shared64vE3val to i64*), align 8
   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
   // CHECK-NEXT: %1 = call i32 @llvm.amdgcn.atomic.inc.i32.p0i32(i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), i32 %0, i32 7, i32 2, i1 false)
   // CHECK-NEXT: store i32 %1, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
   global_val32 = __builtin_amdgcn_atomic_inc32(&global_val32, global_val32, __ATOMIC_SEQ_CST, "workgroup");
 
   // CHECK: %2 = load i32, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
   // CHECK-NEXT: %3 = call i32 @llvm.amdgcn.atomic.dec.i32.p0i32(i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), i32 %2, i32 7, i32 2, i1 false)
   // CHECK-NEXT: store i32 %3, i32* addrspacecast (i32 addrspace(1)* @global_val32 to i32*), align 4
   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
   // CHECK-NEXT: %1 = call i64 @llvm.amdgcn.atomic.inc.i64.p0i64(i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), i64 %0, i32 7, i32 2, i1 false)
   // CHECK-NEXT: store i64 %1, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8
   global_val64 = __builtin_amdgcn_atomic_inc64(&global_val64, global_val64, __ATOMIC_SEQ_CST, "workgroup");
 
   // CHECK: %2 = load i64, i64* addrspacecast (i64 addrspace(1)* @global_val64 to i64*), align 8

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;
+__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;
+  });
+}
+