Move setTargetAttributes after setGVProperties in SetFunctionAttributes

[scott.linder]

Summary

AMDGPU relies on global properties being set before setTargetProperties is called. Existing targets like MIPS which rely on setTargetProperties do not seem to rely on the current behavior, so this patch moves the call later in SetFunctionAttributes.

[rjmccall]

It seems reasonable to me for target hooks to run after global hooks, but can I ask why AMDGPU specifically relies on this?

[scott.linder]

In D60967#1475925, @rjmccall wrote:

It seems reasonable to me for target hooks to run after global hooks, but can I ask why AMDGPU specifically relies on this?

We want to ensure certain symbols have a meaningful visibility. For example, kernel symbols must not have hidden visibility. It's reasonable for the user to arrange for a kernel symbol to have either protected or default visibility, though, so we want our hook to be run after the global hooks have already calculated the global visibility.

[rjmccall]

Shouldn't it be an error if the user tries to give it hidden visibility?

[scott.linder]

In D60967#1476029, @rjmccall wrote:

Shouldn't it be an error if the user tries to give it hidden visibility?

We effectively consider the user explicitly specifying that a symbol is e.g. a kernel to also carry with it visibility information. We don't want to require the user to redundantly specify that a kernel is not hidden, when it is never meaningful for it to be hidden.

[rjmccall]

In D60967#1476057, @scott.linder wrote:

In D60967#1476029, @rjmccall wrote:

Shouldn't it be an error if the user tries to give it hidden visibility?

We effectively consider the user explicitly specifying that a symbol is e.g. a kernel to also carry with it visibility information. We don't want to require the user to redundantly specify that a kernel is not hidden, when it is never meaningful for it to be hidden.

I understand, but if the user explicitly gives it hidden visibility, you should still diagnose that.

Also, shouldn't you just handle this by treating the kernel attribute as a source of explicit visibility at the Sema/AST level?

[atanasyan]

I do not see any problem from MIPS targets point of view.

[scott.linder]

In D60967#1476069, @rjmccall wrote:

In D60967#1476057, @scott.linder wrote:

In D60967#1476029, @rjmccall wrote:

Shouldn't it be an error if the user tries to give it hidden visibility?

We effectively consider the user explicitly specifying that a symbol is e.g. a kernel to also carry with it visibility information. We don't want to require the user to redundantly specify that a kernel is not hidden, when it is never meaningful for it to be hidden.

I understand, but if the user explicitly gives it hidden visibility, you should still diagnose that.

Also, shouldn't you just handle this by treating the kernel attribute as a source of explicit visibility at the Sema/AST level?

I agree that we should diagnose it, and I can update the patch accordingly, but I'm unsure how to go about emitting a diagnostic from this callback. As far as doing this at the AST level, this was my original approach in https://reviews.llvm.org/D53153, however this is really more of an AMDGPU implementation detail. I don't think it is necessarily the case that every OpenCL and Cuda implementation wants/needs require these symbols not have hidden visibility.

If we can involve the target in the AST linkage calculations, or agree that in general the kernel specifier should affect the visibility in this way, along with the __device__ specifier on a variable and the __global__ specifier on a function for Cuda, then moving this up to the AST level makes sense to me.

[rjmccall]

I suspect that other OpenCL and CUDA implementations don't care at all about symbol visibility for device-side code generation, and giving kernel functions default visibility seems like the right thing to do for the (relatively few) things at the AST level that are sensitive to that, like template visibility. Would you mind reaching out to other implementors about that?

This patch seems fine to me regardless.

[scott.linder]

In D60967#1476226, @rjmccall wrote:

I suspect that other OpenCL and CUDA implementations don't care at all about symbol visibility for device-side code generation, and giving kernel functions default visibility seems like the right thing to do for the (relatively few) things at the AST level that are sensitive to that, like template visibility. Would you mind reaching out to other implementors about that?

This patch seems fine to me regardless.

Yes, I can certainly identify who would be interested in terms of OpenCL and Cuda and work on moving this up to the AST.

If you don't object to this patch then is it reasonable for me to submit it? It will get us the required behavior for AMDGPU while I work on the more general solution.

[rjmccall]

Yeah, that's fine.

[scott.linder]

@rjmccall Would you expect similar conflicts in explicit visibility to result in diagnostics? For example, marking a static variable with an explicit visibility attribute doesn't warn, instead the explicit visibility attribute is silently ignored. GCC 7.3 complains with warning: ‘__visibility__’ attribute ignored [-Wattributes]

[rjmccall]

Yeah, that seems like a missing warning.

[scott.linder]

@rjmccall I'm not sure if this is the right place to continue discussing this, but I don't have a patch I am happy with and I would rather not post something half-baked.

Currently for AMDGPU we have the behavior that the user can set the visibility of these symbols with explicit attributes. If we consider the kernel attribute itself as an explicit visibility declaration how do we support this flexibility when we will have effectively mandated a single visibility that the user cannot interact with? Even if we are OK with mandating something like default visibility, we do not currently support preemptible symbols so protected is the optimal visibility. This may not be true of other targets, and it may not even be true of AMDGPU in the future, so hardcoding the visibility of kernel symbols to anything doesn't seem correct. Is something like "not-hidden" reasonable?

[rjmccall]

It seems to fine just forbid hidden. Again, I suspect other targets do not care because they are not using a standard dynamic loader to load the code containing kernel functions.

