From 0a220de9e9ca3e6786df6c03fd37668815805c62 Mon Sep 17 00:00:00 2001 From: Michael Liao Date: Fri, 7 Jun 2019 15:08:29 -0400 Subject: [PATCH] [HIP] Fix visibility for 'extern' device variables. Summary: - Fix a bug which misses the change for a variable to be set with target-specific attributes. Reviewers: yaxunl Subscribers: jvesely, nhaehnle, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D63020 --- clang/lib/CodeGen/CodeGenModule.cpp | 6 +++--- 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 49df82c..be8f389 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 9f44eb0..f23e562 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() -- 2.7.4