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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
yaxunl requested review of this revision.

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.


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,33 @@
+// 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 device variables used by host functions are kept.
+
+// CHECK: @u1
+__device__ int u1;
+
+// CHECK: @u2
+__constant__ int u2;
+
+// Check device variables with used attribute are always kept.
+
+// CHECK: @u3
+__device__ __attribute__((used)) int u3;
+
+int fun1() {
+  return u1 + u2;
+}
+
+// CHECK: @llvm.compiler.used = {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u3
Index: clang/lib/CodeGen/CGCUDANV.cpp
===================================================================
--- clang/lib/CodeGen/CGCUDANV.cpp
+++ clang/lib/CodeGen/CGCUDANV.cpp
@@ -1084,6 +1084,24 @@
 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.
+    for (auto &&Info : DeviceVars) {
+      auto Kind = Info.Flags.getKind();
+      if (!Info.Var->isDeclaration() &&
+          (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.331365.patch
Type: text/x-patch
Size: 2194 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210317/9f8e5562/attachment-0001.bin>