[CUDA] Implement nv_weak attribute for functions

[Hahnfeld]

Summary

This is needed for relocatable device code with CUDA 9 and later.
Before this patch, linking two or more object files resulted in
"Multiple definition" errors for a group of functions from
cuda_device_runtime_api.h which are annoted with "nv_weak".

CUDA headers already used this attribute in earlier releases, but
until CUDA 8.0 the only definitions in cuda_device_runtime_api.h
were conditional under defined(__CUDABE__) which is explicitly
undefined in Clang's wrapper. However since CUDA 9.0 this has
changed to !defined(__CUDACC_RTC__). Trying to add that define
resulted in errors that nvrtc_device_runtime.h could not be found.

Reported by Andrea Bocci!

[Hahnfeld]

Ping

[tra]

IIUIC, nv_weak is a synonym for weak (<rant>why, oh why did they need it?</rant>)
You may need to hunt down and change few other places that deal with the weak attribute.
E.g.: https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/AST/Decl.cpp#L4267
https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/CodeGen/ItaniumCXXABI.cpp#L3045

[aaron.ballman]

In D47201#1119249, @tra wrote:

IIUIC, nv_weak is a synonym for weak (<rant>why, oh why did they need it?</rant>)
You may need to hunt down and change few other places that deal with the weak attribute.
E.g.: https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/AST/Decl.cpp#L4267
https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/CodeGen/ItaniumCXXABI.cpp#L3045

If it is truly a synonym for weak, then a better implementation would be to make no semantic distinction between the two attributes -- just add new spellings to weak. If you need to make minor distinctions between the spellings, you can do it using accessors on the attribute.

include/clang/Basic/Attr.td

1515

No new, undocumented attributes, please.

[Hahnfeld]

In D47201#1119254, @aaron.ballman wrote:

In D47201#1119249, @tra wrote:

IIUIC, nv_weak is a synonym for weak (<rant>why, oh why did they need it?</rant>)
You may need to hunt down and change few other places that deal with the weak attribute.
E.g.: https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/AST/Decl.cpp#L4267
https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/CodeGen/ItaniumCXXABI.cpp#L3045

If it is truly a synonym for weak, then a better implementation would be to make no semantic distinction between the two attributes -- just add new spellings to weak. If you need to make minor distinctions between the spellings, you can do it using accessors on the attribute.

I first went with this approach but then thought it would be better to restrict the new attribute as much as possible. That's why I added a completely new one which is only applicable to functions, but not to variables and CXXRecords. Let me know if you'd prefer nv_weak to be a full alias of weak and I'll revert to what @aaron.ballman suggested.

[aaron.ballman]

In D47201#1119947, @Hahnfeld wrote:

In D47201#1119254, @aaron.ballman wrote:

In D47201#1119249, @tra wrote:

IIUIC, nv_weak is a synonym for weak (<rant>why, oh why did they need it?</rant>)
You may need to hunt down and change few other places that deal with the weak attribute.
E.g.: https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/AST/Decl.cpp#L4267
https://github.com/llvm-project/llvm-project-20170507/blob/master/clang/lib/CodeGen/ItaniumCXXABI.cpp#L3045

If it is truly a synonym for weak, then a better implementation would be to make no semantic distinction between the two attributes -- just add new spellings to weak. If you need to make minor distinctions between the spellings, you can do it using accessors on the attribute.

I first went with this approach but then thought it would be better to restrict the new attribute as much as possible. That's why I added a completely new one which is only applicable to functions, but not to variables and CXXRecords. Let me know if you'd prefer nv_weak to be a full alias of weak and I'll revert to what @aaron.ballman suggested.

I don't know enough about nv_weak's semantics to definitively say one way or the other -- I can find no documentation on this attribute (official or otherwise). However, based purely on the changes made here, I'd likely add the accessors and only go with a single semantic attribute. You can check for the proper subjects by looking at the parsed attribute kind in SemaDeclAttr.cpp and restricting the subjects there (it's a bit less declarative this way, but we don't have a way to map subject lists to spellings like we do for accessors because it's not a common need).

[tra]

