diff options
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 6 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/amdgpu-visibility.cu | 10 |
2 files changed, 13 insertions, 3 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 49df82cea42..be8f389e180 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3575,6 +3575,9 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, } } + if (GV->isDeclaration()) + getTargetCodeGenInfo().setTargetAttributes(D, GV, *this); + LangAS ExpectedAS = D ? D->getType().getAddressSpace() : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default); @@ -3584,9 +3587,6 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace, ExpectedAS, Ty); - if (GV->isDeclaration()) - getTargetCodeGenInfo().setTargetAttributes(D, GV, *this); - return GV; } diff --git a/clang/test/CodeGenCUDA/amdgpu-visibility.cu b/clang/test/CodeGenCUDA/amdgpu-visibility.cu index 9f44eb047f8..f23e562a4f2 100644 --- a/clang/test/CodeGenCUDA/amdgpu-visibility.cu +++ b/clang/test/CodeGenCUDA/amdgpu-visibility.cu @@ -13,6 +13,16 @@ __constant__ int c; __device__ int g; +// CHECK-DEFAULT: @e = external addrspace(1) global +// CHECK-PROTECTED: @e = external protected addrspace(1) global +// CHECK-HIDDEN: @e = external protected addrspace(1) global +extern __device__ int e; + +// dummy one to hold reference to `e`. +__device__ int f() { + return e; +} + // CHECK-DEFAULT: define amdgpu_kernel void @_Z3foov() // CHECK-PROTECTED: define protected amdgpu_kernel void @_Z3foov() // CHECK-HIDDEN: define protected amdgpu_kernel void @_Z3foov() |
