[clang] 5c8911d - [CUDA][HIP] Diagnose reference of host variable

Yaxun Liu via cfe-commits cfe-commits at lists.llvm.org
Wed Dec 2 07:28:42 PST 2020


Author: Yaxun (Sam) Liu
Date: 2020-12-02T10:15:56-05:00
New Revision: 5c8911d0ba3862119d2507aa55b94766263be13b

URL: https://github.com/llvm/llvm-project/commit/5c8911d0ba3862119d2507aa55b94766263be13b
DIFF: https://github.com/llvm/llvm-project/commit/5c8911d0ba3862119d2507aa55b94766263be13b.diff

LOG: [CUDA][HIP] Diagnose reference of host variable

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

Differential Revision: https://reviews.llvm.org/D91281

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

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/Sema/SemaCUDA.cpp
    clang/lib/Sema/SemaExpr.cpp
    clang/test/CodeGenCUDA/function-overload.cu
    clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index f2b2b1d3ab6f..3067c077ddb2 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8145,7 +8145,7 @@ 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">;

diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 12a28ab392f8..0f06adf38f7a 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -743,7 +743,8 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
     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
index 9c2fc1b9e6dd..527605ac4fb8 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -354,6 +354,24 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
 
   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);

diff  --git a/clang/test/CodeGenCUDA/function-overload.cu b/clang/test/CodeGenCUDA/function-overload.cu
index c82b2e96f6c3..9677a5b43b8c 100644
--- 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
index 77ea3d485c8a..16600d15f2c4 100644
--- a/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
+++ b/clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
@@ -124,7 +124,7 @@ __attribute__((device)) void test_shared64() {
   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 @@ __attribute__((device)) void test_global32() {
   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
index 000000000000..cf5514610a42
--- /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<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 at -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;
+  });
+}
+


        


More information about the cfe-commits mailing list