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 ef0f5d01d0ab46df1ea783057843a6479a350a1c Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 15 Mar 2022 13:28:42 -0700 Subject: [PATCH 34/34] Change "customer code" to "user code" in documentation --- clang/include/clang/Basic/AttrDocs.td | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 4868e7e041e71..462df1357fb58 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3044,8 +3044,8 @@ 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. +We do not intend to support this as a general attribute that user code can use, +so we have this attribute in sycl_detail namespace. .. code-block:: c++ @@ -3064,7 +3064,7 @@ 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 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 +general attribute that user code can use, therefore it is wrapped in sycl_detail namespace. .. code-block:: c++