[clang] cac4e2f - [CUDA][HIP] Fix gpu.used.external
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 18 20:10:58 PDT 2022
Author: Yaxun (Sam) Liu
Date: 2022-04-18T23:10:31-04:00
New Revision: cac4e2fe258239489c12dd5f153da3e69e17ebba
URL: https://github.com/llvm/llvm-project/commit/cac4e2fe258239489c12dd5f153da3e69e17ebba
DIFF: https://github.com/llvm/llvm-project/commit/cac4e2fe258239489c12dd5f153da3e69e17ebba.diff
LOG: [CUDA][HIP] Fix gpu.used.external
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
Reviewed by: Joseph Huber, Artem Belevich
Differential Revision: https://reviews.llvm.org/D123946
Added:
Modified:
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGenCUDA/host-used-extern.cu
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index bf5f2d1b42719..9f6e2e5a9f52d 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -599,7 +599,7 @@ void CodeGenModule::Release() {
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);
}
diff --git a/clang/test/CodeGenCUDA/host-used-extern.cu b/clang/test/CodeGenCUDA/host-used-extern.cu
index 02b55cdb76b4c..f430111b566a2 100644
--- a/clang/test/CodeGenCUDA/host-used-extern.cu
+++ b/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();
More information about the cfe-commits
mailing list