lib/CodeGen/CodeGenModule.cpp

if (llvm::Intrinsic::ID IID = F->getIntrinsicID()) {
     // If this is an intrinsic function, set the function's attributes
     // to the intrinsic's attributes.
     F->setAttributes(llvm::Intrinsic::getAttributes(getLLVMContext(), IID));
     return;
   }
 
   const auto *FD = cast<FunctionDecl>(GD.getDecl());
 
-  if (!IsIncompleteFunction) {
+  if (!IsIncompleteFunction)
     SetLLVMFunctionAttributes(GD, getTypes().arrangeGlobalDeclaration(GD), F);
-    // Setup target-specific attributes.
-    if (F->isDeclaration())
-      getTargetCodeGenInfo().setTargetAttributes(FD, F, *this);
-  }
 
   // Add the Returned attribute for "this", except for iOS 5 and earlier
   // where substantial code, including the libstdc++ dylib, was compiled with
   // GCC and does not actually return "this".
   if (!IsThunk && getCXXABI().HasThisReturn(GD) &&
       !(getTriple().isiOS() && getTriple().isOSVersionLT(6))) {
     assert(!F->arg_empty() &&
            F->arg_begin()->getType()
              ->canLosslesslyBitCastTo(F->getReturnType()) &&
            "unexpected this return");
     F->addAttribute(1, llvm::Attribute::Returned);
   }
 
   // Only a few attributes are set on declarations; these may later be
   // overridden by a definition.
 
   setLinkageForGV(F, FD);
   setGVProperties(F, FD);
 
+  // Setup target-specific attributes.
+  if (!IsIncompleteFunction && F->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(FD, F, *this);
+
   if (const auto *CSA = FD->getAttr<CodeSegAttr>())
     F->setSection(CSA->getName());
   else if (const auto *SA = FD->getAttr<SectionAttr>())
      F->setSection(SA->getName());
 
   if (FD->isReplaceableGlobalAllocationFunction()) {
     // A replaceable global allocation function does not act like a builtin by
     // default, only if it is invoked by a new-expression or delete-expression.

test/CodeGenOpenCL/visibility.cl

 // FVIS-PROTECTED: define protected void @func_protected()
 // FVIS-HIDDEN: define protected void @func_protected()
 __attribute__((visibility("protected"))) void func_protected() {}
 // FVIS-DEFAULT: define void @func_default()
 // FVIS-PROTECTED: define void @func_default()
 // FVIS-HIDDEN: define void @func_default()
 __attribute__((visibility("default"))) void func_default() {}
 
+extern kernel void ext_kern();
+__attribute__((visibility("hidden"))) extern kernel void ext_kern_hidden();
+__attribute__((visibility("protected"))) extern kernel void ext_kern_protected();
+__attribute__((visibility("default"))) extern kernel void ext_kern_default();
+
+extern void ext_func();
+__attribute__((visibility("hidden"))) extern void ext_func_hidden();
+__attribute__((visibility("protected"))) extern void ext_func_protected();
+__attribute__((visibility("default"))) extern void ext_func_default();
+
 void use() {
     glob = ext + ext_hidden + ext_protected + ext_default;
+    ext_kern();
+    ext_kern_hidden();
+    ext_kern_protected();
+    ext_kern_default();
+    ext_func();
+    ext_func_hidden();
+    ext_func_protected();
+    ext_func_default();
 }
+
+// FVIS-DEFAULT: declare amdgpu_kernel void @ext_kern()
+// FVIS-PROTECTED: declare protected amdgpu_kernel void @ext_kern()
+// FVIS-HIDDEN: declare protected amdgpu_kernel void @ext_kern()
+
+// FVIS-DEFAULT: declare protected amdgpu_kernel void @ext_kern_hidden()
+// FVIS-PROTECTED: declare protected amdgpu_kernel void @ext_kern_hidden()
+// FVIS-HIDDEN: declare protected amdgpu_kernel void @ext_kern_hidden()
+
+// FVIS-DEFAULT: declare protected amdgpu_kernel void @ext_kern_protected()
+// FVIS-PROTECTED: declare protected amdgpu_kernel void @ext_kern_protected()
+// FVIS-HIDDEN: declare protected amdgpu_kernel void @ext_kern_protected()
+
+// FVIS-DEFAULT: declare amdgpu_kernel void @ext_kern_default()
+// FVIS-PROTECTED: declare amdgpu_kernel void @ext_kern_default()
+// FVIS-HIDDEN: declare amdgpu_kernel void @ext_kern_default()
+
+
+// FVIS-DEFAULT: declare void @ext_func()
+// FVIS-PROTECTED: declare protected void @ext_func()
+// FVIS-HIDDEN: declare hidden void @ext_func()
+
+// FVIS-DEFAULT: declare hidden void @ext_func_hidden()
+// FVIS-PROTECTED: declare hidden void @ext_func_hidden()
+// FVIS-HIDDEN: declare hidden void @ext_func_hidden()
+
+// FVIS-DEFAULT: declare protected void @ext_func_protected()
+// FVIS-PROTECTED: declare protected void @ext_func_protected()
+// FVIS-HIDDEN: declare protected void @ext_func_protected()
+
+// FVIS-DEFAULT: declare void @ext_func_default()
+// FVIS-PROTECTED: declare void @ext_func_default()
+// FVIS-HIDDEN: declare void @ext_func_default()