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

[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
This commit is contained in:
Michael Liao 2019-06-07 15:08:29 -04:00
parent 6cd47f9dd7
commit 0a220de9e9
2 changed files with 13 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

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

10
clang/test/CodeGenCUDA/amdgpu-visibility.cu
View File

@ -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()