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

[yaxunl]

Summary

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.

clang/lib/CodeGen/CodeGenModule.cpp

if (!DeferredVTables.empty()) {
     EmitDeferredVTables();
 
     // Emitting a vtable doesn't directly cause more vtables to
     // become deferred, although it can cause functions to be
     // emitted that then need those vtables.
     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;
 
   // Grab the list of decls to emit. If EmitGlobalDefinition schedules more
   // work, it will not interfere with this.
   std::vector<GlobalDecl> CurDeclsToEmit;
   CurDeclsToEmit.swap(DeferredDeclsToEmit);

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

 // DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456
 
 static constexpr int z2 = 456;
 
 // Test static device variable in inline function, which should not be
 // 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;
   b[1] = &x2;
 }
 
 __global__ void kernel(int *a, const int **b) {
   const static int w = 1;

 }
 
 int* getDeviceSymbol(int *x);
 
 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