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

[tra]

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

41

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

41

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

17042

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

17042

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

22 ↗

(On Diff #335904)

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

22 ↗

(On Diff #335904)

will do

[yaxunl]

Rebase and fix HIP header bug exposed by this patch.

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);
SemaRef.tryCaptureVariable(Var, Loc, Sema::TryCapture_Implicit,
     FunctionScopeIndexToStopAt);
 
   Var->markUsed(SemaRef.Context);
 }
 
 void Sema::MarkCaptureUsedInEnclosingContext(VarDecl *Capture,
                                              SourceLocation Loc,
                                              unsigned CapturingScopeIndex) {
   MarkVarDeclODRUsed(Capture, Loc, *this, &CapturingScopeIndex);
 }
 
 static void
 diagnoseUncapturableValueReference(Sema &S, SourceLocation loc,
                                    ValueDecl *var, DeclContext *DC) {
   DeclContext *VarDC = var->getDeclContext();
 
   //  If the parameter still belongs to the translation unit, then
if (RefersToEnclosingScope) {
         assert(E && "Capture variable should be used in an expression.");
         if (!Var->getType()->isReferenceType() ||
             !Var->isUsableInConstantExpressions(SemaRef.Context))
           LSI->addPotentialCapture(E->IgnoreParens());
       }
     }
     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
 /// (C++ [basic.def.odr]p2, C99 6.9p3).  Note that this should not be
 /// used directly for normal expressions referring to VarDecl.
 void Sema::MarkVariableReferenced(SourceLocation Loc, VarDecl *Var) {
   DoMarkVarDeclReferenced(*this, Loc, Var, nullptr);
 }

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"
 
+struct A {
+  int x;
+  A() {}
+};
+
 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;
 
+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}}
 
 __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_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;
   *out = ref_constant_var;
   *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) {
   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;