scott.linder created this revision.
scott.linder added reviewers: Anastasia, tra, yaxunl, rjmccall.
Herald added subscribers: cfe-commits, tpr.
Herald added a project: clang.
For AMDGPU the visibility of these symbols (OpenCL kernels, CUDA `__global__`
functions, and CUDA `__device__` variables) must not be hidden, as we rely on
them being available in the dynamic symbol table in the final DSO.
This patch implements this by considering language attributes as a source of
explicit visibility, but rather than attributing any one visibility to them
they are simply coerced to be a non-hidden visibility. This allows for the
optimization of using protected visibility when these symbols are known to be
`dso_local`.
This patch also adds diagnostics for explicitly setting a hidden visibility on
these symbols.
I imagine there are a number of issues with the patch in its current state, but
I wanted to get something implemented before reaching out to OpenCL/CUDA
maintainers to see if this is a reasonable change. @Anastasia and @tra I wasn't
certain if you would be good candidates to discuss this change, so please let
me know if I need to keep looking.
Repository:
rC Clang
https://reviews.llvm.org/D61274
Files:
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Basic/Visibility.h
lib/AST/Decl.cpp
lib/CodeGen/TargetInfo.cpp
lib/Sema/SemaDeclAttr.cpp
test/SemaCUDA/visibility-diagnostics.cu
test/SemaOpenCL/visibility-diagnostics.cl
Index: test/SemaOpenCL/visibility-diagnostics.cl
===================================================================
--- /dev/null
+++ test/SemaOpenCL/visibility-diagnostics.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -std=cl2.0 -verify -pedantic -fsyntax-only %s
+
+__attribute__((visibility("hidden"))) kernel void kern_hidden() {} // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) kernel void kern_protected();
+__attribute__((visibility("default"))) kernel void kern_default();
+kernel void kern();
+
+__attribute__((visibility("hidden"))) extern kernel void ext_kern_hidden(); // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) extern kernel void ext_kern_protected();
+__attribute__((visibility("default"))) extern kernel void ext_kern_default();
+extern kernel void ext_kern();
Index: test/SemaCUDA/visibility-diagnostics.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/visibility-diagnostics.cu
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+__attribute__((visibility("hidden"))) __global__ void global_func_hidden() {} // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) __global__ void global_func_protected() {}
+__attribute__((visibility("default"))) __global__ void global_func_default() {}
+__global__ void global_func() {}
+
+__attribute__((visibility("hidden"))) __device__ int device_var_hidden; // expected-warning {{'hidden' visibility on variable with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) __device__ int device_var_protected;
+__attribute__((visibility("default"))) __device__ int device_var_default;
+__device__ int device_var;
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -7375,6 +7375,16 @@
Diag(D->getLocation(), diag::err_designated_init_attr_non_init);
D->dropAttr<ObjCDesignatedInitializerAttr>();
}
+
+ if ((D->hasAttr<VisibilityAttr>() &&
+ D->getAttr<VisibilityAttr>()->getVisibility() ==
+ VisibilityAttr::Hidden) &&
+ (D->hasAttr<OpenCLKernelAttr>() ||
+ (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
+ (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>()))) {
+ Diag(D->getLocation(), diag::warn_attribute_hidden_visibility)
+ << (isa<FunctionDecl>(D) ? 0 : 1);
+ }
}
// Helper for delayed processing TransparentUnion attribute.
Index: lib/CodeGen/TargetInfo.cpp
===================================================================
--- lib/CodeGen/TargetInfo.cpp
+++ lib/CodeGen/TargetInfo.cpp
@@ -7840,23 +7840,8 @@
};
}
-static bool requiresAMDGPUProtectedVisibility(const Decl *D,
- llvm::GlobalValue *GV) {
- if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
- return false;
-
- return D->hasAttr<OpenCLKernelAttr>() ||
- (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
- (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>());
-}
-
void AMDGPUTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
- if (requiresAMDGPUProtectedVisibility(D, GV)) {
- GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
- GV->setDSOLocal(true);
- }
-
if (GV->isDeclaration())
return;
const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);
Index: lib/AST/Decl.cpp
===================================================================
--- lib/AST/Decl.cpp
+++ lib/AST/Decl.cpp
@@ -731,6 +731,19 @@
}
}
+ // We consider OpenCL kernels, __global__ Cuda functions, and __device__ Cuda
+ // variables to have explicit visibility of "non-hidden".
+ if (D->hasAttr<OpenCLKernelAttr>() ||
+ (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
+ (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>())) {
+ Visibility Vis = LV.getVisibility();
+ if (LV.getVisibility() == HiddenVisibility)
+ Vis = Context.getTargetInfo().hasProtectedVisibility()
+ ? ProtectedVisibility
+ : DefaultVisibility;
+ LV.setVisibility(Vis, true);
+ }
+
// C++ [basic.link]p4:
// A name having namespace scope has external linkage if it is the
Index: include/clang/Basic/Visibility.h
===================================================================
--- include/clang/Basic/Visibility.h
+++ include/clang/Basic/Visibility.h
@@ -53,8 +53,6 @@
uint8_t linkage_ : 3;
uint8_t visibility_ : 2;
uint8_t explicit_ : 1;
-
- void setVisibility(Visibility V, bool E) { visibility_ = V; explicit_ = E; }
public:
LinkageInfo() : linkage_(ExternalLinkage), visibility_(DefaultVisibility),
explicit_(false) {}
@@ -86,6 +84,8 @@
void setLinkage(Linkage L) { linkage_ = L; }
+ void setVisibility(Visibility V, bool E) { visibility_ = V; explicit_ = E; }
+
void mergeLinkage(Linkage L) {
setLinkage(minLinkage(getLinkage(), L));
}
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -3362,6 +3362,9 @@
InGroup<IgnoredAttributes>;
def warn_attribute_unknown_visibility : Warning<"unknown visibility %0">,
InGroup<IgnoredAttributes>;
+def warn_attribute_hidden_visibility :
+ Warning<"'hidden' visibility on %select{function|variable}0 with incompatible language attribute will be ignored">,
+ InGroup<IgnoredAttributes>;
def warn_attribute_protected_visibility :
Warning<"target does not support 'protected' visibility; using 'default'">,
InGroup<DiagGroup<"unsupported-visibility">>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
