Hahnfeld created this revision.
Hahnfeld added a reviewer: tra.
Herald added a subscriber: cfe-commits.
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!
Repository:
rC Clang
https://reviews.llvm.org/D47201
Files:
include/clang/Basic/Attr.td
include/clang/Basic/DiagnosticSemaKinds.td
lib/CodeGen/CodeGenModule.cpp
lib/Sema/SemaDecl.cpp
lib/Sema/SemaDeclAttr.cpp
test/CodeGenCUDA/nv_weak.cu
test/SemaCUDA/attr-declspec.cu
test/SemaCUDA/attr-nv_weak.cu
test/SemaCUDA/attributes-on-non-cuda.cu
Index: test/SemaCUDA/attributes-on-non-cuda.cu
===================================================================
--- test/SemaCUDA/attributes-on-non-cuda.cu
+++ test/SemaCUDA/attributes-on-non-cuda.cu
@@ -7,11 +7,12 @@
// 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.
@@ -24,11 +25,11 @@
__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;
Index: test/SemaCUDA/attr-nv_weak.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/attr-nv_weak.cu
@@ -0,0 +1,14 @@
+// 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}}
Index: test/SemaCUDA/attr-declspec.cu
===================================================================
--- test/SemaCUDA/attr-declspec.cu
+++ test/SemaCUDA/attr-declspec.cu
@@ -6,11 +6,12 @@
// 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.)
@@ -23,12 +24,11 @@
__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.
Index: test/CodeGenCUDA/nv_weak.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/nv_weak.cu
@@ -0,0 +1,22 @@
+// 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();
+}
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -6179,6 +6179,9 @@
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;
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -5942,10 +5942,16 @@
// '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);
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -1374,7 +1374,7 @@
// "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);
}
@@ -3442,7 +3442,7 @@
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
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -2757,7 +2757,7 @@
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<
Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -1508,12 +1508,11 @@
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 {
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
