[CUDA][HIP] Mark device var used by host only

[yaxunl]

Summary

Add device variables to llvm.compiler.used if they are
ODR-used by either host or device functions.

This is necessary to prevent them from being
eliminated by whole-program optimization
where the compiler has no way to know a device
variable is used by some host code.

[tra]

clang/lib/CodeGen/CGCUDANV.cpp

1107

Do we want to limit it further to only externally-visible variables?
I think we already externalize the variables we want to be visible across host/device boundary.
If the variable is not visible, there's no point keeping it around as the runtime will not be able to find it in the GPU binary.

[yaxunl]

clang/lib/CodeGen/CGCUDANV.cpp

1107

Good point. Will do

[yaxunl]

revised by Artem's comments

[tra]

clang/test/CodeGenCUDA/host-used-device-var.cu

32–34

I'd rephrase it as 'but is not in llvm.compiler.used'

47

I'd add a comment that we're effectively matching the exact list of the variables here and that ensures that @_ZL2u4 is not among them.

[yaxunl]

clang/test/CodeGenCUDA/host-used-device-var.cu

32–34

will do

47

will do

clang/lib/CodeGen/CGCUDANV.cpp

for (auto &&Info : DeviceVars) {
     }
   }
 }
 
 // Returns module constructor to be added.
 llvm::Function *CGNVCUDARuntime::finalizeModule() {
   if (CGM.getLangOpts().CUDAIsDevice) {
     transformManagedVars();
+
+    // Mark ODR-used device variables as compiler used to prevent it from being
+    // eliminated by optimization. This is necessary for device variables
+    // ODR-used by host functions. Sema correctly marks them as ODR-used no
+    // matter whether they are ODR-used by device or host functions.
+    //
+    // We do not need to do this if the variable has used attribute since it
+    // has already been added.
+    //
+    // Static device variables have been externalized at this point, therefore
+    // variables with LLVM private or internal linkage need not be added.
+    for (auto &&Info : DeviceVars) {
+      auto Kind = Info.Flags.getKind();
+      if (!Info.Var->isDeclaration() &&
+          !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) &&
+          (Kind == DeviceVarFlags::Variable ||
+           Kind == DeviceVarFlags::Surface ||
+           Kind == DeviceVarFlags::Texture) &&
+          Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) {
+        CGM.addCompilerUsedGlobal(Info.Var);
+      }
+    }
     return nullptr;
   }
   return makeModuleCtorFunction();
 }
 
 llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
                                                     GlobalDecl GD) {
   auto Loc = KernelHandles.find(F);

clang/test/CodeGenCUDA/host-used-device-var.cu

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
+// RUN:   | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check device variables used by neither host nor device functioins are not kept.
+
+// CHECK-NOT: @v1
+__device__ int v1;
+
+// CHECK-NOT: @v2
+__constant__ int v2;
+
+// CHECK-NOT: @_ZL2v3
+static __device__ int v3;
+
+// Check device variables used by host functions are kept.
+
+// CHECK-DAG: @u1
+__device__ int u1;
+
+// CHECK-DAG: @u2
+__constant__ int u2;
+
+// Check host-used static device var is in llvm.compiler.used.
+// CHECK-DAG: @_ZL2u3
+static __device__ int u3;
+
+// Check device-used static device var is emitted but is not in llvm.compiler.used.
+// CHECK-DAG: @_ZL2u4
+static __device__ int u4;
+
+// Check device variables with used attribute are always kept.
+// CHECK-DAG: @u5
+__device__ __attribute__((used)) int u5;
+
+int fun1() {
+  return u1 + u2 + u3;
+}
+
+__global__ void kern1(int **x) {
+  *x = &u4;
+}
+// Check the exact list of variables to ensure @_ZL2u4 is not among them.
+// CHECK: @llvm.compiler.used = {{[^@]*}} @_ZL2u3 {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5