Index: clang/docs/LanguageExtensions.rst =================================================================== --- clang/docs/LanguageExtensions.rst +++ clang/docs/LanguageExtensions.rst @@ -3379,54 +3379,79 @@ =========================================================================== The ``#pragma clang attribute`` directive can be used to apply an attribute to -multiple declarations. The ``#pragma clang attribute push`` variation of the -directive pushes a new "scope" of ``#pragma clang attribute`` that attributes -can be added to. The ``#pragma clang attribute (...)`` variation adds an -attribute to that scope, and the ``#pragma clang attribute pop`` variation pops -the scope. You can also use ``#pragma clang attribute push (...)``, which is a -shorthand for when you want to add one attribute to a new scope. Multiple push -directives can be nested inside each other. +multiple declarations. When the directive is used outside of a scope created by +a ``#pragma clang attribute push`` directive, the attribute will be added +to the global scope. The attributes that are used in the ``#pragma clang attribute`` directives can be written using the GNU-style syntax: .. code-block:: c++ - #pragma clang attribute push (__attribute__((annotate("custom"))), apply_to = function) + #pragma clang attribute (__attribute__((annotate("custom"))), apply_to = function) - void function(); // The function now has the annotate("custom") attribute + void function1(); // These functions now both ... - #pragma clang attribute pop + void function2(); // ... have the annotate("custom") attribute. The attributes can also be written using the C++11 style syntax: .. code-block:: c++ - #pragma clang attribute push ([[noreturn]], apply_to = function) - - void function(); // The function now has the [[noreturn]] attribute + #pragma clang attribute ([[noreturn]], apply_to = function) - #pragma clang attribute pop + void function(); // The function now has the [[noreturn]] attribute. The ``__declspec`` style syntax is also supported: .. code-block:: c++ - #pragma clang attribute push (__declspec(dllexport), apply_to = function) + #pragma clang attribute (__declspec(dllexport), apply_to = function) - void function(); // The function now has the __declspec(dllexport) attribute + void function(); // The function now has the __declspec(dllexport) attribute. - #pragma clang attribute pop +The ``#pragma clang attribute push`` variation of the directive pushes a new +"scope" of ``#pragma clang attribute`` that attributes can be added to. The +``#pragma clang attribute pop`` variation pops the scope. Multiple push +directives can be nested inside each other. Attributes added to outer scopes +continue to apply in inner scopes. + +.. code-block:: c++ + + #pragma clang attribute push + #pragma clang attribute (__attribute__((annotate("custom1"))), apply_to = function) + + void function1(); // This function has the annotate("custom1") attribute. + + #pragma clang attribute push + #pragma clang attribute (__attribute__((annotate("custom2"))), apply_to = function) + + void function2(); // This function has both annotate("custom1") and annotate("custom2") attributes. -A single push directive accepts only one attribute regardless of the syntax -used. + #pragma clang attribute pop // pops "custom2" scope + + #pragma clang attribute pop // pops "custom1" scope + + void function3(); // This function does not have the annotate("custom1") or annotate("custom2") attributes. + +You can also use ``#pragma clang attribute push (...)``, which is a shorthand +for when you want to add one attribute to a new scope. A single push directive +accepts only one attribute regardless of the attribute syntax used. + +.. code-block:: c++ + + #pragma clang attribute push (__attribute__((annotate("custom"))), apply_to = function) + + void function(); // The function now has the annotate("custom") attribute. + + #pragma clang attribute pop Because multiple push directives can be nested, if you're writing a macro that expands to ``_Pragma("clang attribute")`` it's good hygiene (though not required) to add a namespace to your push/pop directives. A pop directive with a namespace will pop the innermost push that has that same namespace. This will ensure that another macro's ``pop`` won't inadvertently pop your attribute. Note -that an ``pop`` without a namespace will pop the innermost ``push`` without a +that a ``pop`` without a namespace will pop the innermost ``push`` without a namespace. ``push``es with a namespace can only be popped by ``pop`` with the same namespace. For instance: Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -948,13 +948,10 @@ InGroup; def note_pragma_attribute_region_ends_here : Note< "'#pragma clang attribute push' regions ends here">; -def err_pragma_attribute_no_pop_eof : Error<"unterminated " +def warn_pragma_attribute_no_pop_eof : Warning<"unterminated " "'#pragma clang attribute push' at end of file">; def note_pragma_attribute_applied_decl_here : Note< "when applied to this declaration">; -def err_pragma_attr_attr_no_push : Error< - "'#pragma clang attribute' attribute with no matching " - "'#pragma clang attribute push'">; /// Objective-C parser diagnostics def err_duplicate_class_def : Error< Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -742,6 +742,8 @@ SourceLocation Loc; /// The namespace of this push group. const IdentifierInfo *Namespace; + /// Flag indicating if this is the global group. + bool IsGlobal; SmallVector Entries; }; @@ -10017,6 +10019,10 @@ /// the appropriate attribute. void AddCFAuditedAttribute(Decl *D); + /// AddPragmaAttributeGlobalOnce - Add a Global Attribute Group, which + /// allows usage of '\#pragma clang attribute (...)' without push/pop. + void AddPragmaAttributeGlobalOnce(); + void ActOnPragmaAttributeAttribute(ParsedAttr &Attribute, SourceLocation PragmaLoc, attr::ParsedSubjectMatchRuleSet Rules); Index: clang/lib/Parse/ParsePragma.cpp =================================================================== --- clang/lib/Parse/ParsePragma.cpp +++ clang/lib/Parse/ParsePragma.cpp @@ -1555,6 +1555,8 @@ "Expected #pragma attribute annotation token"); SourceLocation PragmaLoc = Tok.getLocation(); auto *Info = static_cast(Tok.getAnnotationValue()); + // Add the Global attribute when namespace push/pop is unused. + Actions.AddPragmaAttributeGlobalOnce(); if (Info->Action == PragmaAttributeInfo::Pop) { ConsumeAnnotationToken(); Actions.ActOnPragmaAttributePop(PragmaLoc, Info->Namespace); @@ -3523,6 +3525,7 @@ /// /// The syntax is: /// \code +/// #pragma clang attribute (attribute, subject-set) /// #pragma clang attribute push (attribute, subject-set) /// #pragma clang attribute push /// #pragma clang attribute (attribute, subject-set) @@ -3532,7 +3535,7 @@ /// There are also 'namespace' variants of push and pop directives. The bare /// '#pragma clang attribute (attribute, subject-set)' version doesn't require a /// namespace, since it always applies attributes to the most recently pushed -/// group, regardless of namespace. +/// or global group, regardless of namespace. /// \code /// #pragma clang attribute namespace.push (attribute, subject-set) /// #pragma clang attribute namespace.push Index: clang/lib/Sema/SemaAttr.cpp =================================================================== --- clang/lib/Sema/SemaAttr.cpp +++ clang/lib/Sema/SemaAttr.cpp @@ -795,6 +795,14 @@ } // end anonymous namespace +void Sema::AddPragmaAttributeGlobalOnce() { + if (!PragmaAttributeStack.empty()) + return; + PragmaAttributeStack.emplace_back(); + PragmaAttributeStack.back().Namespace = nullptr; + PragmaAttributeStack.back().IsGlobal = true; +} + void Sema::ActOnPragmaAttributeAttribute( ParsedAttr &Attribute, SourceLocation PragmaLoc, attr::ParsedSubjectMatchRuleSet Rules) { @@ -919,11 +927,6 @@ Diagnostic << attrMatcherRuleListToString(ExtraRules); } - if (PragmaAttributeStack.empty()) { - Diag(PragmaLoc, diag::err_pragma_attr_attr_no_push); - return; - } - PragmaAttributeStack.back().Entries.push_back( {PragmaLoc, &Attribute, std::move(SubjectMatchRules), /*IsUsed=*/false}); } @@ -933,11 +936,13 @@ PragmaAttributeStack.emplace_back(); PragmaAttributeStack.back().Loc = PragmaLoc; PragmaAttributeStack.back().Namespace = Namespace; + PragmaAttributeStack.back().IsGlobal = false; } void Sema::ActOnPragmaAttributePop(SourceLocation PragmaLoc, const IdentifierInfo *Namespace) { - if (PragmaAttributeStack.empty()) { + // Only the global namespace exists, pop without matching push. + if (PragmaAttributeStack.size() == 1) { Diag(PragmaLoc, diag::err_pragma_attribute_stack_mismatch) << 1; return; } @@ -948,7 +953,8 @@ // namespace. for (size_t Index = PragmaAttributeStack.size(); Index;) { --Index; - if (PragmaAttributeStack[Index].Namespace == Namespace) { + if (PragmaAttributeStack[Index].Namespace == Namespace && + !PragmaAttributeStack[Index].IsGlobal) { for (const PragmaAttributeEntry &Entry : PragmaAttributeStack[Index].Entries) { if (!Entry.IsUsed) { @@ -971,8 +977,7 @@ } void Sema::AddPragmaAttributes(Scope *S, Decl *D) { - if (PragmaAttributeStack.empty()) - return; + // Expected that at least Global Attribute Group exists. for (auto &Group : PragmaAttributeStack) { for (auto &Entry : Group.Entries) { ParsedAttr *Attribute = Entry.Attribute; @@ -1007,9 +1012,10 @@ } void Sema::DiagnoseUnterminatedPragmaAttribute() { - if (PragmaAttributeStack.empty()) + // There can be a Global Pragma Attribute. + if (PragmaAttributeStack.size() <= 1) return; - Diag(PragmaAttributeStack.back().Loc, diag::err_pragma_attribute_no_pop_eof); + Diag(PragmaAttributeStack.back().Loc, diag::warn_pragma_attribute_no_pop_eof); } void Sema::ActOnPragmaOptimize(bool On, SourceLocation PragmaLoc) { Index: clang/test/Misc/warning-flags.c =================================================================== --- clang/test/Misc/warning-flags.c +++ clang/test/Misc/warning-flags.c @@ -18,7 +18,7 @@ The list of warnings below should NEVER grow. It should gradually shrink to 0. -CHECK: Warnings without flags (68): +CHECK: Warnings without flags (69): CHECK-NEXT: ext_expected_semi_decl_list CHECK-NEXT: ext_explicit_specialization_storage_class @@ -73,6 +73,7 @@ CHECK-NEXT: warn_pp_convert_to_positive CHECK-NEXT: warn_pp_expr_overflow CHECK-NEXT: warn_pp_line_decimal +CHECK-NEXT: warn_pragma_attribute_no_pop_eof CHECK-NEXT: warn_pragma_pack_pop_identifier_and_alignment CHECK-NEXT: warn_pragma_pack_show CHECK-NEXT: warn_property_getter_owning_mismatch Index: clang/test/Parser/pragma-attribute.cpp =================================================================== --- clang/test/Parser/pragma-attribute.cpp +++ clang/test/Parser/pragma-attribute.cpp @@ -112,7 +112,7 @@ #pragma clang attribute push (__attribute__((annotate("test"))), apply_to = function) () // expected-warning {{extra tokens at end of '#pragma clang attribute'}} // expected-error@-1 {{expected unqualified-id}} -// expected-error@-2 {{unterminated '#pragma clang attribute push' at end of file}} +// expected-warning@-2 {{unterminated '#pragma clang attribute push' at end of file}} #pragma clang attribute pop () // expected-warning {{extra tokens at end of '#pragma clang attribute'}} Index: clang/test/Sema/pragma-attribute-global.c =================================================================== --- /dev/null +++ clang/test/Sema/pragma-attribute-global.c @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s +// expected-no-diagnostics + +// Valid to add attribute to global attribute group. +#pragma clang attribute (__attribute__((annotate("func"))), apply_to = function) + +void function(); Index: clang/test/Sema/pragma-attribute.c =================================================================== --- clang/test/Sema/pragma-attribute.c +++ clang/test/Sema/pragma-attribute.c @@ -59,12 +59,10 @@ #pragma clang attribute pop #pragma clang attribute pop // expected-error{{'#pragma clang attribute pop' with no matching '#pragma clang attribute push'}} -#pragma clang attribute (__attribute__((annotate)), apply_to = function) // expected-error{{'#pragma clang attribute' attribute with no matching '#pragma clang attribute push}} - #pragma clang attribute push ([[]], apply_to = function) // A noop #pragma clang attribute pop // expected-error {{'#pragma clang attribute pop' with no matching '#pragma clang attribute push'}} -#pragma clang attribute push (__attribute__((annotate("func"))), apply_to = function) // expected-error {{unterminated '#pragma clang attribute push' at end of file}} +#pragma clang attribute push (__attribute__((annotate("func"))), apply_to = function) // expected-warning {{unterminated '#pragma clang attribute push' at end of file}} void function(); Index: clang/test/SemaCUDA/force-device-globals.cu =================================================================== --- /dev/null +++ clang/test/SemaCUDA/force-device-globals.cu @@ -0,0 +1,92 @@ +// RUN: %clang_cc1 -std=c++14 %s -o - -triple nvptx64-nvidia-cuda \ +// RUN: -fcuda-is-device -verify -fsyntax-only +#include "Inputs/cuda.h" + +static int global_before_pragma = 1; + +#pragma clang attribute pop +// expected-error@-1 {{'#pragma clang attribute pop' with no matching '#pragma clang attribute push'}} + +#pragma clang attribute push (__device__, apply_to = variable(is_global)) + +static int global = 1; +static const int const_global = 1; + +struct S { + static const int static_const_field = 1; + static int static_field; +}; + +namespace NS1 { + namespace NS2 { + int ns_global = 1; + } + + struct S { + static const int ns1_static_const_field = 1; + static int ns1_static_field; + }; +} + +int S::static_field = 1; +int NS1::S::ns1_static_field = 1; + +#pragma clang attribute pop + +__device__ void device_func(int &a) { + static constexpr int local_static_constexpr = 1; + constexpr int local_constexpr = 2; + static int local_static = 3; + a += local_static_constexpr; + a += local_constexpr; + a += local_static; + a += global_before_pragma; + // expected-error@-1 {{reference to __host__ variable 'global_before_pragma' in __device__ function}} + a += global; + a += const_global; + a += S::static_const_field; + a += S::static_field; + a += NS1::S::ns1_static_const_field; + a += NS1::S::ns1_static_field; + a += NS1::NS2::ns_global; +} + +void host_func(int &a) { + static constexpr int local_static_constexpr = 1; + constexpr int local_constexpr = 2; + static int local_static = 3; + a += local_static_constexpr; + a += local_constexpr; + a += local_static; + a += global_before_pragma; + // Note: variables below should not be allowed to read from host. + // Clang allows reading their shadows, but that is not OK. + a += global; + a += const_global; + a += S::static_const_field; + a += S::static_field; + a += NS1::S::ns1_static_const_field; + a += NS1::S::ns1_static_field; + a += NS1::NS2::ns_global; +} + +__host__ __device__ void host_device_func(int &a) { + static constexpr int local_static_constexpr = 1; + constexpr int local_constexpr = 2; + static int local_static = 3; + a += local_static_constexpr; + a += local_constexpr; + a += local_static; + a += global_before_pragma; + // expected-error@-1 {{reference to __host__ variable 'global_before_pragma' in __host__ __device__ function}} + a += global; + a += const_global; + a += S::static_const_field; + a += S::static_field; + a += NS1::S::ns1_static_const_field; + a += NS1::S::ns1_static_field; + a += NS1::NS2::ns_global; +} + +#pragma clang attribute push (__device__, apply_to = variable(is_global)) +// expected-warning@-1 {{unterminated '#pragma clang attribute push' at end of file}}