This is an archive of the discontinued LLVM Phabricator instance.

Differential D61194

[HIP] Fix visibility of `__constant__` variables.
ClosedPublic

Authored by hliao on Apr 26 2019, 8:27 AM.

Details

Reviewers
yaxunl
Commits
rG3820506960e0: [HIP] Fix visibility of `__constant__` variables.
rL359344: [HIP] Fix visibility of `__constant__` variables.
rC359344: [HIP] Fix visibility of `__constant__` variables.
Summary
  • __constant__ variables should not be hidden as the linker may turn them into LOCAL symbols.

Diff Detail

Event Timeline

hliao created this revision.Apr 26 2019, 8:27 AM
Herald added a project: Restricted Project. · View Herald TranscriptApr 26 2019, 8:27 AM
Herald added subscribers: cfe-commits, nhaehnle, jvesely. · View Herald Transcript
Harbormaster completed remote builds in B31035: Diff 196857.Apr 26 2019, 8:27 AM
yaxunl added inline comments.Apr 26 2019, 11:37 AM
clang/lib/CodeGen/TargetInfo.cpp
7851

is format right?

clang/test/CodeGenCUDA/amdgpu-visibility.cu
1

do we need -fapply-global-visibility-to-externs?

hliao marked an inline comment as done.Apr 26 2019, 11:40 AM
hliao added inline comments.
clang/lib/CodeGen/TargetInfo.cpp
7851

yeah, it's the format after clang-format. It seems that clang-format try to apply the same indent for LHS & RHS of a binary operator. With the leading parenthesis, RHS in a newline is indented with one more space.

hliao marked an inline comment as done.Apr 26 2019, 11:44 AM
hliao added inline comments.
clang/test/CodeGenCUDA/amdgpu-visibility.cu
1

that's so far the option applied for both AMDGPU & HIP (targetting non-MS) by default in clang frontend. Check AMDGPUToolChain::addClangTargetOptions and HIPToolChain::addClangTargetOptions

yaxunl accepted this revision.Apr 26 2019, 12:25 PM

LGTM. Thanks!

This revision is now accepted and ready to land.Apr 26 2019, 12:25 PM
Closed by commit rC359344: [HIP] Fix visibility of `__constant__` variables. (authored by hliao). · Explain WhyApr 26 2019, 12:30 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
CodeGen/
3 lines
test/
CodeGenCUDA/
21 lines

Diff 196857

clang/lib/CodeGen/TargetInfo.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 7,841 Lines • ▼ Show 20 Lines
static bool requiresAMDGPUProtectedVisibility(const Decl *D, static bool requiresAMDGPUProtectedVisibility(const Decl *D,
llvm::GlobalValue *GV) { llvm::GlobalValue *GV) {
if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility) if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
return false; return false;
return D->hasAttr<OpenCLKernelAttr>() || return D->hasAttr<OpenCLKernelAttr>() ||
(isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) || (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
(isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>()); (isa<VarDecl>(D) &&
(D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()));
yaxunlUnsubmitted
Not Done
ReplyInline Actions

is format right?

yaxunl: is format right?
hliaoAuthorUnsubmitted
Done
ReplyInline Actions

yeah, it's the format after clang-format. It seems that clang-format try to apply the same indent for LHS & RHS of a binary operator. With the leading parenthesis, RHS in a newline is indented with one more space.

hliao: yeah, it's the format after clang-format. It seems that clang-format try to apply the same…
} }
void AMDGPUTargetCodeGenInfo::setTargetAttributes( void AMDGPUTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
if (requiresAMDGPUProtectedVisibility(D, GV)) { if (requiresAMDGPUProtectedVisibility(D, GV)) {
GV->setVisibility(llvm::GlobalValue::ProtectedVisibility); GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
GV->setDSOLocal(true); GV->setDSOLocal(true);
} }
▲ Show 20 Lines • Show All 1,779 Lines • Show Last 20 Lines

clang/test/CodeGenCUDA/amdgpu-visibility.cu

  • This file was added.
// 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
yaxunlUnsubmitted
Not Done
ReplyInline Actions

do we need -fapply-global-visibility-to-externs?

yaxunl: do we need -fapply-global-visibility-to-externs?
hliaoAuthorUnsubmitted
Done
ReplyInline Actions

that's so far the option applied for both AMDGPU & HIP (targetting non-MS) by default in clang frontend. Check AMDGPUToolChain::addClangTargetOptions and HIPToolChain::addClangTargetOptions

hliao: that's so far the option applied for both AMDGPU & HIP (targetting non-MS) by default in clang…
// 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;
}