[PATCH] D123946: [CUDA][HIP] Fix gpu.used.external

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, jhuber6.
Herald added a subscriber: mattd.
Herald added a project: All.
yaxunl requested review of this revision.

Rename  gpu.used.external as __clang_gpu_used_external as ptxas does not
allow . in global variable name.

Fixes: https://github.com/llvm/llvm-project/issues/54934


https://reviews.llvm.org/D123946

Files:
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/CodeGenCUDA/host-used-extern.cu


Index: clang/test/CodeGenCUDA/host-used-extern.cu
===================================================================
--- clang/test/CodeGenCUDA/host-used-extern.cu
+++ clang/test/CodeGenCUDA/host-used-extern.cu
@@ -11,19 +11,19 @@
 
 #include "Inputs/cuda.h"
 
-// CHECK-LABEL: @gpu.used.external = appending {{.*}}global
+// CHECK-LABEL: @__clang_gpu_used_external = appending {{.*}}global
 // CHECK-DAG: @_Z7kernel1v
 // CHECK-DAG: @_Z7kernel4v
 // CHECK-DAG: @var1
-// CHECK-LABEL: @llvm.compiler.used = {{.*}} @gpu.used.external
-
-// NEG-NOT: @gpu.used.external = {{.*}} @_Z7kernel2v
-// NEG-NOT: @gpu.used.external = {{.*}} @_Z7kernel3v
-// NEG-NOT: @gpu.used.external = {{.*}} @var2
-// NEG-NOT: @gpu.used.external = {{.*}} @var3
-// NORDC-NOT: @gpu.used.external = {{.*}} @_Z7kernel1v
-// NORDC-NOT: @gpu.used.external = {{.*}} @_Z7kernel4v
-// NORDC-NOT: @gpu.used.external = {{.*}} @var1
+// CHECK-LABEL: @llvm.compiler.used = {{.*}} @__clang_gpu_used_external
+
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel2v
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel3v
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @var2
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @var3
+// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel1v
+// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel4v
+// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @var1
 
 __global__ void kernel1();
 
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -599,7 +599,7 @@
 
     auto *GV = new llvm::GlobalVariable(
         getModule(), ATy, false, llvm::GlobalValue::AppendingLinkage,
-        llvm::ConstantArray::get(ATy, UsedArray), "gpu.used.external");
+        llvm::ConstantArray::get(ATy, UsedArray), "__clang_gpu_used_external");
     addCompilerUsedGlobal(GV);
   }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D123946.423401.patch
Type: text/x-patch
Size: 1944 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220418/9635a4d5/attachment-0001.bin>