[CUDA][HIP] Fix gpu.used.external

[yaxunl]

Summary

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

[jhuber6]

LGTM

clang/lib/CodeGen/CodeGenModule.cpp

for (auto D : Context.CUDAExternalDeviceDeclODRUsedByHost) {
       UsedArray.push_back(llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
           GetAddrOfGlobal(GD), Int8PtrTy));
     }
 
     llvm::ArrayType *ATy = llvm::ArrayType::get(Int8PtrTy, UsedArray.size());
 
     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);
   }
 
   emitLLVMUsed();
   if (SanStats)
     SanStats->finish();
 
   if (CodeGenOpts.Autolink &&

clang/test/CodeGenCUDA/host-used-extern.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
 // RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 | FileCheck %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
 // RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 \
 // RUN:   | FileCheck -check-prefix=NEG %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
 // RUN:   -std=c++11 -emit-llvm -o - -target-cpu gfx906 \
 // RUN:   | FileCheck -check-prefixes=NEG,NORDC %s
 
 #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();
 
 // kernel2 is not marked as used since it is a definition.
 __global__ void kernel2() {}
 
 // kernel3 is not marked as used since it is not called by host function.
 __global__ void kernel3();