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

[HIP] Fix visibility of `__constant__` variables. 

Summary:
- `__constant__` variables should not be `hidden` as the linker may turn
  them into `LOCAL` symbols.

Reviewers: yaxunl

Subscribers: jvesely, nhaehnle, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61194

llvm-svn: 359344
This commit is contained in:
Michael Liao 2019-04-26 19:31:48 +00:00
parent 9534e9dbe4
commit 3820506960
2 changed files with 23 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/CodeGen/TargetInfo.cpp
View File

@ -7847,7 +7847,8 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D,
return D->hasAttr<OpenCLKernelAttr>() ||
(isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
(isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>());
(isa<VarDecl>(D) &&
(D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()));
}
void AMDGPUTargetCodeGenInfo::setTargetAttributes(

21
clang/test/CodeGenCUDA/amdgpu-visibility.cu Normal file
View File

@ -0,0 +1,21 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility default -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEFAULT %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility protected -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-PROTECTED %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility hidden -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-HIDDEN %s
#include "Inputs/cuda.h"
// CHECK-DEFAULT: @c = addrspace(4) externally_initialized global
// CHECK-DEFAULT: @g = addrspace(1) externally_initialized global
// CHECK-PROTECTED: @c = protected addrspace(4) externally_initialized global
// CHECK-PROTECTED: @g = protected addrspace(1) externally_initialized global
// CHECK-HIDDEN: @c = protected addrspace(4) externally_initialized global
// CHECK-HIDDEN: @g = protected addrspace(1) externally_initialized global
__constant__ int c;
__device__ int g;
// CHECK-DEFAULT: define amdgpu_kernel void @_Z3foov()
// CHECK-PROTECTED: define protected amdgpu_kernel void @_Z3foov()
// CHECK-HIDDEN: define protected amdgpu_kernel void @_Z3foov()
__global__ void foo() {
g = c;
}