diff options
| author | Michael Liao <michael.hliao@gmail.com> | 2019-04-26 19:31:48 +0000 |
|---|---|---|
| committer | Michael Liao <michael.hliao@gmail.com> | 2019-04-26 19:31:48 +0000 |
| commit | 3820506960e058b04078e8301404474164fb2628 (patch) | |
| tree | b02c5d2116bdd049738a6d399a3dc862b4688a7b | |
| parent | 9534e9dbe44c130acefb48bd87b3e7d78db72def (diff) | |
| download | bcm5719-llvm-3820506960e058b04078e8301404474164fb2628.tar.gz bcm5719-llvm-3820506960e058b04078e8301404474164fb2628.zip | |
[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
| -rw-r--r-- | clang/lib/CodeGen/TargetInfo.cpp | 3 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/amdgpu-visibility.cu | 21 |
2 files changed, 23 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 5c2b3ff353e..432e55da411 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -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( diff --git a/clang/test/CodeGenCUDA/amdgpu-visibility.cu b/clang/test/CodeGenCUDA/amdgpu-visibility.cu new file mode 100644 index 00000000000..9f44eb047f8 --- /dev/null +++ b/clang/test/CodeGenCUDA/amdgpu-visibility.cu @@ -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; +} |
