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 { 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: 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() || ND->isWeakImported())) + (ND->hasAttr() || ND->hasAttr() || ND->isWeakImported())) GV->setLinkage(llvm::GlobalValue::ExternalWeakLinkage); } @@ -3442,7 +3442,7 @@ if (Linkage == GVA_Internal) return llvm::Function::InternalLinkage; - if (D->hasAttr()) { + if (D->hasAttr() || D->hasAttr()) { if (IsConstantVariable) return llvm::GlobalVariable::WeakODRLinkage; else 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()) { if (!ND.isExternallyVisible()) { - S.Diag(Attr->getLocation(), diag::err_attribute_weak_static); + S.Diag(Attr->getLocation(), diag::err_attribute_weak_static) << 0; ND.dropAttr(); } } + if (NvWeakAttr *Attr = ND.getAttr()) { + if (!ND.isExternallyVisible()) { + S.Diag(Attr->getLocation(), diag::err_attribute_weak_static) << 1; + ND.dropAttr(); + } + } if (WeakRefAttr *Attr = ND.getAttr()) { if (ND.isExternallyVisible()) { S.Diag(Attr->getLocation(), diag::err_attribute_weakref_not_static); Index: lib/Sema/SemaDeclAttr.cpp =================================================================== --- lib/Sema/SemaDeclAttr.cpp +++ lib/Sema/SemaDeclAttr.cpp @@ -6179,6 +6179,9 @@ case AttributeList::AT_Weak: handleSimpleAttribute(S, D, AL); break; + case AttributeList::AT_NvWeak: + handleSimpleAttribute(S, D, AL); + break; case AttributeList::AT_WeakRef: handleWeakRefAttr(S, D, AL); break; 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: 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,6 +24,7 @@ __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}; @@ -30,5 +32,3 @@ __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/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/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;