[clang] 301e233 - [CUDA][HIP] Fix static device var used by host code only

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2020-09-23T08:18:19-04:00
New Revision: 301e23305d03cfb4004f845a1d9dfdc5e5931fd8

URL: https://github.com/llvm/llvm-project/commit/301e23305d03cfb4004f845a1d9dfdc5e5931fd8
DIFF: https://github.com/llvm/llvm-project/commit/301e23305d03cfb4004f845a1d9dfdc5e5931fd8.diff

LOG: [CUDA][HIP] Fix static device var used by host code only

A static device variable may be accessed in host code through
cudaMemCpyFromSymbol etc. Currently clang does not
emit the static device variable if it is only referenced by
host code, which causes host code to fail at run time.

This patch fixes that.

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

Added: 
    

Modified: 
    clang/lib/CodeGen/CodeGenModule.cpp
    clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 3ecc8743265c..6a77f6b040a1 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2195,6 +2195,11 @@ void CodeGenModule::EmitDeferred() {
     assert(DeferredVTables.empty());
   }
 
+  // Emit CUDA/HIP static device variables referenced by host code only.
+  if (getLangOpts().CUDA)
+    for (auto V : getContext().CUDAStaticDeviceVarReferencedByHost)
+      DeferredDeclsToEmit.push_back(V);
+
   // Stop if we're out of both deferred vtables and deferred declarations.
   if (DeferredDeclsToEmit.empty())
     return;

diff  --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
index c7beb4c7e1ac..9cb1c6804a48 100644
--- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
+++ b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
@@ -63,6 +63,13 @@ static constexpr int z2 = 456;
 // externalized nor registered.
 // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat
 
+// Check a static device variable referenced by host function only is externalized.
+// DEV-DAG: @_ZL1w = addrspace(1) externally_initialized global i32 0
+// HOST-DAG: @_ZL1w = internal global i32 undef
+// HOST-DAG: @[[DEVNAMEW:[0-9]+]] = {{.*}}c"_ZL1w\00"
+
+static __device__ int w;
+
 inline __device__ void devfun(const int ** b) {
   const static int p = 2;
   b[0] = &p;
@@ -92,11 +99,13 @@ void foo(const int **a) {
   getDeviceSymbol(&x);
   getDeviceSymbol(&x5);
   getDeviceSymbol(&y);
+  getDeviceSymbol(&w);
   z = 123;
   a[0] = &z2;
 }
 
 // HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
 // HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
+// HOST: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]]
 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p