Index: clang/include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticParseKinds.td +++ clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1230,6 +1230,8 @@ "duplicate attribute subject matcher '%0'">; def err_pragma_attribute_expected_period : Error< "expected '.' after pragma attribute namespace %0">; +def err_pragma_attribute_np_expected_period : Error< + "expected '.' after pragma attribute no_pop %0">; def err_pragma_attribute_namespace_on_attribute : Error< "namespace can only apply to 'push' or 'pop' directives">; def note_pragma_attribute_namespace_on_attribute : Note< 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; + /// Allow pragma attribute push without matching pop at eof. + bool AllowNoPopEof; SmallVector Entries; }; @@ -10017,7 +10019,8 @@ SourceLocation PragmaLoc, attr::ParsedSubjectMatchRuleSet Rules); void ActOnPragmaAttributeEmptyPush(SourceLocation PragmaLoc, - const IdentifierInfo *Namespace); + const IdentifierInfo *Namespace, + bool AllowNoPopEof); /// Called on well-formed '\#pragma clang attribute pop'. void ActOnPragmaAttributePop(SourceLocation PragmaLoc, Index: clang/lib/Parse/ParsePragma.cpp =================================================================== --- clang/lib/Parse/ParsePragma.cpp +++ clang/lib/Parse/ParsePragma.cpp @@ -1299,6 +1299,7 @@ ParsedAttributes &Attributes; ActionType Action; const IdentifierInfo *Namespace = nullptr; + bool AllowNoPopEof = false; ArrayRef Tokens; PragmaAttributeInfo(ParsedAttributes &Attributes) : Attributes(Attributes) {} @@ -1563,7 +1564,8 @@ if (Info->Action == PragmaAttributeInfo::Push && Info->Tokens.empty()) { ConsumeAnnotationToken(); - Actions.ActOnPragmaAttributeEmptyPush(PragmaLoc, Info->Namespace); + Actions.ActOnPragmaAttributeEmptyPush(PragmaLoc, Info->Namespace, + Info->AllowNoPopEof); return; } @@ -1716,7 +1718,8 @@ // Handle a mixed push/attribute by desurging to a push, then an attribute. if (Info->Action == PragmaAttributeInfo::Push) - Actions.ActOnPragmaAttributeEmptyPush(PragmaLoc, Info->Namespace); + Actions.ActOnPragmaAttributeEmptyPush(PragmaLoc, Info->Namespace, + Info->AllowNoPopEof); Actions.ActOnPragmaAttributeAttribute(Attribute, PragmaLoc, std::move(SubjectMatchRules)); @@ -3525,6 +3528,13 @@ /// #pragma clang attribute pop /// \endcode /// +/// There is an optional 'no_pop' variant of push directives, where un-popped +/// push directives do not throw an error at the end of TU. +/// \code +/// #pragma clang attribute no_pop.push (attribute, subject-set) +/// #pragma clang attribute no_pop.push +/// \endcode +/// /// 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 @@ -3550,6 +3560,23 @@ auto *Info = new (PP.getPreprocessorAllocator()) PragmaAttributeInfo(AttributesForPragmaAttribute); + // Parse the optional no_pop followed by a period. + if (Tok.is(tok::identifier)) { + IdentifierInfo *II = Tok.getIdentifierInfo(); + if (II->isStr("no_pop")) { + Info->AllowNoPopEof = true; + PP.Lex(Tok); + + if (!Tok.is(tok::period)) { + PP.Diag(Tok.getLocation(), + diag::err_pragma_attribute_np_expected_period) + << II; + return; + } + PP.Lex(Tok); + } + } + // Parse the optional namespace followed by a period. if (Tok.is(tok::identifier)) { IdentifierInfo *II = Tok.getIdentifierInfo(); Index: clang/lib/Sema/SemaAttr.cpp =================================================================== --- clang/lib/Sema/SemaAttr.cpp +++ clang/lib/Sema/SemaAttr.cpp @@ -929,10 +929,12 @@ } void Sema::ActOnPragmaAttributeEmptyPush(SourceLocation PragmaLoc, - const IdentifierInfo *Namespace) { + const IdentifierInfo *Namespace, + bool AllowNoPopEof) { PragmaAttributeStack.emplace_back(); PragmaAttributeStack.back().Loc = PragmaLoc; PragmaAttributeStack.back().Namespace = Namespace; + PragmaAttributeStack.back().AllowNoPopEof = AllowNoPopEof; } void Sema::ActOnPragmaAttributePop(SourceLocation PragmaLoc, @@ -1009,7 +1011,11 @@ void Sema::DiagnoseUnterminatedPragmaAttribute() { if (PragmaAttributeStack.empty()) return; - Diag(PragmaAttributeStack.back().Loc, diag::err_pragma_attribute_no_pop_eof); + + // Don't error if the remaining push allows no_pop. + if (!PragmaAttributeStack.back().AllowNoPopEof) + Diag(PragmaAttributeStack.back().Loc, + diag::err_pragma_attribute_no_pop_eof); } void Sema::ActOnPragmaOptimize(bool On, SourceLocation PragmaLoc) { Index: clang/test/SemaCUDA/force-device-globals.cu =================================================================== --- /dev/null +++ clang/test/SemaCUDA/force-device-globals.cu @@ -0,0 +1,99 @@ +// 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_before_pragma = 3; +#pragma clang attribute push (__device__, apply_to = variable(is_global)) + static int local_static = 3; +#pragma clang attribute pop + a += local_static_constexpr; + a += local_constexpr; + a += local_static_before_pragma; + 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_before_pragma = 3; +#pragma clang attribute push (__device__, apply_to = variable(is_global)) + static int local_static = 3; +#pragma clang attribute pop + a += local_static_constexpr; + a += local_constexpr; + a += local_static_before_pragma; + 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 no_pop.push (__device__, apply_to = variable(is_global))