maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Yaxun (Sam) Liu 2022-04-18 11:08:50 -04:00
parent 80787213ea
commit cac4e2fe25
2 changed files with 10 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -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);
}

18
clang/test/CodeGenCUDA/host-used-extern.cu
View File

@ -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
// CHECK-LABEL: @llvm.compiler.used = {{.*}} @__clang_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
// 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();