From 0251bfb8d236ae234340f63844d8a9891165c7b2 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 7 Feb 2022 11:24:10 -0800 Subject: [PATCH 01/34] [SYCL] Add clang support for device_global --- clang/include/clang/Basic/Attr.td | 23 ++++++++++ clang/include/clang/Basic/AttrDocs.td | 41 +++++++++++++++++ .../clang/Basic/DiagnosticSemaKinds.td | 11 +++++ clang/include/clang/Sema/Sema.h | 15 ++++++ clang/lib/CodeGen/CodeGenModule.cpp | 19 ++++++++ clang/lib/CodeGen/CodeGenModule.h | 4 ++ clang/lib/Sema/SemaDecl.cpp | 6 +++ clang/lib/Sema/SemaDeclAttr.cpp | 35 ++++++++++++++ clang/lib/Sema/SemaExpr.cpp | 9 ++-- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 5 ++ clang/test/CodeGenSYCL/device_global.cpp | 46 +++++++++++++++++++ 11 files changed, 210 insertions(+), 4 deletions(-) create mode 100644 clang/test/CodeGenSYCL/device_global.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 38b7cbe254aa7..08622d320b3e6 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1432,6 +1432,29 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { let SupportsNonconformingLambdaSyntax = 1; } +def SYCLDetailDeviceGlobal: InheritableAttr { + let Spellings = [GNU<"device_global">, + CXX11<"__sycl_detail__", "device_global">]; + let Subjects = SubjectList<[CXXRecord], ErrorDiag>; + let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; + let Documentation = [Undocumented]; +} + +def SYCLDetailGlobalVariableAllowed : InheritableAttr { + let Spellings = [GNU<"global_variable_allowed">, + CXX11<"__sycl_detail__", "global_variable_allowed">]; + let Subjects = SubjectList<[CXXRecord, GlobalStorageNonLocalVar], ErrorDiag>; + let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; + let Documentation = [Undocumented]; +} + +def SYCLUniqueID : InheritableAttr { + let Spellings = [CXX11<"__sycl_detail__", "sycl-unique-id">]; + let Subjects = SubjectList<[GlobalVar]>; + let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; + let Documentation = [Undocumented]; +} + def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intel", "no_global_work_offset">]; let Args = [ExprArgument<"Value", /*optional*/1>]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index f4b6ba92eca29..d8b075a298e59 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3083,6 +3083,47 @@ function. In SYCL 2020 mode, the attribute is not propagated to the kernel. }]; } +def SYCLDetailDeviceGlobalAttrDocs : Documentation { + let Category = DocCatType; + let Heading = "__sycl_detail__::device_global"; + let Content = [{ +This attribute is part of support for SYCL device_global feature. +[[__sycl_detail__::device_global]] attribute is used for checking restrictions on variable declarations using the device_global type, rather than the class name. We do not intend to support this as a general attribute that customer code can use, so we have this attribute in sycl_detail namespace. + +.. code-block:: c++ + + template + struct [[__sycl_detail__::device_global]] device_global {} + + device_global Foo; + }]; +} + +def SYCLDetailGlobalVariableAllowedAttrDocs : Documentation { + let Category = DocCatType; + let Heading = "__sycl_detail__::global_variable_allowed"; + let Content = [{ +This attribute is part of support for SYCL device_global feature. +[[__sycl_detail__::global_variable_allowed]] attribute is used to avoid diagnosing an error when variables of type device_global are referenced in device codei. We do not intend to support this as a general attribute that customer code can use, therefore it is wrapped in sycl_detail namespace. + +.. code-block:: c++ + + template + struct [[__sycl_detail__::device_global]] device_global {} + + device_global Foo; + }]; +} + +def SYCLUniqueIDAttrDocs : Documentation { + let Category = DocCatVariable; + let Heading = "sycl-unique-id"; + let Content = [{ +This attribute is part of support for SYCL device_global feature. +sycl-unique-id is an LLVM IR attribute added to the definition of each device_global variable, which provides a unique string identifier for each device global variable. This attribute uses __builtin_sycl_unique_stable_id to generate the unique string. This cannont be used in customer code. + }]; +} + def SYCLFPGAPipeDocs : Documentation { let Category = DocCatStmt; let Heading = "pipe (read_only, write_only)"; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 91a64b1ad5558..8b15872d87cab 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7099,6 +7099,17 @@ def warn_format_nonliteral : Warning< "format string is not a string literal">, InGroup, DefaultIgnore; +def err_non_static_member_use_not_allowed : Error< + "use of non-static member variable %0 is not allowed">; +def err_not_publicly_accessible: Error< + "member variable %0 not publicly accessible">; +def err_array_of_device_global_not_allowed : Error< + "array of device_global %0 is not allowed">; +def err_shadow_variable_within_same_namespace: Error< + "shadow variable %0 not allowed withing the same enclosing namespace scope">; +def err_namespace_name_shadows_namespace_containing_device_global : Error< + "not allowed: namespace name shadows %0 namespace which contains device_global">; + def err_unexpected_interface : Error< "unexpected interface name %0: expected expression">; def err_ref_non_value : Error<"%0 does not refer to a value">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c391301e30cf9..14c1361cba34e 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10567,6 +10567,21 @@ class Sema final { const SYCLUsesAspectsAttr &A); void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); +/* + void AddSYCLDetailDeviceGlobal(Decl *D, const AttributeCommonInfo &CI, + Expr *E); + SYCLDetailDeviceGlobalAttr * + MergeSYCLDetailDeviceGlobalAttr(Decl *D, const SYCLDetailDeviceGlobalAttr &A); + + void AddSYCLDetailGlobalVariableAllowed(Decl *D, const AttributeCommonInfo &CI, + Expr *E); + SYCLDetailGlobalVariableAllowedAttr * + MergeSYCLDetailGlobalVariableAllowedAttr(Decl *D, const SYCLDetailGlobalVariableAllowedAttr &A); +*/ + void AddSYCLUniqueIDAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E); + SYCLUniqueIDAttr *MergeSYCLUniqueIDAttr(Decl *D, const SYCLUniqueIDAttr &A); + /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index a48c0be947dee..a6bebadb82146 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -31,6 +31,7 @@ #include "clang/AST/DeclCXX.h" #include "clang/AST/DeclObjC.h" #include "clang/AST/DeclTemplate.h" +#include "clang/AST/Expr.h" #include "clang/AST/Mangle.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" @@ -2216,6 +2217,8 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD, GV->addAttribute("rodata-section", SA->getName()); if (auto *SA = D->getAttr()) GV->addAttribute("relro-section", SA->getName()); +// if (auto *SA = D->getAttr()) +// GV->addAttribute("sycl-unique-id","pink"); } if (auto *F = dyn_cast(GO)) { @@ -2839,6 +2842,16 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation())); } +void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV, + const RecordDecl *RD) { + const auto *A = RD->getAttr(); + assert(A && "no device_global attribute"); + const VarDecl *VD = dyn_cast(RD->getParent()); + auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); + GV->addAttribute("sycl-unique-id", builtinString); + //GV->addAttribute("sycl-unique-id", "pink"); +} + bool CodeGenModule::isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn, SourceLocation Loc) const { const auto &NoSanitizeL = getContext().getNoSanitizeList(); @@ -4927,6 +4940,12 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, if (getLangOpts().SYCLIsDevice) addGlobalIntelFPGAAnnotation(D, GV); + if (getLangOpts().SYCLIsDevice) { + const RecordDecl *RD = D->getType()->getAsRecordDecl(); + if (RD && RD->hasAttr()) + addSYCLUniqueID(GV, RD); + } + if (D->getType().isRestrictQualified()) { llvm::LLVMContext &Context = getLLVMContext(); diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index 7b4d874badc3d..4153ca1e088da 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1315,6 +1315,10 @@ class CodeGenModule : public CodeGenTypeCache { /// annotations are emitted during finalization of the LLVM code. void AddGlobalAnnotations(const ValueDecl *D, llvm::GlobalValue *GV); + /// Add "sycl-unique-id" llvm attribute for global variables marked with + /// SYCL device_global attribute + void addSYCLUniqueID(llvm::GlobalVariable *GV, const RecordDecl *RD); + bool isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn, SourceLocation Loc) const; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 79552a10e59b5..06dfef4833ccc 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2781,6 +2781,12 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, NewAttr = S.MergeSYCLUsesAspectsAttr(D, *A); else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeSYCLIntelPipeIOAttr(D, *A); +// else if (const auto *A = dyn_cast(Attr)) +// NewAttr = S.MergeSYCLDetailDeviceGlobalAttr(D, *A); +// else if (const auto *A = dyn_cast(Attr)) +// NewAttr = S.MergeSYCLDetailGlobalVariableAllowedAttr(D, *A); +// else if (const auto *A = dyn_cast(Attr)) +// NewAttr = S.MergeSYCLUniqueIDAttr(D, *A); else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr)) NewAttr = cast(Attr->clone(S.Context)); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 995bea31a318b..248e26a148dcf 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4056,6 +4056,32 @@ Sema::MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A) { return ::new (Context) SYCLIntelLoopFuseAttr(Context, A, A.getValue()); } +static void handleSYCLDetailDeviceGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + //if (D->isATemplateDecl()) { + if (const auto *DeclAttr = D->getAttr()) { + auto *RD = dyn_cast(D); + if (isa(RD) && !S.isUnevaluatedContext()) + S.Diag(AL.getLoc(), diag::err_invalid_non_static_member_use) << AL; + } + + D->addAttr(::new (S.Context) SYCLDetailDeviceGlobalAttr(S.Context, AL)); +} + +static void handleSYCLDetailGlobalVariableAllowedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { +// Avoid diagnosing any erors here, simply accept the + if (const auto *DeclAttr = D->getAttr()) { + if (auto VD = dyn_cast(D)) { + // avoid diagnosing error + } + } + + D->addAttr(::new (S.Context) SYCLDetailGlobalVariableAllowedAttr(S.Context, AL)); +} + +/*static void handleSYCLUniqueIDAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E) { +}*/ + static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) { // If no attribute argument is specified, set to default value '1'. Expr *E = A.isArgExpr(0) @@ -10388,6 +10414,15 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim: handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLDetailDeviceGlobal: + handleSYCLDetailDeviceGlobalAttr(S, D, AL); + break; + case ParsedAttr::AT_SYCLDetailGlobalVariableAllowed: + handleSYCLDetailGlobalVariableAllowedAttr(S, D, AL); + break; +// case ParsedAttr::AT_SYCLUniqueID: +// handleSYCLUniqueIDAttr(S, D, AL); +// break; case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset: handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index f59bcc982cdbc..3daac4df0436c 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -227,16 +227,17 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, (!isUnevaluatedContext() && !isConstantEvaluated()); bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD); // Non-const statics are not allowed in SYCL except for ESIMD or with the - // SYCLGlobalVar attribute. + // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && - !VD->hasAttr()) + (!VD->hasAttr() || !VD->hasAttr())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the - // SYCLGlobalVar attribute. + // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage() && !VD->hasAttr()) + VD->hasGlobalStorage() && (!VD->hasAttr() + || !VD->hasAttr())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index a883c8a41cc16..24987a0d4fbd1 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -129,6 +129,11 @@ struct no_alias { template class instance {}; }; } // namespace property +// Global type decorated with attributes +template +struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global +//struct __attribute__((device_global)) device_global { // sycl::ext::oneapi:device_global +}; } // namespace oneapi } // namespace ext diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp new file mode 100644 index 0000000000000..4c5554d31514d --- /dev/null +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -std=c++17 -emit-llvm %s -o - | FileCheck %s +#include "Inputs/sycl.hpp" + +using namespace sycl::ext::oneapi; +static device_global Foo; + +device_global a; // OK +static device_global b; // OK +//inline device_global c; // OK + +struct Foo { + static device_global d; // OK +}; +device_global Foo::d; + +struct Bar { + device_global e; // ILLEGAL: non-static member variable not +}; // allowed + +//struct Baz { +// private: +// static device_global f; // ILLEGAL: not publicly accessible from +//}; // namespace scope +//device_global Baz::f; + +//device_global g; // OK +//device_global h[4]; // ILLEGAL: array of "device_global" not + // allowed + +//device_global same_name; // OK +//namespace foo { +// device_global same_name; // OK +//} +//namespace { +// device_global same_name; // OK +//} + +//inline namespace other { +// device_global same_name; // ILLEGAL: shadows "device_global" variable +//} // with same name in enclosing namespace scope + +//inline namespace { +// namespace foo { // ILLEGAL: namespace name shadows "::foo" +// } // namespace which contains "device_global" + // variable. +//} From 117de59a6008c7adc8dec72a182292db445db2a6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 16 Feb 2022 12:40:29 -0800 Subject: [PATCH 02/34] Remove unused lines i.e., Merge attribute method calls and SYCLUniqueID --- clang/include/clang/Basic/Attr.td | 7 ------- clang/include/clang/Basic/AttrDocs.td | 9 --------- clang/include/clang/Sema/Sema.h | 14 -------------- clang/lib/Sema/SemaDecl.cpp | 6 ------ clang/lib/Sema/SemaDeclAttr.cpp | 7 ------- 5 files changed, 43 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 08622d320b3e6..f841359da02ab 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1448,13 +1448,6 @@ def SYCLDetailGlobalVariableAllowed : InheritableAttr { let Documentation = [Undocumented]; } -def SYCLUniqueID : InheritableAttr { - let Spellings = [CXX11<"__sycl_detail__", "sycl-unique-id">]; - let Subjects = SubjectList<[GlobalVar]>; - let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Documentation = [Undocumented]; -} - def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intel", "no_global_work_offset">]; let Args = [ExprArgument<"Value", /*optional*/1>]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index d8b075a298e59..2d6fa203c529b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3115,15 +3115,6 @@ This attribute is part of support for SYCL device_global feature. }]; } -def SYCLUniqueIDAttrDocs : Documentation { - let Category = DocCatVariable; - let Heading = "sycl-unique-id"; - let Content = [{ -This attribute is part of support for SYCL device_global feature. -sycl-unique-id is an LLVM IR attribute added to the definition of each device_global variable, which provides a unique string identifier for each device global variable. This attribute uses __builtin_sycl_unique_stable_id to generate the unique string. This cannont be used in customer code. - }]; -} - def SYCLFPGAPipeDocs : Documentation { let Category = DocCatStmt; let Heading = "pipe (read_only, write_only)"; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 14c1361cba34e..691cca10e54c4 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10567,20 +10567,6 @@ class Sema final { const SYCLUsesAspectsAttr &A); void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); -/* - void AddSYCLDetailDeviceGlobal(Decl *D, const AttributeCommonInfo &CI, - Expr *E); - SYCLDetailDeviceGlobalAttr * - MergeSYCLDetailDeviceGlobalAttr(Decl *D, const SYCLDetailDeviceGlobalAttr &A); - - void AddSYCLDetailGlobalVariableAllowed(Decl *D, const AttributeCommonInfo &CI, - Expr *E); - SYCLDetailGlobalVariableAllowedAttr * - MergeSYCLDetailGlobalVariableAllowedAttr(Decl *D, const SYCLDetailGlobalVariableAllowedAttr &A); -*/ - void AddSYCLUniqueIDAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); - SYCLUniqueIDAttr *MergeSYCLUniqueIDAttr(Decl *D, const SYCLUniqueIDAttr &A); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 06dfef4833ccc..79552a10e59b5 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2781,12 +2781,6 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, NewAttr = S.MergeSYCLUsesAspectsAttr(D, *A); else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeSYCLIntelPipeIOAttr(D, *A); -// else if (const auto *A = dyn_cast(Attr)) -// NewAttr = S.MergeSYCLDetailDeviceGlobalAttr(D, *A); -// else if (const auto *A = dyn_cast(Attr)) -// NewAttr = S.MergeSYCLDetailGlobalVariableAllowedAttr(D, *A); -// else if (const auto *A = dyn_cast(Attr)) -// NewAttr = S.MergeSYCLUniqueIDAttr(D, *A); else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr)) NewAttr = cast(Attr->clone(S.Context)); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 248e26a148dcf..3c8b5b5f13786 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4078,10 +4078,6 @@ static void handleSYCLDetailGlobalVariableAllowedAttr(Sema &S, Decl *D, const Pa D->addAttr(::new (S.Context) SYCLDetailGlobalVariableAllowedAttr(S.Context, AL)); } -/*static void handleSYCLUniqueIDAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E) { -}*/ - static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) { // If no attribute argument is specified, set to default value '1'. Expr *E = A.isArgExpr(0) @@ -10420,9 +10416,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLDetailGlobalVariableAllowed: handleSYCLDetailGlobalVariableAllowedAttr(S, D, AL); break; -// case ParsedAttr::AT_SYCLUniqueID: -// handleSYCLUniqueIDAttr(S, D, AL); -// break; case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset: handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL); break; From a96d57af9e67beaff70c3260c4135d96111dd0e1 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 16 Feb 2022 13:07:38 -0800 Subject: [PATCH 03/34] Address some comments, fix format, remove unused lines --- clang/include/clang/Basic/AttrDocs.td | 10 ++++++++-- clang/lib/CodeGen/CodeGenModule.cpp | 4 ---- clang/lib/Sema/Sema.cpp | 3 ++- clang/lib/Sema/SemaExpr.cpp | 8 +++++--- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 2 +- 5 files changed, 16 insertions(+), 11 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2d6fa203c529b..d3ba9a1dbe6cb 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3088,7 +3088,10 @@ def SYCLDetailDeviceGlobalAttrDocs : Documentation { let Heading = "__sycl_detail__::device_global"; let Content = [{ This attribute is part of support for SYCL device_global feature. -[[__sycl_detail__::device_global]] attribute is used for checking restrictions on variable declarations using the device_global type, rather than the class name. We do not intend to support this as a general attribute that customer code can use, so we have this attribute in sycl_detail namespace. +[[__sycl_detail__::device_global]] attribute is used for checking restrictions +on variable declarations using the device_global type instead of the class name. +We do not intend to support this as a general attribute that customer code can +use, so we have this attribute in sycl_detail namespace. .. code-block:: c++ @@ -3104,7 +3107,10 @@ def SYCLDetailGlobalVariableAllowedAttrDocs : Documentation { let Heading = "__sycl_detail__::global_variable_allowed"; let Content = [{ This attribute is part of support for SYCL device_global feature. -[[__sycl_detail__::global_variable_allowed]] attribute is used to avoid diagnosing an error when variables of type device_global are referenced in device codei. We do not intend to support this as a general attribute that customer code can use, therefore it is wrapped in sycl_detail namespace. +[[__sycl_detail__::global_variable_allowed]] attribute is used to avoid +diagnosing an error when variables of type device_global are referenced in +device code. We do not intend to support this as a general attribute that +customer code can use, therefore it is wrapped in sycl_detail namespace. .. code-block:: c++ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index a6bebadb82146..7ff84f30598c8 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -31,7 +31,6 @@ #include "clang/AST/DeclCXX.h" #include "clang/AST/DeclObjC.h" #include "clang/AST/DeclTemplate.h" -#include "clang/AST/Expr.h" #include "clang/AST/Mangle.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" @@ -2217,8 +2216,6 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD, GV->addAttribute("rodata-section", SA->getName()); if (auto *SA = D->getAttr()) GV->addAttribute("relro-section", SA->getName()); -// if (auto *SA = D->getAttr()) -// GV->addAttribute("sycl-unique-id","pink"); } if (auto *F = dyn_cast(GO)) { @@ -2849,7 +2846,6 @@ void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD = dyn_cast(RD->getParent()); auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); GV->addAttribute("sycl-unique-id", builtinString); - //GV->addAttribute("sycl-unique-id", "pink"); } bool CodeGenModule::isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn, diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index b4b1fc2e6b1b0..7651b5ca40234 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1691,7 +1691,8 @@ class DeferredDiagnosticsEmitter void visitUsedDecl(SourceLocation Loc, Decl *D) { if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { - if (!S.checkAllowedSYCLInitializer(VD)) { + if (!S.checkAllowedSYCLInitializer(VD) && + !VD->hasAttr()) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; return; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 3daac4df0436c..003b2c9ad4b6f 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -230,14 +230,16 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && - (!VD->hasAttr() || !VD->hasAttr())) + (!VD->hasAttr() || + !VD->hasAttr())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage() && (!VD->hasAttr() - || !VD->hasAttr())) + VD->hasGlobalStorage() && + (!VD->hasAttr() || + !VD->hasAttr())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 24987a0d4fbd1..c62fe597d229e 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -132,7 +132,7 @@ struct no_alias { // Global type decorated with attributes template struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global -//struct __attribute__((device_global)) device_global { // sycl::ext::oneapi:device_global + device_global() = default; }; } // namespace oneapi } // namespace ext From f2230afcd9c997034aa362ace49d71f6fe7c9a4b Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 17 Feb 2022 08:09:06 -0800 Subject: [PATCH 04/34] Add Sema test; address more comments --- clang/include/clang/Basic/Attr.td | 6 +-- clang/include/clang/Sema/Sema.h | 1 - clang/test/CodeGenSYCL/Inputs/sycl.hpp | 4 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 6 +++ clang/test/SemaSYCL/device_global.cpp | 65 ++++++++++++++++++++++++++ 5 files changed, 76 insertions(+), 6 deletions(-) create mode 100644 clang/test/SemaSYCL/device_global.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index f841359da02ab..e3e04cdb744aa 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1437,15 +1437,15 @@ def SYCLDetailDeviceGlobal: InheritableAttr { CXX11<"__sycl_detail__", "device_global">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Documentation = [Undocumented]; + let Documentation = [SYCLDetailDeviceGlobalAttrDocs]; } def SYCLDetailGlobalVariableAllowed : InheritableAttr { let Spellings = [GNU<"global_variable_allowed">, CXX11<"__sycl_detail__", "global_variable_allowed">]; - let Subjects = SubjectList<[CXXRecord, GlobalStorageNonLocalVar], ErrorDiag>; + let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Documentation = [Undocumented]; + let Documentation = [SYCLDetailGlobalVariableAllowedAttrDocs]; } def SYCLIntelNoGlobalWorkOffset : InheritableAttr { diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 691cca10e54c4..c391301e30cf9 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10567,7 +10567,6 @@ class Sema final { const SYCLUsesAspectsAttr &A); void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); - /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index c62fe597d229e..9341b0d3e38a1 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -129,10 +129,10 @@ struct no_alias { template class instance {}; }; } // namespace property -// Global type decorated with attributes +// device_global type decorated with attributes template struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global - device_global() = default; + device_global() {} }; } // namespace oneapi } // namespace ext diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index f3e27d2fd928c..ff77aaa732ac8 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -64,6 +64,12 @@ namespace ext { namespace oneapi { template class accessor_property_list {}; + +// device_global type decorated with attributes +template +struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global + device_global() {} +}; } // namespace oneapi } // namespace ext diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp new file mode 100644 index 0000000000000..aa874789058bc --- /dev/null +++ b/clang/test/SemaSYCL/device_global.cpp @@ -0,0 +1,65 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump %s | FileCheck %s +#include "Inputs/sycl.hpp" + +using namespace sycl::ext::oneapi; + +device_global a; // OK +static device_global b; // OK +inline device_global c; // OK + +struct Foo { + static device_global d; // OK +}; +device_global Foo::d; + +//struct Bar { +// device_global e; // ILLEGAL: non-static member variable not +//}; // allowed + +//struct Baz { +// private: +// static device_global f; // ILLEGAL: not publicly accessible from +//}; // namespace scope +//device_global Baz::f; + +//device_global g; // OK +//device_global h[4]; // ILLEGAL: array of "device_global" not + // allowed + +//device_global same_name; // OK +//namespace foo { +// device_global same_name; // OK +//} +//namespace { +// device_global same_name; // OK +//} + +//inline namespace other { +// device_global same_name; // ILLEGAL: shadows "device_global" variable +//} // with same name in enclosing namespace scope + +//inline namespace { +// namespace foo { // ILLEGAL: namespace name shadows "::foo" +// } // namespace which contains "device_global" + // variable. +//} +// +// CHECK: ClassTemplateDecl {{.*}} device_global +// CHECK: CXXRecordDecl {{.*}} struct device_global definition +// CHECK: SYCLDetailDeviceGlobalAttr {{.*}} +// CHECK: SYCLDetailGlobalVariableAllowedAttr {{.*}} +// CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition +// CHECK: SYCLDetailDeviceGlobalAttr {{.*}} +// CHECK: SYCLDetailGlobalVariableAllowedAttr {{.*}} + +// CHECK: VarDecl {{.*}} a 'device_global':'sycl::ext::oneapi::device_global' callinit +// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' +// CHECK: VarDecl {{.*}} b 'device_global':'sycl::ext::oneapi::device_global' static callinit +// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' +// CHECK: VarDecl {{.*}} c 'device_global':'sycl::ext::oneapi::device_global' inline callinit +// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' +// CHECK: CXXRecordDecl {{.*}} struct Foo definition +// CHECK: VarDecl {{.*}} d 'device_global':'sycl::ext::oneapi::device_global' static +// CHECK: VarDecl {{.*}} d 'device_global':'sycl::ext::oneapi::device_global' callinit +// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' + From 0fb176dd80e5d968857112ab50241886a0243ba9 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Fri, 25 Feb 2022 16:19:52 +0300 Subject: [PATCH 05/34] Some fixes in order to pass CodeGen test - Check type of the variable, not the variable itself when checking presence of the attribute - Fix generation of LLVM IR attribute - Add the CodeGen test --- clang/lib/CodeGen/CodeGenModule.cpp | 7 +- clang/lib/CodeGen/CodeGenModule.h | 2 +- clang/lib/Sema/Sema.cpp | 13 +++- clang/lib/Sema/SemaExpr.cpp | 21 ++++-- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 9 ++- clang/test/CodeGenSYCL/device_global.cpp | 86 +++++++++++++----------- 6 files changed, 85 insertions(+), 53 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 7ff84f30598c8..8537dcee1c2a0 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2840,10 +2840,7 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, } void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV, - const RecordDecl *RD) { - const auto *A = RD->getAttr(); - assert(A && "no device_global attribute"); - const VarDecl *VD = dyn_cast(RD->getParent()); + const VarDecl *VD) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); GV->addAttribute("sycl-unique-id", builtinString); } @@ -4939,7 +4936,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, if (getLangOpts().SYCLIsDevice) { const RecordDecl *RD = D->getType()->getAsRecordDecl(); if (RD && RD->hasAttr()) - addSYCLUniqueID(GV, RD); + addSYCLUniqueID(GV, D); } if (D->getType().isRestrictQualified()) { diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index 4153ca1e088da..c86a0bedf0346 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1317,7 +1317,7 @@ class CodeGenModule : public CodeGenTypeCache { /// Add "sycl-unique-id" llvm attribute for global variables marked with /// SYCL device_global attribute - void addSYCLUniqueID(llvm::GlobalVariable *GV, const RecordDecl *RD); + void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD); bool isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn, SourceLocation Loc) const; diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 7651b5ca40234..c0185529481c6 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1620,6 +1620,17 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) { FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD); } } +static bool isSyclGlobalVariableAllowedType(QualType Ty) { + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + if (!RecTy) + return false; + if (auto *CTSD = dyn_cast(RecTy)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RD = Template->getTemplatedDecl()) + return RD->hasAttr(); + } + return RecTy->hasAttr(); +} namespace { @@ -1692,7 +1703,7 @@ class DeferredDiagnosticsEmitter if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { if (!S.checkAllowedSYCLInitializer(VD) && - !VD->hasAttr()) { + !isSyclGlobalVariableAllowedType(VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; return; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 003b2c9ad4b6f..a0b3573645c0f 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -202,6 +202,18 @@ void Sema::MaybeSuggestAddingStaticToDecl(const FunctionDecl *Cur) { } } +static bool isSyclGlobalVariableAllowedType(QualType Ty) { + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + if (!RecTy) + return false; + if (auto *CTSD = dyn_cast(RecTy)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RD = Template->getTemplatedDecl()) + return RD->hasAttr(); + } + return RecTy->hasAttr(); +} + /// Determine whether the use of this declaration is valid, and /// emit any corresponding diagnostics. /// @@ -230,16 +242,15 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && - (!VD->hasAttr() || - !VD->hasAttr())) + !VD->hasAttr() && + !isSyclGlobalVariableAllowedType(VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage() && - (!VD->hasAttr() || - !VD->hasAttr())) + VD->hasGlobalStorage() && !VD->hasAttr() && + !isSyclGlobalVariableAllowedType(VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 9341b0d3e38a1..b557927dab460 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -129,10 +129,17 @@ struct no_alias { template class instance {}; }; } // namespace property + // device_global type decorated with attributes template -struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global +class [[__sycl_detail__::device_global]] +[[__sycl_detail__::global_variable_allowed]] device_global { +public : + const T & get() const noexcept { return *Data; } device_global() {} + operator T&() noexcept { return *Data; } +private: + T *Data; }; } // namespace oneapi } // namespace ext diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 4c5554d31514d..5596bed20d187 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -1,46 +1,52 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -std=c++17 -emit-llvm %s -o - | FileCheck %s -#include "Inputs/sycl.hpp" +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s +#include "sycl.hpp" using namespace sycl::ext::oneapi; -static device_global Foo; +using namespace cl::sycl; +queue q; -device_global a; // OK -static device_global b; // OK -//inline device_global c; // OK +device_global A; +static device_global B; struct Foo { - static device_global d; // OK + static device_global C; }; -device_global Foo::d; - -struct Bar { - device_global e; // ILLEGAL: non-static member variable not -}; // allowed - -//struct Baz { -// private: -// static device_global f; // ILLEGAL: not publicly accessible from -//}; // namespace scope -//device_global Baz::f; - -//device_global g; // OK -//device_global h[4]; // ILLEGAL: array of "device_global" not - // allowed - -//device_global same_name; // OK -//namespace foo { -// device_global same_name; // OK -//} -//namespace { -// device_global same_name; // OK -//} - -//inline namespace other { -// device_global same_name; // ILLEGAL: shadows "device_global" variable -//} // with same name in enclosing namespace scope - -//inline namespace { -// namespace foo { // ILLEGAL: namespace name shadows "::foo" -// } // namespace which contains "device_global" - // variable. -//} +device_global Foo::C; +// CHECK: @A = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]] +// CHECK: @_ZL1B = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]] +// CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]] + +device_global same_name; +namespace NS { + device_global same_name; +} +// CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]] +// CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]] + +void foo() { + q.submit([&](handler &h) { + h.single_task([=]() {(void)A; (void)B; (void)Foo::C; (void)same_name; (void)NS::same_name; }); + }); +} + +namespace { + device_global same_name; +} +// CHECK: @_ZN12_GLOBAL__N_19same_nameE = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ANON_NS_ATTRS:[0-9]+]] + +namespace { +void bar() { + q.submit([&](handler &h) { + h.single_task([=]() { int A = same_name; }); + }); +} + +} + + +// CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } +// CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" } +// CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" } +// CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" } +// CHECK: attributes #[[SAME_NAME_NS_ATTRS]] = { "sycl-unique-id"="_ZN2NS9same_nameE" } +// CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" } From e4c15c41888c994588ef6676ce266d6eb39ae773 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 25 Feb 2022 14:15:49 -0800 Subject: [PATCH 06/34] Refactor isSyclGlobalVariableAllowedType --- clang/include/clang/Sema/Sema.h | 2 ++ clang/lib/Sema/Sema.cpp | 5 +++-- clang/lib/Sema/SemaExpr.cpp | 12 ------------ 3 files changed, 5 insertions(+), 14 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c391301e30cf9..7b7119573112c 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13125,6 +13125,8 @@ class Sema final { SourceLocation BuiltinLoc, SourceLocation RParenLoc); + bool isSyclGlobalVariableAllowedType(QualType Ty); + private: bool SemaBuiltinPrefetch(CallExpr *TheCall); bool SemaBuiltinAllocaWithAlign(CallExpr *TheCall); diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index c0185529481c6..f1a289d60a371 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1620,7 +1620,8 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) { FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD); } } -static bool isSyclGlobalVariableAllowedType(QualType Ty) { + +bool Sema::isSyclGlobalVariableAllowedType(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; @@ -1703,7 +1704,7 @@ class DeferredDiagnosticsEmitter if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { if (!S.checkAllowedSYCLInitializer(VD) && - !isSyclGlobalVariableAllowedType(VD->getType())) { + !S.isSyclGlobalVariableAllowedType(VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; return; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index a0b3573645c0f..ab322f41befbe 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -202,18 +202,6 @@ void Sema::MaybeSuggestAddingStaticToDecl(const FunctionDecl *Cur) { } } -static bool isSyclGlobalVariableAllowedType(QualType Ty) { - const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); - if (!RecTy) - return false; - if (auto *CTSD = dyn_cast(RecTy)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); - } - return RecTy->hasAttr(); -} - /// Determine whether the use of this declaration is valid, and /// emit any corresponding diagnostics. /// From af8a294777491a7e99a9ccfa2922f8d489dd8d42 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 18 Feb 2022 15:47:48 -0800 Subject: [PATCH 07/34] Fix some test cases; address some comments --- clang/include/clang/Basic/Attr.td | 10 ++-- clang/include/clang/Basic/AttrDocs.td | 4 +- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/lib/Sema/Sema.cpp | 4 +- clang/lib/Sema/SemaDecl.cpp | 3 +- clang/lib/Sema/SemaDeclAttr.cpp | 32 ++++++------- clang/lib/Sema/SemaExpr.cpp | 4 +- clang/test/CodeGenSYCL/device_global.cpp | 8 +++- clang/test/SemaSYCL/Inputs/sycl.hpp | 8 +++- clang/test/SemaSYCL/device_global.cpp | 60 ++++++++++++++++-------- 10 files changed, 85 insertions(+), 50 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index e3e04cdb744aa..507561d36d79a 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1432,20 +1432,22 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { let SupportsNonconformingLambdaSyntax = 1; } -def SYCLDetailDeviceGlobal: InheritableAttr { +def SYCLDeviceGlobal: InheritableAttr { let Spellings = [GNU<"device_global">, CXX11<"__sycl_detail__", "device_global">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Documentation = [SYCLDetailDeviceGlobalAttrDocs]; + let Documentation = [SYCLDeviceGlobalAttrDocs]; + let SimpleHandler = 1; } -def SYCLDetailGlobalVariableAllowed : InheritableAttr { +def SYCLGlobalVariableAllowed : InheritableAttr { let Spellings = [GNU<"global_variable_allowed">, CXX11<"__sycl_detail__", "global_variable_allowed">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Documentation = [SYCLDetailGlobalVariableAllowedAttrDocs]; + let Documentation = [SYCLGlobalVariableAllowedAttrDocs]; + let SimpleHandler = 1; } def SYCLIntelNoGlobalWorkOffset : InheritableAttr { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index d3ba9a1dbe6cb..39c334c92075a 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3083,7 +3083,7 @@ function. In SYCL 2020 mode, the attribute is not propagated to the kernel. }]; } -def SYCLDetailDeviceGlobalAttrDocs : Documentation { +def SYCLDeviceGlobalAttrDocs : Documentation { let Category = DocCatType; let Heading = "__sycl_detail__::device_global"; let Content = [{ @@ -3102,7 +3102,7 @@ use, so we have this attribute in sycl_detail namespace. }]; } -def SYCLDetailGlobalVariableAllowedAttrDocs : Documentation { +def SYCLGlobalVariableAllowedAttrDocs : Documentation { let Category = DocCatType; let Heading = "__sycl_detail__::global_variable_allowed"; let Content = [{ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 8537dcee1c2a0..bd4775b4f1f76 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4935,7 +4935,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, if (getLangOpts().SYCLIsDevice) { const RecordDecl *RD = D->getType()->getAsRecordDecl(); - if (RD && RD->hasAttr()) + if (RD && RD->hasAttr()) addSYCLUniqueID(GV, D); } diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index f1a289d60a371..1d998399210fe 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1628,9 +1628,9 @@ bool Sema::isSyclGlobalVariableAllowedType(QualType Ty) { if (auto *CTSD = dyn_cast(RecTy)) { ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); + return RD->hasAttr(); } - return RecTy->hasAttr(); + return RecTy->hasAttr(); } namespace { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 79552a10e59b5..91f4fb891fefd 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7395,7 +7395,8 @@ NamedDecl *Sema::ActOnVariableDeclarator( // Static variables declared inside SYCL device code must be const or // constexpr if (getLangOpts().SYCLIsDevice) - if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context)) + if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && + NewVD->getType()->getAsRecordDecl()->hasAttr()) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 3c8b5b5f13786..50e67592db1b7 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4056,26 +4056,26 @@ Sema::MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A) { return ::new (Context) SYCLIntelLoopFuseAttr(Context, A, A.getValue()); } -static void handleSYCLDetailDeviceGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - //if (D->isATemplateDecl()) { - if (const auto *DeclAttr = D->getAttr()) { - auto *RD = dyn_cast(D); - if (isa(RD) && !S.isUnevaluatedContext()) - S.Diag(AL.getLoc(), diag::err_invalid_non_static_member_use) << AL; - } - - D->addAttr(::new (S.Context) SYCLDetailDeviceGlobalAttr(S.Context, AL)); +static void handleSYCLDeviceGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + auto *RD = dyn_cast(D); + if (auto *CTSD = dyn_cast(D)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RDec = Template->getTemplatedDecl()) + if (isa(RDec) && !S.isUnevaluatedContext()) + S.Diag(AL.getLoc(), diag::err_invalid_non_static_member_use) << AL; + } + D->addAttr(::new (S.Context) SYCLDeviceGlobalAttr(S.Context, AL)); } -static void handleSYCLDetailGlobalVariableAllowedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { +static void handleSYCLGlobalVariableAllowedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Avoid diagnosing any erors here, simply accept the - if (const auto *DeclAttr = D->getAttr()) { + if (const auto *DeclAttr = D->getAttr()) { if (auto VD = dyn_cast(D)) { // avoid diagnosing error } } - D->addAttr(::new (S.Context) SYCLDetailGlobalVariableAllowedAttr(S.Context, AL)); + D->addAttr(::new (S.Context) SYCLGlobalVariableAllowedAttr(S.Context, AL)); } static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) { @@ -10410,11 +10410,11 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim: handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL); break; - case ParsedAttr::AT_SYCLDetailDeviceGlobal: - handleSYCLDetailDeviceGlobalAttr(S, D, AL); + case ParsedAttr::AT_SYCLDeviceGlobal: + handleSYCLDeviceGlobalAttr(S, D, AL); break; - case ParsedAttr::AT_SYCLDetailGlobalVariableAllowed: - handleSYCLDetailGlobalVariableAllowedAttr(S, D, AL); + case ParsedAttr::AT_SYCLGlobalVariableAllowed: + handleSYCLGlobalVariableAllowedAttr(S, D, AL); break; case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset: handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL); diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index ab322f41befbe..656b38036f6f7 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -227,7 +227,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, (!isUnevaluatedContext() && !isConstantEvaluated()); bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD); // Non-const statics are not allowed in SYCL except for ESIMD or with the - // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. + // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && !VD->hasAttr() && @@ -235,7 +235,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the - // SYCLGlobalVar or SYCLDetailGlobalVariableAllowed attribute. + // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr() && !isSyclGlobalVariableAllowedType(VD->getType())) diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 5596bed20d187..8cbbf801c3910 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -25,7 +25,13 @@ namespace NS { void foo() { q.submit([&](handler &h) { - h.single_task([=]() {(void)A; (void)B; (void)Foo::C; (void)same_name; (void)NS::same_name; }); + h.single_task([=]() { + (void)A; + (void)B; + (void)Foo::C; + (void)same_name; + (void)NS::same_name; + }); }); } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index ff77aaa732ac8..b4a4123673d19 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -67,8 +67,14 @@ class accessor_property_list {}; // device_global type decorated with attributes template -struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { // sycl::ext::oneapi:device_global +struct [[__sycl_detail__::device_global]] +[[__sycl_detail__::global_variable_allowed]] device_global { +public: + const T & get() const noexcept { return *Data; } device_global() {} + operator T&() noexcept { return *Data; } +private: + T *Data; }; } // namespace oneapi } // namespace ext diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index aa874789058bc..11f743bc53be1 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -3,9 +3,10 @@ using namespace sycl::ext::oneapi; -device_global a; // OK -static device_global b; // OK -inline device_global c; // OK +device_global glob; // OK +static device_global static_glob; // OK +inline device_global inline_glob; // OK +static const device_global static_const_glob; struct Foo { static device_global d; // OK @@ -22,17 +23,17 @@ device_global Foo::d; //}; // namespace scope //device_global Baz::f; -//device_global g; // OK +device_global not_array; // OK //device_global h[4]; // ILLEGAL: array of "device_global" not // allowed -//device_global same_name; // OK -//namespace foo { -// device_global same_name; // OK -//} -//namespace { -// device_global same_name; // OK -//} +device_global same_name; // OK +namespace foo { + device_global same_name; // OK +} +namespace { + device_global same_name; // OK +} //inline namespace other { // device_global same_name; // ILLEGAL: shadows "device_global" variable @@ -43,23 +44,42 @@ device_global Foo::d; // } // namespace which contains "device_global" // variable. //} + +int main() { + cl::sycl::kernel_single_task([=]() { + (void)glob; + (void)static_glob; + (void)inline_glob; + (void)static_const_glob; + (void)Foo::d; + }); + + cl::sycl::kernel_single_task([]() { +// static device_global non_const_static; + }); +} // // CHECK: ClassTemplateDecl {{.*}} device_global // CHECK: CXXRecordDecl {{.*}} struct device_global definition -// CHECK: SYCLDetailDeviceGlobalAttr {{.*}} -// CHECK: SYCLDetailGlobalVariableAllowedAttr {{.*}} +// CHECK: SYCLDeviceGlobalAttr {{.*}} +// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} // CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition -// CHECK: SYCLDetailDeviceGlobalAttr {{.*}} -// CHECK: SYCLDetailGlobalVariableAllowedAttr {{.*}} +// CHECK: SYCLDeviceGlobalAttr {{.*}} +// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} -// CHECK: VarDecl {{.*}} a 'device_global':'sycl::ext::oneapi::device_global' callinit +// CHECK: VarDecl {{.*}} used glob 'device_global':'sycl::ext::oneapi::device_global' callinit // CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} b 'device_global':'sycl::ext::oneapi::device_global' static callinit +// CHECK: VarDecl {{.*}} used static_glob 'device_global':'sycl::ext::oneapi::device_global' static callinit // CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} c 'device_global':'sycl::ext::oneapi::device_global' inline callinit +// CHECK: VarDecl {{.*}} used inline_glob 'device_global':'sycl::ext::oneapi::device_global' inline callinit // CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' +// CHECK: VarDecl {{.*}} used static_const_glob 'const device_global':'const sycl::ext::oneapi::device_global' static callinit +// CHECK: CXXConstructExpr {{.*}} 'const device_global':'const sycl::ext::oneapi::device_global' 'void ()' // CHECK: CXXRecordDecl {{.*}} struct Foo definition -// CHECK: VarDecl {{.*}} d 'device_global':'sycl::ext::oneapi::device_global' static +// CHECK: VarDecl {{.*}} used d 'device_global':'sycl::ext::oneapi::device_global' static // CHECK: VarDecl {{.*}} d 'device_global':'sycl::ext::oneapi::device_global' callinit // CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' - +// CHECK: VarDecl {{.*}} not_array 'device_global':'sycl::ext::oneapi::device_global' callinit +// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' +// CHECK: VarDecl {{.*}} same_name 'device_global':'sycl::ext::oneapi::device_global' callinit +// From feb841be2942f4849f2289fdc17ffa81770cfd98 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 28 Feb 2022 21:02:20 -0800 Subject: [PATCH 08/34] Remove explicit attribute handling; change diagnostic message; add test --- .../clang/Basic/DiagnosticSemaKinds.td | 6 ++-- clang/lib/Sema/SemaDecl.cpp | 21 ++++++++++++-- clang/lib/Sema/SemaDeclAttr.cpp | 28 ------------------- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 11 ++++++++ clang/test/CodeGenSYCL/device_global.cpp | 6 ++++ 5 files changed, 38 insertions(+), 34 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 8b15872d87cab..f98230cbf3af0 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7099,10 +7099,10 @@ def warn_format_nonliteral : Warning< "format string is not a string literal">, InGroup, DefaultIgnore; -def err_non_static_member_use_not_allowed : Error< - "use of non-static member variable %0 is not allowed">; +def err_sycl_device_global_incorrect_scope : Error< + "`device_global` variables must be static or declared at namespace scope">; def err_not_publicly_accessible: Error< - "member variable %0 not publicly accessible">; + "member variable %0 not publicly accessible from namespace scope">; def err_array_of_device_global_not_allowed : Error< "array of device_global %0 is not allowed">; def err_shadow_variable_within_same_namespace: Error< diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 91f4fb891fefd..bcbe62991ad3c 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7057,6 +7057,17 @@ static void copyAttrFromTypedefToDecl(Sema &S, Decl *D, const TypedefType *TT) { D->addAttr(Clone); } } +static bool isSyclDeviceGlobalType(QualType Ty) { + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + if (!RecTy) + return false; + if (auto *CTSD = dyn_cast(RecTy)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RD = Template->getTemplatedDecl()) + return RD->hasAttr(); + } + return RecTy->hasAttr(); +} NamedDecl *Sema::ActOnVariableDeclarator( Scope *S, Declarator &D, DeclContext *DC, TypeSourceInfo *TInfo, @@ -7394,11 +7405,15 @@ NamedDecl *Sema::ActOnVariableDeclarator( // Static variables declared inside SYCL device code must be const or // constexpr - if (getLangOpts().SYCLIsDevice) - if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && - NewVD->getType()->getAsRecordDecl()->hasAttr()) + if (getLangOpts().SYCLIsDevice) { + if (isSyclDeviceGlobalType(NewVD->getType()) && + SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage()) { + Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); + } + if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context)) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; + } switch (D.getDeclSpec().getConstexprSpecifier()) { case ConstexprSpecKind::Unspecified: diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 50e67592db1b7..995bea31a318b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4056,28 +4056,6 @@ Sema::MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A) { return ::new (Context) SYCLIntelLoopFuseAttr(Context, A, A.getValue()); } -static void handleSYCLDeviceGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - auto *RD = dyn_cast(D); - if (auto *CTSD = dyn_cast(D)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RDec = Template->getTemplatedDecl()) - if (isa(RDec) && !S.isUnevaluatedContext()) - S.Diag(AL.getLoc(), diag::err_invalid_non_static_member_use) << AL; - } - D->addAttr(::new (S.Context) SYCLDeviceGlobalAttr(S.Context, AL)); -} - -static void handleSYCLGlobalVariableAllowedAttr(Sema &S, Decl *D, const ParsedAttr &AL) { -// Avoid diagnosing any erors here, simply accept the - if (const auto *DeclAttr = D->getAttr()) { - if (auto VD = dyn_cast(D)) { - // avoid diagnosing error - } - } - - D->addAttr(::new (S.Context) SYCLGlobalVariableAllowedAttr(S.Context, AL)); -} - static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) { // If no attribute argument is specified, set to default value '1'. Expr *E = A.isArgExpr(0) @@ -10410,12 +10388,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim: handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL); break; - case ParsedAttr::AT_SYCLDeviceGlobal: - handleSYCLDeviceGlobalAttr(S, D, AL); - break; - case ParsedAttr::AT_SYCLGlobalVariableAllowed: - handleSYCLGlobalVariableAllowedAttr(S, D, AL); - break; case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset: handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL); break; diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index b557927dab460..1d71073ba41c0 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -141,6 +141,17 @@ public : private: T *Data; }; + +// decorated with only global_variable_allowed attribute +template +class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { +public : + const T & get() const noexcept { return *Data; } + only_global_var_allowed() {} + operator T&() noexcept { return *Data; } +private: + T *Data; +}; } // namespace oneapi } // namespace ext diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 8cbbf801c3910..5e89ea120abf5 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -23,6 +23,11 @@ namespace NS { // CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]] // CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]] +// check that we don't generate `sycl-unique-id` IR attribute if class does not use +// [[__sycl_detail__::device_global]] +only_global_var_allowed no_device_global; +// CHECK: @no_device_global = addrspace(1) global %"class.cl::sycl::ext::oneapi::only_global_var_allowed" zeroinitializer, align 8{{$}} + void foo() { q.submit([&](handler &h) { h.single_task([=]() { @@ -31,6 +36,7 @@ void foo() { (void)Foo::C; (void)same_name; (void)NS::same_name; + (void)no_device_global; }); }); } From df337a7de8f9164bc63c7d82ed6e0f6120920aa5 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 28 Feb 2022 23:31:29 -0800 Subject: [PATCH 09/34] Add diagnostic to test --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaDecl.cpp | 5 +++-- clang/test/SemaSYCL/device_global.cpp | 29 ++++++++++++++++----------- 3 files changed, 21 insertions(+), 14 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 7b7119573112c..471c69a254c36 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13126,6 +13126,7 @@ class Sema final { SourceLocation RParenLoc); bool isSyclGlobalVariableAllowedType(QualType Ty); + bool isSyclDeviceGlobalType(QualType Ty); private: bool SemaBuiltinPrefetch(CallExpr *TheCall); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index bcbe62991ad3c..f8256703bd7e8 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7057,7 +7057,7 @@ static void copyAttrFromTypedefToDecl(Sema &S, Decl *D, const TypedefType *TT) { D->addAttr(Clone); } } -static bool isSyclDeviceGlobalType(QualType Ty) { +bool Sema::isSyclDeviceGlobalType(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; @@ -7410,7 +7410,8 @@ NamedDecl *Sema::ActOnVariableDeclarator( SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage()) { Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); } - if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context)) + if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && + !isSyclGlobalVariableAllowedType(NewVD->getType())) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; } diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 11f743bc53be1..81280144cdb01 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump -verify %s | FileCheck %s #include "Inputs/sycl.hpp" using namespace sycl::ext::oneapi; @@ -13,18 +13,18 @@ struct Foo { }; device_global Foo::d; -//struct Bar { -// device_global e; // ILLEGAL: non-static member variable not -//}; // allowed +struct Bar { + device_global e; // ILLEGAL: non-static member variable not +}; // allowed -//struct Baz { -// private: -// static device_global f; // ILLEGAL: not publicly accessible from -//}; // namespace scope -//device_global Baz::f; +struct Baz { + private: + static device_global f; // ILLEGAL: not publicly accessible from +}; // namespace scope +device_global Baz::f; device_global not_array; // OK -//device_global h[4]; // ILLEGAL: array of "device_global" not +device_global h[4]; // ILLEGAL: array of "device_global" not // allowed device_global same_name; // OK @@ -54,8 +54,13 @@ int main() { (void)Foo::d; }); - cl::sycl::kernel_single_task([]() { -// static device_global non_const_static; + cl::sycl::kernel_single_task([]() { + // expected-error@+1{{`device_global` variables must be static or declared at namespace scope}} + device_global non_static; + + // expect no error on non_const_static declaration if decorated with + // [[__sycl_detail__::global_variable_allowed]] + static device_global non_const_static; }); } // From 748e8cefbd63f1697620ded0634079f49b89d959 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 1 Mar 2022 08:36:48 -0800 Subject: [PATCH 10/34] Address Mariya's comments --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 11 ----------- clang/test/CodeGenSYCL/device_global.cpp | 19 +++++++++++++++++++ clang/test/SemaSYCL/device_global.cpp | 16 ---------------- 3 files changed, 19 insertions(+), 27 deletions(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 1d71073ba41c0..b557927dab460 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -141,17 +141,6 @@ public : private: T *Data; }; - -// decorated with only global_variable_allowed attribute -template -class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { -public : - const T & get() const noexcept { return *Data; } - only_global_var_allowed() {} - operator T&() noexcept { return *Data; } -private: - T *Data; -}; } // namespace oneapi } // namespace ext diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 5e89ea120abf5..23ee803190851 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -1,6 +1,25 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s #include "sycl.hpp" +namespace cl { +namespace sycl { +namespace ext { +namespace oneapi { +// decorated with only global_variable_allowed attribute +template +class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { +public : + const T & get() const noexcept { return *Data; } + only_global_var_allowed() {} + operator T&() noexcept { return *Data; } +private: + T *Data; +}; +} // namespace oneapi +} // namespace ext +} // namespace sycl +} // namespace cl + using namespace sycl::ext::oneapi; using namespace cl::sycl; queue q; diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 81280144cdb01..7c735255f8218 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -72,19 +72,3 @@ int main() { // CHECK: SYCLDeviceGlobalAttr {{.*}} // CHECK: SYCLGlobalVariableAllowedAttr {{.*}} -// CHECK: VarDecl {{.*}} used glob 'device_global':'sycl::ext::oneapi::device_global' callinit -// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} used static_glob 'device_global':'sycl::ext::oneapi::device_global' static callinit -// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} used inline_glob 'device_global':'sycl::ext::oneapi::device_global' inline callinit -// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} used static_const_glob 'const device_global':'const sycl::ext::oneapi::device_global' static callinit -// CHECK: CXXConstructExpr {{.*}} 'const device_global':'const sycl::ext::oneapi::device_global' 'void ()' -// CHECK: CXXRecordDecl {{.*}} struct Foo definition -// CHECK: VarDecl {{.*}} used d 'device_global':'sycl::ext::oneapi::device_global' static -// CHECK: VarDecl {{.*}} d 'device_global':'sycl::ext::oneapi::device_global' callinit -// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} not_array 'device_global':'sycl::ext::oneapi::device_global' callinit -// CHECK: CXXConstructExpr {{.*}} 'device_global':'sycl::ext::oneapi::device_global' 'void ()' -// CHECK: VarDecl {{.*}} same_name 'device_global':'sycl::ext::oneapi::device_global' callinit -// From d84428cbeea2d1d58f234e9798e9387ae6e1d9b6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 1 Mar 2022 11:18:21 -0800 Subject: [PATCH 11/34] Fix typo; Refactor methods; --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/include/clang/Sema/Sema.h | 4 ++-- clang/lib/Sema/Sema.cpp | 9 +++++---- clang/lib/Sema/SemaDecl.cpp | 15 ++------------- clang/lib/Sema/SemaExpr.cpp | 4 ++-- 5 files changed, 12 insertions(+), 22 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f98230cbf3af0..635c1ca10686e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7106,7 +7106,7 @@ def err_not_publicly_accessible: Error< def err_array_of_device_global_not_allowed : Error< "array of device_global %0 is not allowed">; def err_shadow_variable_within_same_namespace: Error< - "shadow variable %0 not allowed withing the same enclosing namespace scope">; + "shadow variable %0 not allowed within the same enclosing namespace scope">; def err_namespace_name_shadows_namespace_containing_device_global : Error< "not allowed: namespace name shadows %0 namespace which contains device_global">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 471c69a254c36..acdfd091d0e85 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13125,8 +13125,8 @@ class Sema final { SourceLocation BuiltinLoc, SourceLocation RParenLoc); - bool isSyclGlobalVariableAllowedType(QualType Ty); - bool isSyclDeviceGlobalType(QualType Ty); + template + bool isSyclGlobalType(QualType Ty); private: bool SemaBuiltinPrefetch(CallExpr *TheCall); diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 1d998399210fe..bdcadb1a4ee95 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1621,16 +1621,17 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) { } } -bool Sema::isSyclGlobalVariableAllowedType(QualType Ty) { +template +bool Sema::isSyclGlobalType(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; if (auto *CTSD = dyn_cast(RecTy)) { ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); + return RD->hasAttr(); } - return RecTy->hasAttr(); + return RecTy->hasAttr(); } namespace { @@ -1704,7 +1705,7 @@ class DeferredDiagnosticsEmitter if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { if (!S.checkAllowedSYCLInitializer(VD) && - !S.isSyclGlobalVariableAllowedType(VD->getType())) { + !S.isSyclGlobalType(VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; return; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index f8256703bd7e8..31a829aea6bcd 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7057,17 +7057,6 @@ static void copyAttrFromTypedefToDecl(Sema &S, Decl *D, const TypedefType *TT) { D->addAttr(Clone); } } -bool Sema::isSyclDeviceGlobalType(QualType Ty) { - const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); - if (!RecTy) - return false; - if (auto *CTSD = dyn_cast(RecTy)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); - } - return RecTy->hasAttr(); -} NamedDecl *Sema::ActOnVariableDeclarator( Scope *S, Declarator &D, DeclContext *DC, TypeSourceInfo *TInfo, @@ -7406,12 +7395,12 @@ NamedDecl *Sema::ActOnVariableDeclarator( // Static variables declared inside SYCL device code must be const or // constexpr if (getLangOpts().SYCLIsDevice) { - if (isSyclDeviceGlobalType(NewVD->getType()) && + if (isSyclGlobalType(NewVD->getType()) && SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage()) { Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); } if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && - !isSyclGlobalVariableAllowedType(NewVD->getType())) + !isSyclGlobalType(NewVD->getType())) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; } diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 656b38036f6f7..0429a8bbaa43b 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -231,14 +231,14 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && !VD->hasAttr() && - !isSyclGlobalVariableAllowedType(VD->getType())) + !isSyclGlobalType(VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr() && - !isSyclGlobalVariableAllowedType(VD->getType())) + !isSyclGlobalType(VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. From 034315826f5f717013bf2c0781b660b5493a3362 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 2 Mar 2022 16:44:27 +0300 Subject: [PATCH 12/34] Fix quotes, update diag messages, report private members --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 6 +++--- clang/lib/Sema/SemaDeclCXX.cpp | 10 ++++++++++ clang/test/SemaSYCL/device_global.cpp | 3 ++- 3 files changed, 15 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 635c1ca10686e..54f799d69e065 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7100,9 +7100,9 @@ def warn_format_nonliteral : Warning< InGroup, DefaultIgnore; def err_sycl_device_global_incorrect_scope : Error< - "`device_global` variables must be static or declared at namespace scope">; -def err_not_publicly_accessible: Error< - "member variable %0 not publicly accessible from namespace scope">; + "'device_global' variables must be static or declared at namespace scope">; +def err_sycl_device_global_not_publicly_accessible: Error< + "'device_global' member variable %0 is not publicly accessible from namespace scope">; def err_array_of_device_global_not_allowed : Error< "array of device_global %0 is not allowed">; def err_shadow_variable_within_same_namespace: Error< diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 16cdb7e577237..9acdafdfff8a6 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3529,6 +3529,16 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, } } + if (getLangOpts().SYCLIsDevice) { + if (auto Value = dyn_cast(Member)) { + if (isSyclDeviceGlobalType(Value->getType()) && + Value->getAccess() != AS_public) { + Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) + << Value; + } + } + } + return Member; } diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 7c735255f8218..ef03dfef3252a 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -19,6 +19,7 @@ struct Bar { struct Baz { private: + // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}} static device_global f; // ILLEGAL: not publicly accessible from }; // namespace scope device_global Baz::f; @@ -55,7 +56,7 @@ int main() { }); cl::sycl::kernel_single_task([]() { - // expected-error@+1{{`device_global` variables must be static or declared at namespace scope}} + // expected-error@+1{{'device_global' variables must be static or declared at namespace scope}} device_global non_static; // expect no error on non_const_static declaration if decorated with From 6d6d7dd01b2753d2fcad29be62ae7068ad7012d3 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 09:57:49 -0800 Subject: [PATCH 13/34] Change DeviceGlobalType checking call --- clang/lib/Sema/SemaDeclCXX.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 9acdafdfff8a6..e352fa026aea3 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3531,7 +3531,7 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { - if (isSyclDeviceGlobalType(Value->getType()) && + if (isSyclGlobalType(Value->getType()) && Value->getAccess() != AS_public) { Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) << Value; From 679b5f02a36576b026770098fc602cb234c6964d Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 10:52:37 -0800 Subject: [PATCH 14/34] Move isSyclGlobalType definition to header --- clang/include/clang/Sema/Sema.h | 14 ++++++++++++-- clang/lib/Sema/Sema.cpp | 13 ------------- 2 files changed, 12 insertions(+), 15 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index acdfd091d0e85..8a1fc516d7abd 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13125,8 +13125,18 @@ class Sema final { SourceLocation BuiltinLoc, SourceLocation RParenLoc); - template - bool isSyclGlobalType(QualType Ty); + template + bool isSyclGlobalType(QualType Ty) { + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + if (!RecTy) + return false; + if (auto *CTSD = dyn_cast(RecTy)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RD = Template->getTemplatedDecl()) + return RD->hasAttr(); + } + return RecTy->hasAttr(); + } private: bool SemaBuiltinPrefetch(CallExpr *TheCall); diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index bdcadb1a4ee95..5bb87e3871516 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1621,19 +1621,6 @@ static void emitCallStackNotes(Sema &S, FunctionDecl *FD) { } } -template -bool Sema::isSyclGlobalType(QualType Ty) { - const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); - if (!RecTy) - return false; - if (auto *CTSD = dyn_cast(RecTy)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); - } - return RecTy->hasAttr(); -} - namespace { /// Helper class that emits deferred diagnostic messages if an entity directly From 5ecd5456d394994ba2af1b73b04da6dea2dfe0eb Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 13:01:21 -0800 Subject: [PATCH 15/34] Address latest comments --- clang/include/clang/Basic/AttrDocs.td | 13 +++++++++---- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 ---- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/Sema.cpp | 3 ++- clang/lib/Sema/SemaDecl.cpp | 7 ++++--- clang/lib/Sema/SemaDeclCXX.cpp | 3 ++- clang/lib/Sema/SemaExpr.cpp | 6 ++++-- clang/test/SemaSYCL/device_global.cpp | 10 ---------- 8 files changed, 22 insertions(+), 26 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 39c334c92075a..a2d1860931e08 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3090,6 +3090,10 @@ def SYCLDeviceGlobalAttrDocs : Documentation { This attribute is part of support for SYCL device_global feature. [[__sycl_detail__::device_global]] attribute is used for checking restrictions on variable declarations using the device_global type instead of the class name. +Global or static variables of type decorated with this attribute have +`sycl-unique-id`, an LLVM IR attribute, added to the definition of each such +variable, which provides a unique string identifier using +__builtin_sycl_unique_stable_id. We do not intend to support this as a general attribute that customer code can use, so we have this attribute in sycl_detail namespace. @@ -3108,14 +3112,15 @@ def SYCLGlobalVariableAllowedAttrDocs : Documentation { let Content = [{ This attribute is part of support for SYCL device_global feature. [[__sycl_detail__::global_variable_allowed]] attribute is used to avoid -diagnosing an error when variables of type device_global are referenced in -device code. We do not intend to support this as a general attribute that -customer code can use, therefore it is wrapped in sycl_detail namespace. +diagnosing an error when global or static variables of type decorated with this +attribute are referenced in device code. We do not intend to support this as a +general attribute that customer code can use, therefore it is wrapped in +sycl_detail namespace. .. code-block:: c++ template - struct [[__sycl_detail__::device_global]] device_global {} + struct [[__sycl_detail__::global_variable_allowed]] device_global {} device_global Foo; }]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 54f799d69e065..b34abb5890377 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7105,10 +7105,6 @@ def err_sycl_device_global_not_publicly_accessible: Error< "'device_global' member variable %0 is not publicly accessible from namespace scope">; def err_array_of_device_global_not_allowed : Error< "array of device_global %0 is not allowed">; -def err_shadow_variable_within_same_namespace: Error< - "shadow variable %0 not allowed within the same enclosing namespace scope">; -def err_namespace_name_shadows_namespace_containing_device_global : Error< - "not allowed: namespace name shadows %0 namespace which contains device_global">; def err_unexpected_interface : Error< "unexpected interface name %0: expected expression">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 8a1fc516d7abd..81d318c005eca 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13126,7 +13126,7 @@ class Sema final { SourceLocation RParenLoc); template - bool isSyclGlobalType(QualType Ty) { + bool isDecoratedWithSyclAttribute(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 5bb87e3871516..49d9149b439e2 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1692,7 +1692,8 @@ class DeferredDiagnosticsEmitter if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { if (!S.checkAllowedSYCLInitializer(VD) && - !S.isSyclGlobalType(VD->getType())) { + !S.isDecoratedWithSyclAttribute( + VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; return; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 31a829aea6bcd..6834e6692b93e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7395,12 +7395,13 @@ NamedDecl *Sema::ActOnVariableDeclarator( // Static variables declared inside SYCL device code must be const or // constexpr if (getLangOpts().SYCLIsDevice) { - if (isSyclGlobalType(NewVD->getType()) && - SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage()) { + if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() && + isDecoratedWithSyclAttribute(NewVD->getType()) { Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); } if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && - !isSyclGlobalType(NewVD->getType())) + !isDecoratedWithSyclAttribute( + NewVD->getType())) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; } diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index e352fa026aea3..9254066ee02c6 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3531,7 +3531,8 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { - if (isSyclGlobalType(Value->getType()) && + if (isDecoratedWithSyclAttribute( + Value->getType()) && Value->getAccess() != AS_public) { Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) << Value; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 0429a8bbaa43b..eab4d2c2bf676 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -231,14 +231,16 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && !VD->hasAttr() && - !isSyclGlobalType(VD->getType())) + !isDecoratedWithSyclAttribute( + VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr() && - !isSyclGlobalType(VD->getType())) + !isDecoratedWithSyclAttribute( + VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index ef03dfef3252a..80e6b71e6459a 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -36,16 +36,6 @@ namespace { device_global same_name; // OK } -//inline namespace other { -// device_global same_name; // ILLEGAL: shadows "device_global" variable -//} // with same name in enclosing namespace scope - -//inline namespace { -// namespace foo { // ILLEGAL: namespace name shadows "::foo" -// } // namespace which contains "device_global" - // variable. -//} - int main() { cl::sycl::kernel_single_task([=]() { (void)glob; From aef34f9ed1af18321425e12d5b33aa800e73444e Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 14:06:55 -0800 Subject: [PATCH 16/34] Add back attribute to CodeGenSYCL/Inputs/sycl.hpp after merge --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index facf6dfa2fd16..6d72116688de4 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -130,8 +130,10 @@ struct no_alias { }; } // namespace property +// device_global type decorated with attributes template -class [[__sycl_detail__::device_global]] device_global { +class [[__sycl_detail__::device_global]] +[[__sycl_detail__::global_variable_allowed]] device_global { public: const T &get() const noexcept { return *Data; } device_global() {} From 4cac80e7deaec398135a25a3e2e63837fd32f3e6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 15:03:31 -0800 Subject: [PATCH 17/34] Remove additional definition of SYCLDeviceGlobal --- clang/include/clang/Basic/Attr.td | 17 ++++------------- 1 file changed, 4 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fa2f6c984ac9d..68ae04d8d461f 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1261,15 +1261,6 @@ def SYCLUsesAspects : InheritableAttr { let Documentation = [Undocumented]; } -def SYCLDeviceGlobal : InheritableAttr { - let Spellings = [CXX11<"__sycl_detail__", "device_global">]; - let Subjects = SubjectList<[CXXRecord], ErrorDiag>; - let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - // Only used internally by the SYCL implementation - let Documentation = [Undocumented]; - let SimpleHandler = 1; -} - // Marks functions which must not be vectorized via horizontal SIMT widening, // e.g. because the function is already vectorized. Used to mark SYCL // explicit SIMD kernels and functions. @@ -1449,19 +1440,19 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { } def SYCLDeviceGlobal: InheritableAttr { - let Spellings = [GNU<"device_global">, - CXX11<"__sycl_detail__", "device_global">]; + let Spellings = [CXX11<"__sycl_detail__", "device_global">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; +// Only used internally let Documentation = [SYCLDeviceGlobalAttrDocs]; let SimpleHandler = 1; } def SYCLGlobalVariableAllowed : InheritableAttr { - let Spellings = [GNU<"global_variable_allowed">, - CXX11<"__sycl_detail__", "global_variable_allowed">]; + let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; +// Only used internally let Documentation = [SYCLGlobalVariableAllowedAttrDocs]; let SimpleHandler = 1; } From 3f79c5e27fb481920ac398d0f08c43e200f5d799 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 15:34:27 -0800 Subject: [PATCH 18/34] Add test description; Remove unsupported test cases --- clang/lib/Sema/SemaDecl.cpp | 4 ++-- clang/test/CodeGenSYCL/device_global.cpp | 4 ++++ clang/test/SemaSYCL/device_global.cpp | 9 +++------ 3 files changed, 9 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 73e2705e976c6..0216bfccb4bf8 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7404,9 +7404,9 @@ NamedDecl *Sema::ActOnVariableDeclarator( // constexpr if (getLangOpts().SYCLIsDevice) { if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() && - isDecoratedWithSyclAttribute(NewVD->getType()) { + isDecoratedWithSyclAttribute(NewVD->getType())) Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); - } + if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && !isDecoratedWithSyclAttribute( NewVD->getType())) diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 23ee803190851..edbc6941539ed 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -1,6 +1,10 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s #include "sycl.hpp" +// Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the +// global variable whose type is decorated with device_global attribute, and that a +// unique string is generated. +// namespace cl { namespace sycl { namespace ext { diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 80e6b71e6459a..d57f6c8489cc4 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -1,6 +1,9 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump -verify %s | FileCheck %s #include "Inputs/sycl.hpp" +// Test cases below check for valid usage of device_global and +// global_variable_allowed attributes, and that they are being correctly +// generated in the AST using namespace sycl::ext::oneapi; device_global glob; // OK @@ -13,10 +16,6 @@ struct Foo { }; device_global Foo::d; -struct Bar { - device_global e; // ILLEGAL: non-static member variable not -}; // allowed - struct Baz { private: // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}} @@ -25,8 +24,6 @@ struct Baz { device_global Baz::f; device_global not_array; // OK -device_global h[4]; // ILLEGAL: array of "device_global" not - // allowed device_global same_name; // OK namespace foo { From 2664869d3208ed8f25a3e78c7595cccc6269cdeb Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 15:50:58 -0800 Subject: [PATCH 19/34] Fix format --- clang/include/clang/Sema/Sema.h | 3 +-- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 3 +-- clang/test/CodeGenSYCL/device_global.cpp | 14 +++++++------- clang/test/SemaSYCL/device_global.cpp | 23 +++++++++++------------ 4 files changed, 20 insertions(+), 23 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 70d7ec02c4cde..c1fc1b4b3eb0a 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13163,8 +13163,7 @@ class Sema final { SourceLocation BuiltinLoc, SourceLocation RParenLoc); - template - bool isDecoratedWithSyclAttribute(QualType Ty) { + template bool isDecoratedWithSyclAttribute(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 6d72116688de4..49af3e9c9807a 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -132,8 +132,7 @@ struct no_alias { // device_global type decorated with attributes template -class [[__sycl_detail__::device_global]] -[[__sycl_detail__::global_variable_allowed]] device_global { +class [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { public: const T &get() const noexcept { return *Data; } device_global() {} diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index edbc6941539ed..688d3945b7ce6 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -13,9 +13,10 @@ namespace oneapi { template class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { public : - const T & get() const noexcept { return *Data; } + const T &get() const noexcept { return *Data; } only_global_var_allowed() {} - operator T&() noexcept { return *Data; } + operator T &() noexcept { return *Data; } + private: T *Data; }; @@ -41,7 +42,7 @@ device_global Foo::C; device_global same_name; namespace NS { - device_global same_name; +device_global same_name; } // CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]] // CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]] @@ -60,12 +61,12 @@ void foo() { (void)same_name; (void)NS::same_name; (void)no_device_global; - }); + }); }); } namespace { - device_global same_name; +device_global same_name; } // CHECK: @_ZN12_GLOBAL__N_19same_nameE = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ANON_NS_ATTRS:[0-9]+]] @@ -75,8 +76,7 @@ void bar() { h.single_task([=]() { int A = same_name; }); }); } - -} +} // namespace // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index d57f6c8489cc4..2d2ed400af6b0 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -6,31 +6,31 @@ // generated in the AST using namespace sycl::ext::oneapi; -device_global glob; // OK -static device_global static_glob; // OK -inline device_global inline_glob; // OK +device_global glob; // OK +static device_global static_glob; // OK +inline device_global inline_glob; // OK static const device_global static_const_glob; struct Foo { - static device_global d; // OK + static device_global d; // OK }; device_global Foo::d; struct Baz { - private: +private: // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}} - static device_global f; // ILLEGAL: not publicly accessible from -}; // namespace scope + static device_global f; +}; device_global Baz::f; -device_global not_array; // OK +device_global not_array; // OK -device_global same_name; // OK +device_global same_name; // OK namespace foo { - device_global same_name; // OK +device_global same_name; // OK } namespace { - device_global same_name; // OK +device_global same_name; // OK } int main() { @@ -59,4 +59,3 @@ int main() { // CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition // CHECK: SYCLDeviceGlobalAttr {{.*}} // CHECK: SYCLGlobalVariableAllowedAttr {{.*}} - From 0469d18f07f1e288e546f7fe651edd0bfa34eee0 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 16:26:27 -0800 Subject: [PATCH 20/34] clang-format again! --- clang/test/CodeGenSYCL/device_global.cpp | 3 +-- clang/test/SemaSYCL/Inputs/sycl.hpp | 8 ++++---- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 688d3945b7ce6..f878eeba566b5 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -12,7 +12,7 @@ namespace oneapi { // decorated with only global_variable_allowed attribute template class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { -public : +public: const T &get() const noexcept { return *Data; } only_global_var_allowed() {} operator T &() noexcept { return *Data; } @@ -78,7 +78,6 @@ void bar() { } } // namespace - // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } // CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" } // CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 992b0c1b5f4a6..64de62be9a195 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -68,12 +68,12 @@ class accessor_property_list {}; // device_global type decorated with attributes template -struct [[__sycl_detail__::device_global]] -[[__sycl_detail__::global_variable_allowed]] device_global { +struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global { public: - const T & get() const noexcept { return *Data; } + const T &get() const noexcept { return *Data; } device_global() {} - operator T&() noexcept { return *Data; } + operator T &() noexcept { return *Data; } + private: T *Data; }; From d25c8163c34b61cf882b5645aae1386633529ea0 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Mar 2022 18:04:02 -0800 Subject: [PATCH 21/34] Fix lit tests; Address comments --- clang/include/clang/Basic/Attr.td | 6 +-- .../clang/Basic/DiagnosticSemaKinds.td | 2 - clang/lib/CodeGen/CodeGenModule.cpp | 7 +++- clang/lib/CodeGen/CodeGenModule.h | 4 -- clang/lib/Sema/SemaSYCL.cpp | 22 ++-------- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 1 + clang/test/CodeGenSYCL/device_global.cpp | 41 +++++++++---------- ...a-attribute-supported-attributes-list.test | 1 + clang/test/SemaSYCL/device_global.cpp | 2 +- 9 files changed, 35 insertions(+), 51 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 68ae04d8d461f..8917901d591d2 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1257,7 +1257,7 @@ def SYCLUsesAspects : InheritableAttr { let Subjects = SubjectList<[CXXRecord, Function], ErrorDiag>; let Args = [VariadicExprArgument<"Aspects">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - // Only used internally by the SYCL implementation + // Only used internally by SYCL implementation let Documentation = [Undocumented]; } @@ -1443,7 +1443,7 @@ def SYCLDeviceGlobal: InheritableAttr { let Spellings = [CXX11<"__sycl_detail__", "device_global">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; -// Only used internally +// Only used internally by SYCL implementation let Documentation = [SYCLDeviceGlobalAttrDocs]; let SimpleHandler = 1; } @@ -1452,7 +1452,7 @@ def SYCLGlobalVariableAllowed : InheritableAttr { let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; -// Only used internally +// Only used internally by SYCL implementation let Documentation = [SYCLGlobalVariableAllowedAttrDocs]; let SimpleHandler = 1; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a5c7f7cba870b..3c521df759a54 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7098,8 +7098,6 @@ def err_sycl_device_global_incorrect_scope : Error< "'device_global' variables must be static or declared at namespace scope">; def err_sycl_device_global_not_publicly_accessible: Error< "'device_global' member variable %0 is not publicly accessible from namespace scope">; -def err_array_of_device_global_not_allowed : Error< - "array of device_global %0 is not allowed">; def err_unexpected_interface : Error< "unexpected interface name %0: expected expression">; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 0b0af02fdc847..3889965fa71da 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2846,7 +2846,10 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation())); } -void CodeGenModule::addSYCLUniqueID(llvm::GlobalVariable *GV, + // Add "sycl-unique-id" llvm IR attribute for global variables marked with + // SYCL device_global attribute, and return a unique string using + // __builtin_sycl_unique_stable_id. +static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); GV->addAttribute("sycl-unique-id", builtinString); @@ -4948,6 +4951,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, if (getLangOpts().SYCLIsDevice) addGlobalIntelFPGAAnnotation(D, GV); + // If VarDecl has a type decorated with SYCL device_global attribute, emit IR + // attribute 'sycl-unique-id'. if (getLangOpts().SYCLIsDevice) { const RecordDecl *RD = D->getType()->getAsRecordDecl(); if (RD && RD->hasAttr()) diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index f2af0f4cae092..2e542c0333ec4 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1308,10 +1308,6 @@ class CodeGenModule : public CodeGenTypeCache { /// annotations are emitted during finalization of the LLVM code. void AddGlobalAnnotations(const ValueDecl *D, llvm::GlobalValue *GV); - /// Add "sycl-unique-id" llvm attribute for global variables marked with - /// SYCL device_global attribute - void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD); - bool isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn, SourceLocation Loc) const; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4a94c25bbfd17..3ab67c646abde 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -123,10 +123,6 @@ class Util { /// specialization id class. static bool isSyclSpecIdType(QualType Ty); - /// Checks whether given clang type is a full specialization of the SYCL - /// device_global class. - static bool isSyclDeviceGlobalType(QualType Ty); - /// Checks whether given clang type is a full specialization of the SYCL /// kernel_handler class. static bool isSyclKernelHandlerType(QualType Ty); @@ -4896,7 +4892,7 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) { return; // Step 1: ensure that this is of the correct type template specialization. if (!Util::isSyclSpecIdType(VD->getType()) && - !Util::isSyclDeviceGlobalType(VD->getType())) { + !S.isDecoratedWithSyclAttribute(VD->getType())) { // Handle the case where this could be a deduced type, such as a deduction // guide. We have to do this here since this function, unlike most of the // rest of this file, is called during Sema instead of after it. We will @@ -5076,7 +5072,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { // Skip if this isn't a SpecIdType or DeviceGlobal. This can happen if it // was a deduced type. if (!Util::isSyclSpecIdType(VD->getType()) && - !Util::isSyclDeviceGlobalType(VD->getType())) + !S.isDecoratedWithSyclAttribute(VD->getType())) continue; // Skip if we've already visited this. @@ -5090,7 +5086,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { Visited.insert(VD); std::string TopShim = EmitShims(OS, ShimCounter, Policy, VD); - if (Util::isSyclDeviceGlobalType(VD->getType())) { + if (S.isDecoratedWithSyclAttribute(VD->getType())) { DeviceGlobalsEmitted = true; DeviceGlobOS << "device_global_map::add("; DeviceGlobOS << "(void *)&"; @@ -5189,18 +5185,6 @@ bool Util::isSyclSpecIdType(QualType Ty) { return matchQualifiedTypeName(Ty, Scopes); } -bool Util::isSyclDeviceGlobalType(QualType Ty) { - const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); - if (!RecTy) - return false; - if (auto *CTSD = dyn_cast(RecTy)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); - } - return RecTy->hasAttr(); -} - bool Util::isSyclKernelHandlerType(QualType Ty) { std::array Scopes = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 49af3e9c9807a..7200b51695d2d 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -136,6 +136,7 @@ class [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allo public: const T &get() const noexcept { return *Data; } device_global() {} + operator T &() noexcept { return *Data; } private: T *Data; diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index f878eeba566b5..662736686930d 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -4,26 +4,6 @@ // Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the // global variable whose type is decorated with device_global attribute, and that a // unique string is generated. -// -namespace cl { -namespace sycl { -namespace ext { -namespace oneapi { -// decorated with only global_variable_allowed attribute -template -class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { -public: - const T &get() const noexcept { return *Data; } - only_global_var_allowed() {} - operator T &() noexcept { return *Data; } - -private: - T *Data; -}; -} // namespace oneapi -} // namespace ext -} // namespace sycl -} // namespace cl using namespace sycl::ext::oneapi; using namespace cl::sycl; @@ -47,10 +27,27 @@ device_global same_name; // CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]] // CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]] +// decorated with only global_variable_allowed attribute +template +class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed { +public: + const T &get() const noexcept { return *Data; } + only_global_var_allowed() {} + operator T &() noexcept { return *Data; } + +private: + T *Data; +}; + // check that we don't generate `sycl-unique-id` IR attribute if class does not use // [[__sycl_detail__::device_global]] only_global_var_allowed no_device_global; -// CHECK: @no_device_global = addrspace(1) global %"class.cl::sycl::ext::oneapi::only_global_var_allowed" zeroinitializer, align 8{{$}} +// CHECK: @no_device_global = addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}} + +inline namespace Bar { +device_global InlineNS; +} +// CHECK: @_ZN3Bar8InlineNSE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #[[BAR_INLINENS_ATTRS:[0-9]+]] void foo() { q.submit([&](handler &h) { @@ -61,6 +58,7 @@ void foo() { (void)same_name; (void)NS::same_name; (void)no_device_global; + (void)Bar::InlineNS; }); }); } @@ -83,4 +81,5 @@ void bar() { // CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" } // CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" } // CHECK: attributes #[[SAME_NAME_NS_ATTRS]] = { "sycl-unique-id"="_ZN2NS9same_nameE" } +// CHECK: attributes #[[BAR_INLINENS_ATTRS]] = { "sycl-unique-id"="_ZN3Bar8InlineNSE" } // CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" } diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 4c1c3e14bca59..59908bbe226f7 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -160,6 +160,7 @@ // CHECK-NEXT: SYCLDeviceGlobal (SubjectMatchRule_record) // CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) +// CHECK-NEXT: SYCLGlobalVariableAllowed (SubjectMatchRule_record) // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function) diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 2d2ed400af6b0..324aba7999321 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -3,7 +3,7 @@ // Test cases below check for valid usage of device_global and // global_variable_allowed attributes, and that they are being correctly -// generated in the AST +// generated in the AST. using namespace sycl::ext::oneapi; device_global glob; // OK From 19f62a51891746b01452864e6ccb3114030a4717 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 4 Mar 2022 13:08:12 -0800 Subject: [PATCH 22/34] Fix lint --- clang/lib/CodeGen/CodeGenModule.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3889965fa71da..fe026809b620d 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2846,11 +2846,10 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation())); } - // Add "sycl-unique-id" llvm IR attribute for global variables marked with - // SYCL device_global attribute, and return a unique string using - // __builtin_sycl_unique_stable_id. -static void addSYCLUniqueID(llvm::GlobalVariable *GV, - const VarDecl *VD) { +// Add "sycl-unique-id" llvm IR attribute for global variables marked with +// SYCL device_global attribute, and return a unique string using +// __builtin_sycl_unique_stable_id. +static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); GV->addAttribute("sycl-unique-id", builtinString); } From 0354ea5183a2592633c0a27102b20dae54e0cedf Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 4 Mar 2022 13:16:19 -0800 Subject: [PATCH 23/34] Lint again :( --- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index fe026809b620d..14aff38ad0d74 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2847,7 +2847,7 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, } // Add "sycl-unique-id" llvm IR attribute for global variables marked with -// SYCL device_global attribute, and return a unique string using +// SYCL device_global attribute, and return a unique string using // __builtin_sycl_unique_stable_id. static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); From b45159d59d253ceeb58bb93cccc62d8979a1d4da Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 7 Mar 2022 13:58:50 -0800 Subject: [PATCH 24/34] Fix build failure; Add comment; Fix indentation --- clang/include/clang/Basic/Attr.td | 4 ++-- clang/lib/CodeGen/CodeGenModule.cpp | 5 +++-- clang/lib/Sema/SemaDeclCXX.cpp | 2 ++ 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 8917901d591d2..b544ff38a54d2 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1443,7 +1443,7 @@ def SYCLDeviceGlobal: InheritableAttr { let Spellings = [CXX11<"__sycl_detail__", "device_global">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; -// Only used internally by SYCL implementation + // Only used internally by SYCL implementation let Documentation = [SYCLDeviceGlobalAttrDocs]; let SimpleHandler = 1; } @@ -1452,7 +1452,7 @@ def SYCLGlobalVariableAllowed : InheritableAttr { let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">]; let Subjects = SubjectList<[CXXRecord], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; -// Only used internally by SYCL implementation + // Only used internally by SYCL implementation let Documentation = [SYCLGlobalVariableAllowedAttrDocs]; let SimpleHandler = 1; } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 14aff38ad0d74..3f343f1ebfbd9 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2849,7 +2849,8 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, // Add "sycl-unique-id" llvm IR attribute for global variables marked with // SYCL device_global attribute, and return a unique string using // __builtin_sycl_unique_stable_id. -static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD) { +static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD, + ASTContext &Context) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); GV->addAttribute("sycl-unique-id", builtinString); } @@ -4955,7 +4956,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, if (getLangOpts().SYCLIsDevice) { const RecordDecl *RD = D->getType()->getAsRecordDecl(); if (RD && RD->hasAttr()) - addSYCLUniqueID(GV, D); + addSYCLUniqueID(GV, D, Context); } if (D->getType().isRestrictQualified()) { diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index f68d1302bc29d..ea2cc5981097c 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3529,6 +3529,8 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, } } + // Emit diagnostic if a private member of type decorated with device_global + // attribute is accessed if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { if (isDecoratedWithSyclAttribute( From 3baff38d0a3a7905aed8a5a83344b1c970a74af9 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 9 Mar 2022 10:52:49 -0800 Subject: [PATCH 25/34] Fix lit test SemaSYCL/explicit-cast-to-generic.cpp --- clang/test/SemaSYCL/Inputs/sycl.hpp | 2 +- clang/test/SemaSYCL/explicit-cast-to-generic.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 64de62be9a195..abc4358d739ff 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -365,7 +365,7 @@ template class multi_ptr { pointer_t m_Pointer; public: - multi_ptr(T *Ptr) : m_Pointer((pointer_t)(Ptr)) {} + multi_ptr(T *Ptr) : m_Pointer((pointer_t)(Ptr)) {} // #MultiPtrConstructor pointer_t get() { return m_Pointer; } }; diff --git a/clang/test/SemaSYCL/explicit-cast-to-generic.cpp b/clang/test/SemaSYCL/explicit-cast-to-generic.cpp index 6fdcf14f3e296..f8a1583a8669b 100644 --- a/clang/test/SemaSYCL/explicit-cast-to-generic.cpp +++ b/clang/test/SemaSYCL/explicit-cast-to-generic.cpp @@ -15,7 +15,7 @@ void __attribute__((sycl_device)) onDeviceUsages(multi_ptr::pointer_t' (aka '__private int *') potentially leads to an invalid address space cast in the resulting code}} + // expected-warning@#MultiPtrConstructor {{explicit cast from 'int *' to 'sycl::multi_ptr::pointer_t' (aka '__private int *') potentially leads to an invalid address space cast in the resulting code}} // expected-note@+1 {{called by 'onDeviceUsages'}} auto P = multi_ptr{F.get()}; From 5bced1a651bb78c46754af0c26a75a006950232f Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 9 Mar 2022 18:33:52 +0300 Subject: [PATCH 26/34] Emit generic addrspace in llvm.used and llvm.global_ctors Also adds a couple of test cases with templates --- clang/lib/CodeGen/CodeGenModule.cpp | 25 +++++++++++++++++------- clang/test/CodeGenSYCL/device_global.cpp | 18 +++++++++++++++++ 2 files changed, 36 insertions(+), 7 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3f343f1ebfbd9..920c466237f56 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1594,10 +1594,14 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) { llvm::FunctionType* CtorFTy = llvm::FunctionType::get(VoidTy, false); llvm::Type *CtorPFTy = llvm::PointerType::get(CtorFTy, TheModule.getDataLayout().getProgramAddressSpace()); + llvm::PointerType *TargetType = VoidPtrTy; + if (getLangOpts().SYCLIsDevice) + TargetType = llvm::IntegerType::getInt8PtrTy( + getLLVMContext(), getContext().getTargetAddressSpace(LangAS::Default)); // Get the type of a ctor entry, { i32, void ()*, i8* }. llvm::StructType *CtorStructTy = llvm::StructType::get( - Int32Ty, CtorPFTy, VoidPtrTy); + Int32Ty, CtorPFTy, TargetType); // Construct the constructor and destructor arrays. ConstantInitBuilder builder(*this); @@ -1606,10 +1610,11 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) { auto ctor = ctors.beginStruct(CtorStructTy); ctor.addInt(Int32Ty, I.Priority); ctor.add(llvm::ConstantExpr::getBitCast(I.Initializer, CtorPFTy)); - if (I.AssociatedData) - ctor.add(llvm::ConstantExpr::getBitCast(I.AssociatedData, VoidPtrTy)); - else - ctor.addNullPointer(VoidPtrTy); + if (I.AssociatedData) { + ctor.add(llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( + I.AssociatedData, TargetType)); + } else + ctor.addNullPointer(TargetType); ctor.finishAndAddTo(ctors); } @@ -2429,18 +2434,24 @@ static void emitUsed(CodeGenModule &CGM, StringRef Name, if (List.empty()) return; + llvm::PointerType *TargetType = CGM.Int8PtrTy; + if (CGM.getLangOpts().SYCLIsDevice) + TargetType = llvm::IntegerType::getInt8PtrTy( + CGM.getLLVMContext(), + CGM.getContext().getTargetAddressSpace(LangAS::Default)); + // Convert List to what ConstantArray needs. SmallVector UsedArray; UsedArray.resize(List.size()); for (unsigned i = 0, e = List.size(); i != e; ++i) { UsedArray[i] = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( - cast(&*List[i]), CGM.Int8PtrTy); + cast(&*List[i]), TargetType); } if (UsedArray.empty()) return; - llvm::ArrayType *ATy = llvm::ArrayType::get(CGM.Int8PtrTy, UsedArray.size()); + llvm::ArrayType *ATy = llvm::ArrayType::get(TargetType, UsedArray.size()); auto *GV = new llvm::GlobalVariable( CGM.getModule(), ATy, false, llvm::GlobalValue::AppendingLinkage, diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 662736686930d..c6150010f7d49 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -49,6 +49,17 @@ device_global InlineNS; } // CHECK: @_ZN3Bar8InlineNSE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #[[BAR_INLINENS_ATTRS:[0-9]+]] +template struct TS { +public: + static device_global d; +}; +template<> device_global TS::d{}; +// CHECK: @_ZN2TSIiE1dE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[TEMPLATED_WRAPPER_ATTRS:[0-9]+]] + +template +device_global templ_dev_global; +// CHECK: @[[TEMPL_DEV_GLOB:[a-zA-Z0-9_]+]] = linkonce_odr addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 8 #[[TEMPL_DEV_GLOB_ATTRS:[0-9]+]] + void foo() { q.submit([&](handler &h) { h.single_task([=]() { @@ -59,6 +70,8 @@ void foo() { (void)NS::same_name; (void)no_device_global; (void)Bar::InlineNS; + auto AA = TS::d.get(); + auto val = templ_dev_global.get(); }); }); } @@ -76,10 +89,15 @@ void bar() { } } // namespace +// CHECK: @llvm.global_ctors = appending global [2 x { i32, void ()*, i8 addrspace(4)* }] [{ i32, void ()*, i8 addrspace(4)* } { i32 65535, void ()* @__cxx_global_var_init{{.*}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @[[TEMPL_DEV_GLOB]] to i8 addrspace(1)*) to i8 addrspace(4)*) }, { i32, void ()*, i8 addrspace(4)* } { i32 65535, void ()* @_GLOBAL__sub_I_device_global.cpp, i8 addrspace(4)* null }] +// CHECK: @llvm.used = appending global [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @[[TEMPL_DEV_GLOB]] to i8 addrspace(1)*) to i8 addrspace(4)*)], section "llvm.metadata" + // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } // CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" } // CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" } // CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" } // CHECK: attributes #[[SAME_NAME_NS_ATTRS]] = { "sycl-unique-id"="_ZN2NS9same_nameE" } // CHECK: attributes #[[BAR_INLINENS_ATTRS]] = { "sycl-unique-id"="_ZN3Bar8InlineNSE" } +// CHECK: attributes #[[TEMPLATED_WRAPPER_ATTRS]] = { "sycl-unique-id"="_ZN2TSIiE1dE" } +// CHECK: attributes #[[TEMPL_DEV_GLOB_ATTRS]] = { "sycl-unique-id"="_Z16templ_dev_globalIiE" } // CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" } From 505a4f2019e3f8a6e20183a590dca67e4312c671 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 9 Mar 2022 12:29:27 -0800 Subject: [PATCH 27/34] Fix format --- clang/lib/CodeGen/CodeGenModule.cpp | 9 ++++----- clang/test/CodeGenSYCL/device_global.cpp | 2 +- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 920c466237f56..e5dd06b6c0847 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1600,8 +1600,8 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) { getLLVMContext(), getContext().getTargetAddressSpace(LangAS::Default)); // Get the type of a ctor entry, { i32, void ()*, i8* }. - llvm::StructType *CtorStructTy = llvm::StructType::get( - Int32Ty, CtorPFTy, TargetType); + llvm::StructType *CtorStructTy = + llvm::StructType::get(Int32Ty, CtorPFTy, TargetType); // Construct the constructor and destructor arrays. ConstantInitBuilder builder(*this); @@ -2444,9 +2444,8 @@ static void emitUsed(CodeGenModule &CGM, StringRef Name, SmallVector UsedArray; UsedArray.resize(List.size()); for (unsigned i = 0, e = List.size(); i != e; ++i) { - UsedArray[i] = - llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( - cast(&*List[i]), TargetType); + UsedArray[i] = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( + cast(&*List[i]), TargetType); } if (UsedArray.empty()) diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index c6150010f7d49..c1b68db290282 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -53,7 +53,7 @@ template struct TS { public: static device_global d; }; -template<> device_global TS::d{}; +template <> device_global TS::d{}; // CHECK: @_ZN2TSIiE1dE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[TEMPLATED_WRAPPER_ATTRS:[0-9]+]] template From 4905d70a613c9de7c0c954534db51824e1bdbe62 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 10 Mar 2022 12:37:26 -0800 Subject: [PATCH 28/34] Add comments; rename method; add separate AST test --- clang/include/clang/Sema/Sema.h | 3 ++- clang/lib/Sema/Sema.cpp | 2 +- clang/lib/Sema/SemaDecl.cpp | 12 ++++++++---- clang/lib/Sema/SemaDeclCXX.cpp | 4 ++-- clang/lib/Sema/SemaExpr.cpp | 4 ++-- clang/lib/Sema/SemaSYCL.cpp | 9 ++++++--- clang/test/SemaSYCL/device_global.cpp | 10 +--------- clang/test/SemaSYCL/device_global_ast.cpp | 17 +++++++++++++++++ 8 files changed, 39 insertions(+), 22 deletions(-) create mode 100644 clang/test/SemaSYCL/device_global_ast.cpp diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c1fc1b4b3eb0a..e6119c24630e0 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13163,7 +13163,8 @@ class Sema final { SourceLocation BuiltinLoc, SourceLocation RParenLoc); - template bool isDecoratedWithSyclAttribute(QualType Ty) { + template + bool isTypeDecoratedWithDeclAttribute(QualType Ty) { const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); if (!RecTy) return false; diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 903a4795528be..1c5df431e07dd 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1704,7 +1704,7 @@ class DeferredDiagnosticsEmitter if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { if (!S.checkAllowedSYCLInitializer(VD) && - !S.isDecoratedWithSyclAttribute( + !S.isTypeDecoratedWithDeclAttribute( VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 0216bfccb4bf8..cb0eb43d8c8b6 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7400,15 +7400,19 @@ NamedDecl *Sema::ActOnVariableDeclarator( NewVD->setTSCSpec(TSCS); } - // Static variables declared inside SYCL device code must be const or - // constexpr + // Global variables with types decorated with device_global attribute must be + // static if they are declared in SYCL device code. if (getLangOpts().SYCLIsDevice) { if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() && - isDecoratedWithSyclAttribute(NewVD->getType())) + isTypeDecoratedWithDeclAttribute( + NewVD->getType())) Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); + // Static variables declared inside SYCL device code must be const or + // constexpr unless their types are decorated with global_variable_allowed + // attribute. if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && - !isDecoratedWithSyclAttribute( + !isTypeDecoratedWithDeclAttribute( NewVD->getType())) SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index ea2cc5981097c..a1e763ee73cbb 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3530,10 +3530,10 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, } // Emit diagnostic if a private member of type decorated with device_global - // attribute is accessed + // attribute is accessed. if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { - if (isDecoratedWithSyclAttribute( + if (isTypeDecoratedWithDeclAttribute( Value->getType()) && Value->getAccess() != AS_public) { Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index d762c7cb0c257..eecef13454109 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -231,7 +231,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && !VD->hasAttr() && - !isDecoratedWithSyclAttribute( + !isTypeDecoratedWithDeclAttribute( VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; @@ -239,7 +239,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr() && - !isDecoratedWithSyclAttribute( + !isTypeDecoratedWithDeclAttribute( VD->getType())) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3ab67c646abde..0db660f1dfd8b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4892,7 +4892,8 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) { return; // Step 1: ensure that this is of the correct type template specialization. if (!Util::isSyclSpecIdType(VD->getType()) && - !S.isDecoratedWithSyclAttribute(VD->getType())) { + !S.isTypeDecoratedWithDeclAttribute( + VD->getType())) { // Handle the case where this could be a deduced type, such as a deduction // guide. We have to do this here since this function, unlike most of the // rest of this file, is called during Sema instead of after it. We will @@ -5072,7 +5073,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { // Skip if this isn't a SpecIdType or DeviceGlobal. This can happen if it // was a deduced type. if (!Util::isSyclSpecIdType(VD->getType()) && - !S.isDecoratedWithSyclAttribute(VD->getType())) + !S.isTypeDecoratedWithDeclAttribute( + VD->getType())) continue; // Skip if we've already visited this. @@ -5086,7 +5088,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { Visited.insert(VD); std::string TopShim = EmitShims(OS, ShimCounter, Policy, VD); - if (S.isDecoratedWithSyclAttribute(VD->getType())) { + if (S.isTypeDecoratedWithDeclAttribute( + VD->getType())) { DeviceGlobalsEmitted = true; DeviceGlobOS << "device_global_map::add("; DeviceGlobOS << "(void *)&"; diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 324aba7999321..f6567a95a6d57 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump -verify %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" // Test cases below check for valid usage of device_global and @@ -51,11 +51,3 @@ int main() { static device_global non_const_static; }); } -// -// CHECK: ClassTemplateDecl {{.*}} device_global -// CHECK: CXXRecordDecl {{.*}} struct device_global definition -// CHECK: SYCLDeviceGlobalAttr {{.*}} -// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} -// CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition -// CHECK: SYCLDeviceGlobalAttr {{.*}} -// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} diff --git a/clang/test/SemaSYCL/device_global_ast.cpp b/clang/test/SemaSYCL/device_global_ast.cpp new file mode 100644 index 0000000000000..6e618808f7048 --- /dev/null +++ b/clang/test/SemaSYCL/device_global_ast.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump %s | FileCheck %s +#include "Inputs/sycl.hpp" + +// Test cases below check that DeviceGlobalAttr and GlobalVariableAllowedAttr +// are correctly emitted. + +using namespace sycl::ext::oneapi; + +device_global glob; +// CHECK: ClassTemplateDecl {{.*}} device_global +// CHECK: CXXRecordDecl {{.*}} struct device_global definition +// CHECK: SYCLDeviceGlobalAttr {{.*}} +// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} +// CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition +// CHECK: SYCLDeviceGlobalAttr {{.*}} +// CHECK: SYCLGlobalVariableAllowedAttr {{.*}} + From 3817bf05e52028eed4e349ba6fe48f6e03b67306 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 10 Mar 2022 15:25:05 -0800 Subject: [PATCH 29/34] Fix lint that git-clang-format didn't catch :( --- clang/test/SemaSYCL/device_global_ast.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/SemaSYCL/device_global_ast.cpp b/clang/test/SemaSYCL/device_global_ast.cpp index 6e618808f7048..e39a41ab7fa24 100644 --- a/clang/test/SemaSYCL/device_global_ast.cpp +++ b/clang/test/SemaSYCL/device_global_ast.cpp @@ -14,4 +14,3 @@ device_global glob; // CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition // CHECK: SYCLDeviceGlobalAttr {{.*}} // CHECK: SYCLGlobalVariableAllowedAttr {{.*}} - From 1c53934b473317419bf4f2e73eceffff985ff0e6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 11 Mar 2022 12:52:05 -0800 Subject: [PATCH 30/34] Add case where device_global attributes are applied to the wrong subject --- clang/test/SemaSYCL/device_global.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index f6567a95a6d57..3caf3eb54c1aa 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -33,6 +33,10 @@ namespace { device_global same_name; // OK } +// expected-error@+2{{'device_global' attribute only applies to classes}} +// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}} +[[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] int integer; + int main() { cl::sycl::kernel_single_task([=]() { (void)glob; From e489d50c71e7a3cea2ca6fecf3d9611f03c1985a Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 11 Mar 2022 13:08:36 -0800 Subject: [PATCH 31/34] Attribute doesn't apply to this type --- clang/test/SemaSYCL/device_global.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 3caf3eb54c1aa..66296eaf923dd 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -37,6 +37,12 @@ device_global same_name; // OK // expected-error@+1{{'global_variable_allowed' attribute only applies to classes}} [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] int integer; +// expected-error@+2{{'device_global' attribute only applies to classes}} +// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}} +[[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] int *pointer; + +union [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] a_union; + int main() { cl::sycl::kernel_single_task([=]() { (void)glob; From 0709448c000346eeb6aa3761e1988f30506a5f49 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 11 Mar 2022 14:45:02 -0800 Subject: [PATCH 32/34] Add comments in CodeGenModule.cpp --- clang/lib/CodeGen/CodeGenModule.cpp | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index e5dd06b6c0847..25ec4a6611935 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1595,6 +1595,9 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) { llvm::Type *CtorPFTy = llvm::PointerType::get(CtorFTy, TheModule.getDataLayout().getProgramAddressSpace()); llvm::PointerType *TargetType = VoidPtrTy; + // Get target type when templated global variables are used, + // to emit them correctly in the target (default) address space and avoid + // emitting them in a private address space. if (getLangOpts().SYCLIsDevice) TargetType = llvm::IntegerType::getInt8PtrTy( getLLVMContext(), getContext().getTargetAddressSpace(LangAS::Default)); @@ -1610,10 +1613,11 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) { auto ctor = ctors.beginStruct(CtorStructTy); ctor.addInt(Int32Ty, I.Priority); ctor.add(llvm::ConstantExpr::getBitCast(I.Initializer, CtorPFTy)); - if (I.AssociatedData) { + // Emit appropriate bitcasts for pointers of different address spaces. + if (I.AssociatedData) ctor.add(llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( I.AssociatedData, TargetType)); - } else + else ctor.addNullPointer(TargetType); ctor.finishAndAddTo(ctors); } @@ -2433,7 +2437,9 @@ static void emitUsed(CodeGenModule &CGM, StringRef Name, // Don't create llvm.used if there is no need. if (List.empty()) return; - + // Get target type when templated global variables are used, + // to emit them correctly in the target (default) address space and avoid + // emitting them in a private address space. llvm::PointerType *TargetType = CGM.Int8PtrTy; if (CGM.getLangOpts().SYCLIsDevice) TargetType = llvm::IntegerType::getInt8PtrTy( From d4647f676317e525bf29f83ff22d04a4a6b01f4a Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 14 Mar 2022 10:09:15 -0700 Subject: [PATCH 33/34] Change comments --- clang/lib/CodeGen/CodeGenModule.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 25ec4a6611935..fe88f5c99f5c9 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2437,9 +2437,9 @@ static void emitUsed(CodeGenModule &CGM, StringRef Name, // Don't create llvm.used if there is no need. if (List.empty()) return; - // Get target type when templated global variables are used, - // to emit them correctly in the target (default) address space and avoid - // emitting them in a private address space. + // For SYCL emit pointers in the default address space which is a superset of + // other address spaces, so that casts from any other address spaces will be + // valid. llvm::PointerType *TargetType = CGM.Int8PtrTy; if (CGM.getLangOpts().SYCLIsDevice) TargetType = llvm::IntegerType::getInt8PtrTy( @@ -2862,9 +2862,9 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D, Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation())); } -// Add "sycl-unique-id" llvm IR attribute for global variables marked with -// SYCL device_global attribute, and return a unique string using -// __builtin_sycl_unique_stable_id. +// Add "sycl-unique-id" llvm IR attribute that has a unique string generated +// by __builtin_sycl_unique_stable_id for global variables marked with +// SYCL device_global attribute. static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD, ASTContext &Context) { auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD); From e4ab3bd7b29a4ad77e4f60c75e72d366f5e693bb Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 14 Mar 2022 20:11:02 -0700 Subject: [PATCH 34/34] More diagnostic work --- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 34 ++++++++++++++++++- 1 file changed, 33 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 590ab0c355e79..383cbc5e7cf50 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1517,8 +1517,29 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D, // Only add this if we aren't instantiating a variable template. We'll end up // adding the VarTemplateSpecializationDecl later. - if (!InstantiatingVarTemplate) + if (!InstantiatingVarTemplate) { SemaRef.addSyclVarDecl(Var); + if (SemaRef.getLangOpts().SYCLIsDevice) { + if (SemaRef.isDecoratedWithDeclAttribute(Var->getType())) { + if (!Var->hasGlobalStorage() || Var->isLocalVarDeclOrParm()) { + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope); + } + + if (Var->isStaticLocal()) { + const DeclContext *DeclCtx = Var->getDeclContext(); + while (!DeclCtx->isTranslationUnit()) { + if (isa(DeclCtx)) { + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope); + break; + } + DeclCtx = DeclCtx->getParent(); + } + } + } + } + } return Var; } @@ -1607,6 +1628,17 @@ Decl *TemplateDeclInstantiator::VisitFieldDecl(FieldDecl *D) { Field->setImplicit(D->isImplicit()); Field->setAccess(D->getAccess()); + // Static members are not processed here, so error out if we have a device + // global without checking access modifier. + if (SemaRef.getLangOpts().SYCLIsDevice) { + if (auto Value = dyn_cast(Field)) { + if (SemaRef.isDecoratedWithDeclAttribute(Value->getType())) { + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope) + << Value; + } + } + } Owner->addDecl(Field); return Field;