[PATCH] D98814: [CUDA][HIP] Mark device var used by host only

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 332157.
yaxunl marked an inline comment as done.
yaxunl added a comment.

revised by Artem's comments


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D98814/new/

https://reviews.llvm.org/D98814

Files:
  clang/lib/CodeGen/CGCUDANV.cpp
  clang/test/CodeGenCUDA/host-used-device-var.cu


Index: clang/test/CodeGenCUDA/host-used-device-var.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/host-used-device-var.cu
@@ -0,0 +1,46 @@
+// 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 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: @llvm.compiler.used = {{[^@]*}} @_ZL2u3 {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5
Index: clang/lib/CodeGen/CGCUDANV.cpp
===================================================================
--- clang/lib/CodeGen/CGCUDANV.cpp
+++ clang/lib/CodeGen/CGCUDANV.cpp
@@ -1084,6 +1084,28 @@
 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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D98814.332157.patch
Type: text/x-patch
Size: 2803 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210321/a34f655d/attachment.bin>