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(); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