I've experimented a bit and I think that we may not need this patch at all.
As far as I can tell, nv_weak is only applicable to device functions. It's ignored for global kernels and is apparently forbidden for data.
For device functions nvcc produces .weak attribute in PTX.

Using plain old attribute((weak)) does exactly the same. Considering that nv_weak is only used inside CUDA SDK headers, substituting weak in place of nv_weak will result in correct PTX, which is all we really need. I don't see much benefit turning it into full blown attribute just to mimic an internal CUDA implementation detail we don't care all that much about.

Now, replacing it in CUDA headers for all CUDA versions we support may be tricky. Let me give it a try. I'll send a patch, if I manage to make it work.

[Hahnfeld]

See D47804

include/clang/Basic/Attr.td

 }
 
 def NoThrow : InheritableAttr {
   let Spellings = [GCC<"nothrow">, Declspec<"nothrow">];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [NoThrowDocs];
 }
 
-def NvWeak : IgnoredAttr {
-  // No Declspec spelling of this attribute; the CUDA headers use
-  // __attribute__((nv_weak)) unconditionally. Does not receive an [[]]
-  // spelling because it is a CUDA attribute.
-  let Spellings = [GNU<"nv_weak">];
+def NvWeak : InheritableAttr {
+  let Spellings = [GNU<"nv_weak">, Declspec<"__nv_weak__">];
+  let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
+  let Documentation = [Undocumented];
 }
 
 def ObjCBridge : InheritableAttr {
   let Spellings = [Clang<"objc_bridge">];
   let Subjects = SubjectList<[Record, TypedefName], ErrorDiag>;
   let Args = [IdentifierArgument<"BridgedType">];
   let Documentation = [Undocumented];
 }

include/clang/Basic/DiagnosticSemaKinds.td

def warn_ptr_independentclass_attribute : Warning<
   InGroup<IndependentClassAttribute>;
 def warn_attribute_weak_on_local : Warning<
   "__weak attribute cannot be specified on an automatic variable when ARC "
   "is not enabled">,
   InGroup<IgnoredAttributes>;
 def warn_weak_identifier_undeclared : Warning<
   "weak identifier %0 never declared">;
 def err_attribute_weak_static : Error<
-  "weak declaration cannot have internal linkage">;
+  "%select{weak|nv_weak}0 declaration cannot have internal linkage">;
 def err_attribute_selectany_non_extern_data : Error<
   "'selectany' can only be applied to data items with external linkage">;
 def err_declspec_thread_on_thread_variable : Error<
   "'__declspec(thread)' applied to variable that already has a "
   "thread-local storage specifier">;
 def err_attribute_dll_not_extern : Error<
   "%q0 must have external linkage when declared %q1">;
 def err_attribute_dll_thread_local : Error<

lib/CodeGen/CodeGenModule.cpp

 
 static void setLinkageForGV(llvm::GlobalValue *GV, const NamedDecl *ND) {
   // Set linkage and visibility in case we never see a definition.
   LinkageInfo LV = ND->getLinkageAndVisibility();
   // Don't set internal linkage on declarations.
   // "extern_weak" is overloaded in LLVM; we probably should have
   // separate linkage types for this.
   if (isExternallyVisible(LV.getLinkage()) &&
-      (ND->hasAttr<WeakAttr>() || ND->isWeakImported()))
+      (ND->hasAttr<WeakAttr>() || ND->hasAttr<NvWeakAttr>() || ND->isWeakImported()))
     GV->setLinkage(llvm::GlobalValue::ExternalWeakLinkage);
 }
 
 void CodeGenModule::CreateFunctionTypeMetadata(const FunctionDecl *FD,
                                                llvm::Function *F) {
   // Only if we are checking indirect calls.
   if (!LangOpts.Sanitize.has(SanitizerKind::CFIICall))
     return;
static bool isVarDeclStrongDefinition(const ASTContext &Context,
   return false;
 }
 
 llvm::GlobalValue::LinkageTypes CodeGenModule::getLLVMLinkageForDeclarator(
     const DeclaratorDecl *D, GVALinkage Linkage, bool IsConstantVariable) {
   if (Linkage == GVA_Internal)
     return llvm::Function::InternalLinkage;
 
-  if (D->hasAttr<WeakAttr>()) {
+  if (D->hasAttr<WeakAttr>() || D->hasAttr<NvWeakAttr>()) {
     if (IsConstantVariable)
       return llvm::GlobalVariable::WeakODRLinkage;
     else
       return llvm::GlobalVariable::WeakAnyLinkage;
   }
 
   // We are guaranteed to have a strong definition somewhere else,
   // so we can use available_externally linkage.

lib/Sema/SemaDecl.cpp

 static void checkAttributesAfterMerging(Sema &S, NamedDecl &ND) {
   // Ensure that an auto decl is deduced otherwise the checks below might cache
   // the wrong linkage.
   assert(S.ParsingInitForAutoVars.count(&ND) == 0);
 
   // 'weak' only applies to declarations with external linkage.
   if (WeakAttr *Attr = ND.getAttr<WeakAttr>()) {
     if (!ND.isExternallyVisible()) {
-      S.Diag(Attr->getLocation(), diag::err_attribute_weak_static);
+      S.Diag(Attr->getLocation(), diag::err_attribute_weak_static) << 0;
       ND.dropAttr<WeakAttr>();
     }
   }
+  if (NvWeakAttr *Attr = ND.getAttr<NvWeakAttr>()) {
+    if (!ND.isExternallyVisible()) {
+      S.Diag(Attr->getLocation(), diag::err_attribute_weak_static) << 1;
+      ND.dropAttr<NvWeakAttr>();
+    }
+  }
   if (WeakRefAttr *Attr = ND.getAttr<WeakRefAttr>()) {
     if (ND.isExternallyVisible()) {
       S.Diag(Attr->getLocation(), diag::err_attribute_weakref_not_static);
       ND.dropAttr<WeakRefAttr>();
       ND.dropAttr<AliasAttr>();
     }
   }
 

lib/Sema/SemaDeclAttr.cpp

case AttributeList::AT_WarnUnused:
     handleSimpleAttribute<WarnUnusedAttr>(S, D, AL);
     break;
   case AttributeList::AT_WarnUnusedResult:
     handleWarnUnusedResult(S, D, AL);
     break;
   case AttributeList::AT_Weak:
     handleSimpleAttribute<WeakAttr>(S, D, AL);
     break;
+  case AttributeList::AT_NvWeak:
+    handleSimpleAttribute<NvWeakAttr>(S, D, AL);
+    break;
   case AttributeList::AT_WeakRef:
     handleWeakRefAttr(S, D, AL);
     break;
   case AttributeList::AT_WeakImport:
     handleWeakImportAttr(S, D, AL);
     break;
   case AttributeList::AT_TransparentUnion:
     handleTransparentUnionAttr(S, D, AL);

test/CodeGenCUDA/nv_weak.cu

+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN:     -o - %s | FileCheck %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm \
+// RUN:     -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-DAG: declare extern_weak i32 @_Z2f1v()
+extern
+#if defined(__CUDA_ARCH__)
+__device__
+#endif
+int f1() __attribute__((nv_weak));
+
+// CHECK-DAG: define weak i32 @_Z2f2v()
+#if defined(__CUDA_ARCH__)
+__device__
+#endif
+int f2() __attribute__((nv_weak)) {
+  return f1();
+}

test/SemaCUDA/attr-declspec.cu

 // Test the __declspec spellings of CUDA attributes.
 //
 // RUN: %clang_cc1 -fsyntax-only -fms-extensions -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fms-extensions -fcuda-is-device -verify %s
 // Now pretend that we're compiling a C file. There should be warnings.
 // RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s
 
 #if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'__device__' attribute ignored}}
-// expected-warning@+12 {{'__global__' attribute ignored}}
-// expected-warning@+12 {{'__constant__' attribute ignored}}
-// expected-warning@+12 {{'__shared__' attribute ignored}}
-// expected-warning@+12 {{'__host__' attribute ignored}}
+// expected-warning@+13 {{'__device__' attribute ignored}}
+// expected-warning@+13 {{'__global__' attribute ignored}}
+// expected-warning@+13 {{'__constant__' attribute ignored}}
+// expected-warning@+13 {{'__shared__' attribute ignored}}
+// expected-warning@+13 {{'__host__' attribute ignored}}
+// expected-warning@+13 {{'__nv_weak__' attribute ignored}}
 //
 // (Currently we don't for the other attributes. They are implemented with
 // IgnoredAttr, which is ignored irrespective of any LangOpts.)
 #else
 // expected-no-diagnostics
 #endif
 
 __declspec(__device__) void f_device();
 __declspec(__global__) void f_global();
 __declspec(__constant__) int* g_constant;
 __declspec(__shared__) float *g_shared;
 __declspec(__host__) void f_host();
+__declspec(__nv_weak__) void f_nv_weak();
 __declspec(__device_builtin__) void f_device_builtin();
 typedef __declspec(__device_builtin__) const void *t_device_builtin;
 enum __declspec(__device_builtin__) e_device_builtin {E};
 __declspec(__device_builtin__) int v_device_builtin;
 __declspec(__cudart_builtin__) void f_cudart_builtin();
 __declspec(__device_builtin_surface_type__) unsigned long long surface_var;
 __declspec(__device_builtin_texture_type__) unsigned long long texture_var;
-
-// Note that there's no __declspec spelling of nv_weak.

test/SemaCUDA/attr-nv_weak.cu

+// RUN: %clang_cc1 -verify -fsyntax-only %s
+
+extern int f0() __attribute__((nv_weak));
+extern int g0 __attribute__((nv_weak)); // expected-warning {{'nv_weak' attribute only applies to functions}}
+int f1() __attribute__((nv_weak));
+int g1 __attribute__((nv_weak)); // expected-warning {{'nv_weak' attribute only applies to functions}}
+
+
+struct __attribute__((nv_weak)) s0 {}; // expected-warning {{'nv_weak' attribute only applies to functions}}
+
+static int f() __attribute__((nv_weak)); // expected-error {{nv_weak declaration cannot have internal linkage}}
+
+static void pr14946_f();
+void pr14946_f() __attribute__((nv_weak)); // expected-error {{nv_weak declaration cannot have internal linkage}}

test/SemaCUDA/attributes-on-non-cuda.cu

 // Tests that CUDA attributes are warnings when compiling C files, but not when
 // compiling CUDA files.
 //
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
 // Now pretend that we're compiling a C file. There should be warnings.
 // RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s
 
 #if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'device' attribute ignored}}
-// expected-warning@+12 {{'global' attribute ignored}}
-// expected-warning@+12 {{'constant' attribute ignored}}
-// expected-warning@+12 {{'shared' attribute ignored}}
-// expected-warning@+12 {{'host' attribute ignored}}
+// expected-warning@+13 {{'device' attribute ignored}}
+// expected-warning@+13 {{'global' attribute ignored}}
+// expected-warning@+13 {{'constant' attribute ignored}}
+// expected-warning@+13 {{'shared' attribute ignored}}
+// expected-warning@+13 {{'host' attribute ignored}}
+// expected-warning@+13 {{'nv_weak' attribute ignored}}
 //
 // NOTE: IgnoredAttr in clang which is used for the rest of
 // attributes ignores LangOpts, so there are no warnings.
 #else
 // expected-no-diagnostics
 #endif
 
 __attribute__((device)) void f_device();
 __attribute__((global)) void f_global();
 __attribute__((constant)) int* g_constant;
 __attribute__((shared)) float *g_shared;
 __attribute__((host)) void f_host();
+__attribute__((nv_weak)) void f_nv_weak();
 __attribute__((device_builtin)) void f_device_builtin();
 typedef __attribute__((device_builtin)) const void *t_device_builtin;
 enum __attribute__((device_builtin)) e_device_builtin {E};
 __attribute__((device_builtin)) int v_device_builtin;
 __attribute__((cudart_builtin)) void f_cudart_builtin();
-__attribute__((nv_weak)) void f_nv_weak();
 __attribute__((device_builtin_surface_type)) unsigned long long surface_var;
 __attribute__((device_builtin_texture_type)) unsigned long long texture_var;