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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
yaxunl requested review of this revision.

https://reviews.llvm.org/D98193

Files:
  clang/lib/Sema/SemaExpr.cpp
  clang/test/SemaCUDA/device-use-host-var.cu


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<typename F>
 __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) {
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<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);
@@ -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<FunctionDecl>(SemaRef.CurContext);
+    auto Target = SemaRef.IdentifyCUDATarget(FD);
+    if (FD && Target != Sema::CFT_Host) {
+      if (Var && Var->hasGlobalStorage() && !Var->hasAttr<CUDADeviceAttr>() &&
+          !Var->hasAttr<CUDAConstantAttr>() &&
+          !Var->hasAttr<CUDASharedAttr>() &&
+          !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

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits