[Sema][AST] Explicit visibility for OpenCL/CUDA kernels/variables

[scott.linder]

Summary

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.

[tra]

A kernel functions in CUDA is actually two different functions. One is the real kernel we compile for the GPU, another is a host-side stub that launches the device-side kernel.

On device side both clang and nvcc currently silently ignore hidden visibility and force the kernel to always be visible:
https://godbolt.org/z/xrPMGc
This is needed because the kernel must be externally visible in the device-side executable for the host-side code to execute it.
The device-side executable and its symbols are isolated from the host DSO they are encapsulated in, so whether the hidden attribute is ignored or not on device side is independent of the visibility the kernel symbol gets on the host side.

On the host side there's no particular reason to give kernel (or, rather, its host-side stub) a special treatment. Setting hidden on it should be fine, if someone needs it for whatever reason.
Most likely users who may have applied hidden to a kernel would do so in order to avoid exposing kernel symbols outside of a DSO and the attribute will do that job just fine.

I think the warning is not going to buy anything for CUDA. The hidden attribute effectively applies to the host side only, where it should work correctly and where it is potentially useful. I'd rather not impose restrictions that are not necessary, even if it's just a warning.

[Anastasia]

Seems reasonable for OpenCL kernels. You might want to add an AST dump test to check that the visibility is being set correctly in case it's being printed in AST.

[rjmccall]

Okay. So it sounds like this should either be a device-only rule, with no warning in mixed-mode languages like CUDA, or we should take a different approach.

[yaxunl]

lib/AST/Decl.cpp

738

we also need this for __constant__ variables.

[Anastasia]

lib/AST/Decl.cpp

738

And what about __global and __constant program scope variables in OpenCL?

include/clang/Basic/DiagnosticSemaKinds.td

def warn_transparent_union_attribute_zero_fields : Warning<
   "transparent union definition must contain at least one field; "
   "transparent_union attribute ignored">,
   InGroup<IgnoredAttributes>;
 def warn_attribute_type_not_supported : Warning<
   "%0 attribute argument not supported: %1">,
   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">>;
 def err_mismatched_visibility: Error<"visibility does not match previous declaration">;
 def note_previous_attribute : Note<"previous attribute is here">;
 def note_conflicting_attribute : Note<"conflicting attribute is here">;
 def note_attribute : Note<"attribute is here">;
 def err_mismatched_ms_inheritance : Error<

include/clang/Basic/Visibility.h

 inline Visibility minVisibility(Visibility L, Visibility R) {
   return L < R ? L : R;
 }
 
 class LinkageInfo {
   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) {}
   LinkageInfo(Linkage L, Visibility V, bool E)
     : linkage_(L), visibility_(V), explicit_(E) {
     assert(getLinkage() == L && getVisibility() == V &&
            isVisibilityExplicit() == E && "Enum truncated!");
   }
public:
   }
 
   Linkage getLinkage() const { return (Linkage)linkage_; }
   Visibility getVisibility() const { return (Visibility)visibility_; }
   bool isVisibilityExplicit() const { return explicit_; }
 
   void setLinkage(Linkage L) { linkage_ = L; }
 
+  void setVisibility(Visibility V, bool E) { visibility_ = V; explicit_ = E; }
+
   void mergeLinkage(Linkage L) {
     setLinkage(minLinkage(getLinkage(), L));
   }
   void mergeLinkage(LinkageInfo other) {
     mergeLinkage(other.getLinkage());
   }
 
   void mergeExternalVisibility(Linkage L) {

lib/AST/Decl.cpp

if (!LV.isVisibilityExplicit()) {
 
       // If we're paying attention to global visibility, apply
       // -finline-visibility-hidden if this is an inline method.
       if (useInlineVisibilityHidden(D))
         LV.mergeVisibility(HiddenVisibility, /*visibilityExplicit=*/false);
     }
   }
 
+  // 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
   //   name of
   //
   //     - an object or reference, unless it has internal linkage; or
   if (const auto *Var = dyn_cast<VarDecl>(D)) {
     // GCC applies the following optimization to variables and static

lib/CodeGen/TargetInfo.cpp

public:
   createEnqueuedBlockKernel(CodeGenFunction &CGF,
                             llvm::Function *BlockInvokeFunc,
                             llvm::Value *BlockLiteral) const override;
   bool shouldEmitStaticExternCAliases() const override;
   void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
 };
 }
 
-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);
   if (!FD)
     return;
 
   llvm::Function *F = cast<llvm::Function>(GV);
 

lib/Sema/SemaDeclAttr.cpp

void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
   // family, and it can be applied after objc_designated_initializer. This is a
   // bit of a hack, but we need it to be compatible with versions of clang that
   // processed the attribute list in the wrong order.
   if (D->hasAttr<ObjCDesignatedInitializerAttr>() &&
       cast<ObjCMethodDecl>(D)->getMethodFamily() != OMF_init) {
     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.
 void Sema::ProcessDeclAttributeDelayed(Decl *D,
                                        const ParsedAttributesView &AttrList) {
   for (const ParsedAttr &AL : AttrList)
     if (AL.getKind() == ParsedAttr::AT_TransparentUnion) {
       handleTransparentUnionAttr(*this, D, AL);

test/SemaCUDA/visibility-diagnostics.cu

+// 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;

test/SemaOpenCL/visibility-diagnostics.cl

+// 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();