From 682277f050217fcb60d03f8d25a4ae7b26825446 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Wed, 12 May 2021 19:33:29 -0700 Subject: [PATCH 01/17] [clang] Add sycl global var attribute Normally global variables are disallowed within kernels, but the presence of this new sycl_global_var attribute will cause sema to allow that particular global variable. --- clang/include/clang/Basic/Attr.td | 7 +++++++ clang/lib/Sema/SemaDeclAttr.cpp | 3 +++ clang/lib/Sema/SemaExpr.cpp | 4 ++-- .../Misc/pragma-attribute-supported-attributes-list.test | 1 + clang/test/SemaSYCL/invalid-kernel-arguments.cpp | 5 +++++ 5 files changed, 18 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a05fe355d314d..0ac90992c273b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1197,6 +1197,13 @@ def SYCLDevice : InheritableAttr { let Documentation = [SYCLDeviceDocs]; } +def SYCLGlobalVar : InheritableAttr { + let Spellings = [GNU<"sycl_global_var">]; + let Subjects = SubjectList<[GlobalVar]>; + let LangOpts = [SYCLIsDevice]; + let Documentation = [Undocumented]; +} + def SYCLKernel : InheritableAttr { let Spellings = [Clang<"sycl_kernel">]; let Subjects = SubjectList<[FunctionTmpl]>; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 548fa364d08d9..fc8dd52915771 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -9171,6 +9171,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLDeviceIndirectlyCallable: handleSYCLDeviceIndirectlyCallableAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLGlobalVar: + handleSimpleAttribute(S, D, AL); + break; case ParsedAttr::AT_SYCLRegisterNum: handleSYCLRegisterNumAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index cf3bc4300247b..fb13f908ac896 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -223,9 +223,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; - // Non-const globals are allowed for SYCL explicit SIMD. + // Non-const globals are allowed for SYCL explicit SIMD or with attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage()) + VD->hasGlobalStorage() && !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/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 80f2ac881ea0b..c3c38bd645b83 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -152,6 +152,7 @@ // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) +// CHECK-NEXT: SYCLGlobalVar (SubjectMatchRule_variable_is_global) // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function) diff --git a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp index dd0679ffef511..2f8288b0d2fc1 100644 --- a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp +++ b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp @@ -48,6 +48,9 @@ class H { int A; }; +int GlobalNoAttribute; +__attribute__((sycl_global_var)) int GlobalWithAttribute; + int main() { A Obj{}; D Obj1{}; @@ -65,6 +68,8 @@ int main() { (void)Obj4; (void)Obj5; (void)Obj6; + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + (void)GlobalWithAttribute; }); return 0; } From 2c6f8fa90bac8305752300dc98ab407546f8f804 Mon Sep 17 00:00:00 2001 From: "Mott, Jeffrey T" <59898758+jtmott-intel@users.noreply.github.com> Date: Thu, 13 May 2021 07:58:46 -0700 Subject: [PATCH 02/17] Update clang/lib/Sema/SemaExpr.cpp Clarified attribute name in comment. Co-authored-by: premanandrao <47116977+premanandrao@users.noreply.github.com> --- clang/lib/Sema/SemaExpr.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index fb13f908ac896..7603dcee28278 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -223,7 +223,8 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; - // Non-const globals are allowed for SYCL explicit SIMD or with attribute. + // Non-const globals are allowed for SYCL explicit SIMD or with the + // SYCLGlobalVar attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr()) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) From 4184aa49f9e44eb1a7f66abc391a8bfa037ec8ff Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Thu, 13 May 2021 09:44:02 -0700 Subject: [PATCH 03/17] Use SimpleHandler = 1 and revert changes to SemaDeclAttr --- clang/include/clang/Basic/Attr.td | 1 + clang/lib/Sema/SemaDeclAttr.cpp | 3 --- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 0ac90992c273b..d487b791d6a95 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1202,6 +1202,7 @@ def SYCLGlobalVar : InheritableAttr { let Subjects = SubjectList<[GlobalVar]>; let LangOpts = [SYCLIsDevice]; let Documentation = [Undocumented]; + let SimpleHandler = 1; } def SYCLKernel : InheritableAttr { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index fc8dd52915771..548fa364d08d9 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -9171,9 +9171,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLDeviceIndirectlyCallable: handleSYCLDeviceIndirectlyCallableAttr(S, D, AL); break; - case ParsedAttr::AT_SYCLGlobalVar: - handleSimpleAttribute(S, D, AL); - break; case ParsedAttr::AT_SYCLRegisterNum: handleSYCLRegisterNumAttr(S, D, AL); break; From 3f7887ff0a28aa5dd1e91a2e56651fe84102d3b9 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Thu, 13 May 2021 12:52:39 -0700 Subject: [PATCH 04/17] Wrapped attribute in pre-defined macro --- clang/lib/Frontend/InitPreprocessor.cpp | 1 + clang/test/SemaSYCL/invalid-kernel-arguments.cpp | 2 ++ 2 files changed, 3 insertions(+) diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index fff38ed036d96..93147861f060b 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1145,6 +1145,7 @@ static void InitializePredefinedMacros(const TargetInfo &TI, // SYCL device compiler which doesn't produce host binary. if (LangOpts.SYCLIsDevice) { Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1"); + Builder.defineMacro("__SYCL_GLOBAL_VAR__", "__attribute__((sycl_global_var))"); Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))"); // Enable __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ macro for diff --git a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp index 2f8288b0d2fc1..eba0abebf65f5 100644 --- a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp +++ b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp @@ -50,6 +50,7 @@ class H { int GlobalNoAttribute; __attribute__((sycl_global_var)) int GlobalWithAttribute; +__SYCL_GLOBAL_VAR__ int GlobalWithAttributeViaMacro; int main() { A Obj{}; @@ -70,6 +71,7 @@ int main() { (void)Obj6; (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} (void)GlobalWithAttribute; + (void)GlobalWithAttributeViaMacro; }); return 0; } From 63c4f17d781bb708df355d6bc1c426238a0ba8a7 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Fri, 14 May 2021 20:16:18 -0700 Subject: [PATCH 05/17] Add tests and docs, and remove predefined macro --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Basic/AttrDocs.td | 10 +++++++++ clang/lib/Frontend/InitPreprocessor.cpp | 1 - clang/test/SemaSYCL/attr-syclglobalvar.cpp | 21 +++++++++++++++++++ .../SemaSYCL/invalid-kernel-arguments.cpp | 7 ------- 5 files changed, 32 insertions(+), 9 deletions(-) create mode 100644 clang/test/SemaSYCL/attr-syclglobalvar.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index d487b791d6a95..fa1446d827505 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1201,7 +1201,7 @@ def SYCLGlobalVar : InheritableAttr { let Spellings = [GNU<"sycl_global_var">]; let Subjects = SubjectList<[GlobalVar]>; let LangOpts = [SYCLIsDevice]; - let Documentation = [Undocumented]; + let Documentation = [SYCLGlobalVarDocs]; let SimpleHandler = 1; } diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 6e033f9e083b2..2ba8da3e3de9e 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3155,6 +3155,16 @@ implicitly inherit this attribute. }]; } +def SYCLGlobalVarDocs : Documentation { + let Category = DocCatFunction; + let Heading = "SYCL Global Var"; + let Content = [{ +This attribute can only be applied to global variables and indicates that the +variable can be used in a SYCL kernel. Access to global variable is expected to +be only from device-side. + }]; +} + def RISCVInterruptDocs : Documentation { let Category = DocCatFunction; let Heading = "interrupt (RISCV)"; diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 93147861f060b..fff38ed036d96 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1145,7 +1145,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI, // SYCL device compiler which doesn't produce host binary. if (LangOpts.SYCLIsDevice) { Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1"); - Builder.defineMacro("__SYCL_GLOBAL_VAR__", "__attribute__((sycl_global_var))"); Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))"); // Enable __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ macro for diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp new file mode 100644 index 0000000000000..b97858b9be817 --- /dev/null +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -0,0 +1,21 @@ +// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s + +#include "Inputs/sycl.hpp" + +__attribute__((sycl_global_var)) int GlobalWithAttribute; + +int GlobalNoAttribute; + +// expected-error@+1 {{attribute takes no arguments}} +__attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; + +// expected-warning@+1 {{attribute only applies to global variables}} +__attribute__((sycl_global_var)) void F() { + // expected-warning@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)GlobalWithAttribute; + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + }); +} \ No newline at end of file diff --git a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp index eba0abebf65f5..dd0679ffef511 100644 --- a/clang/test/SemaSYCL/invalid-kernel-arguments.cpp +++ b/clang/test/SemaSYCL/invalid-kernel-arguments.cpp @@ -48,10 +48,6 @@ class H { int A; }; -int GlobalNoAttribute; -__attribute__((sycl_global_var)) int GlobalWithAttribute; -__SYCL_GLOBAL_VAR__ int GlobalWithAttributeViaMacro; - int main() { A Obj{}; D Obj1{}; @@ -69,9 +65,6 @@ int main() { (void)Obj4; (void)Obj5; (void)Obj6; - (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} - (void)GlobalWithAttribute; - (void)GlobalWithAttributeViaMacro; }); return 0; } From 8b721d9239305193e78bd90d718f30d1741937e4 Mon Sep 17 00:00:00 2001 From: "Mott, Jeffrey T" <59898758+jtmott-intel@users.noreply.github.com> Date: Mon, 17 May 2021 09:55:19 -0700 Subject: [PATCH 06/17] Update clang/include/clang/Basic/AttrDocs.td Tweak docs verbiage. Co-authored-by: Aaron Ballman --- clang/include/clang/Basic/AttrDocs.td | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2ba8da3e3de9e..1201d4da25946 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3159,8 +3159,8 @@ def SYCLGlobalVarDocs : Documentation { let Category = DocCatFunction; let Heading = "SYCL Global Var"; let Content = [{ -This attribute can only be applied to global variables and indicates that the -variable can be used in a SYCL kernel. Access to global variable is expected to +This attribute can only be applied to a global variable and indicates that the +variable can be used in a SYCL kernel. Access to the global variable is expected to be only from device-side. }]; } From 08d0c6758aae04f74461c2c5e93e527c07380dbc Mon Sep 17 00:00:00 2001 From: "Mott, Jeffrey T" <59898758+jtmott-intel@users.noreply.github.com> Date: Mon, 17 May 2021 10:14:37 -0700 Subject: [PATCH 07/17] Update clang/include/clang/Basic/AttrDocs.td Remove custom heading. Co-authored-by: Aaron Ballman --- clang/include/clang/Basic/AttrDocs.td | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 1201d4da25946..6e195f179fb4b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3157,7 +3157,6 @@ implicitly inherit this attribute. def SYCLGlobalVarDocs : Documentation { let Category = DocCatFunction; - let Heading = "SYCL Global Var"; let Content = [{ This attribute can only be applied to a global variable and indicates that the variable can be used in a SYCL kernel. Access to the global variable is expected to From 5c0f8aff13d060c329afc0298666d685b21fd2e6 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Mon, 17 May 2021 20:39:32 -0700 Subject: [PATCH 08/17] Expand documentation and tests. --- clang/include/clang/Basic/AttrDocs.td | 21 ++++++++++-- clang/test/SemaSYCL/attr-syclglobalvar.cpp | 37 ++++++++++++++++++++-- 2 files changed, 53 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 6e195f179fb4b..e63852db0e681 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3158,9 +3158,24 @@ implicitly inherit this attribute. def SYCLGlobalVarDocs : Documentation { let Category = DocCatFunction; let Content = [{ -This attribute can only be applied to a global variable and indicates that the -variable can be used in a SYCL kernel. Access to the global variable is expected to -be only from device-side. +Normally a SYCL kernel cannot access a global variable, but there are cases +where we might want a global to be allocated and accessed on a SYCL device. This +attribute is only available to SYCL device compiler (that is, when +-fsycl-is-device), and only applies to global variables. It affects semantic +checks to allow use of a marked global within a SYCL kernel. + +.. code-block:: c++ + + #ifdef __SYCL_DEVICE_ONLY__ + __attribute__((sycl_global_var)) int Var; + #endif + + void F1(cl::sycl::handler& CGH) { + CGH.parallel_for_impl([=] () { + Var = 42; // device code + }); + } + }]; } diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp index b97858b9be817..82adab7192e62 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -4,18 +4,51 @@ __attribute__((sycl_global_var)) int GlobalWithAttribute; -int GlobalNoAttribute; +__attribute__((sycl_global_var)) extern int ExternGlobalWithAttribute; + +namespace NS { + __attribute__((sycl_global_var)) int NSGlobalWithAttribute; +} + +union U { + int InstanceMember; +}; + +__attribute__((sycl_global_var)) U GlobalUnion; + +struct S { + __attribute__((sycl_global_var)) static int StaticMember; + + // expected-warning@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; +int S::StaticMember = 0; + +__attribute__((sycl_global_var)) S GlobalStruct; + +__attribute__((sycl_global_var)) static S StaticGlobal; // expected-error@+1 {{attribute takes no arguments}} __attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; +int GlobalNoAttribute; + // expected-warning@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) void F() { + __attribute__((sycl_global_var)) static int StaticLocalVar; + // expected-warning@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) int Local; cl::sycl::kernel_single_task([=] () { (void)GlobalWithAttribute; + (void)ExternGlobalWithAttribute; + (void)NS::NSGlobalWithAttribute; + (void)GlobalUnion.InstanceMember; + (void)S::StaticMember; + (void)GlobalStruct.InstanceMember; + (void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} }); -} \ No newline at end of file +} From a1ac1c5664612b96e7d91c1cfe525a33e1714747 Mon Sep 17 00:00:00 2001 From: "Mott, Jeffrey T" <59898758+jtmott-intel@users.noreply.github.com> Date: Tue, 18 May 2021 11:51:38 -0700 Subject: [PATCH 09/17] Update clang/include/clang/Basic/AttrDocs.td Tweak docs verbiage. Co-authored-by: Aaron Ballman --- clang/include/clang/Basic/AttrDocs.td | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index e63852db0e681..1ec7872ca43db 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3158,10 +3158,10 @@ implicitly inherit this attribute. def SYCLGlobalVarDocs : Documentation { let Category = DocCatFunction; let Content = [{ -Normally a SYCL kernel cannot access a global variable, but there are cases -where we might want a global to be allocated and accessed on a SYCL device. This -attribute is only available to SYCL device compiler (that is, when --fsycl-is-device), and only applies to global variables. It affects semantic +Normally, a SYCL kernel cannot access a global variable, but there are cases +where it is desirable to use a global variable allocated and accessed on a SYCL device. This +attribute is only available to a SYCL device compiler (that is, when passing +``-fsycl-is-device``) and only applies to global variables. It affects semantic checks to allow use of a marked global within a SYCL kernel. .. code-block:: c++ From a046f3908c124b2ce552741245d9a9d4440352c4 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Wed, 19 May 2021 10:23:01 -0700 Subject: [PATCH 10/17] Add anonymous union to tests --- clang/test/SemaSYCL/attr-syclglobalvar.cpp | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp index 82adab7192e62..52bb906fdef65 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -10,12 +10,6 @@ namespace NS { __attribute__((sycl_global_var)) int NSGlobalWithAttribute; } -union U { - int InstanceMember; -}; - -__attribute__((sycl_global_var)) U GlobalUnion; - struct S { __attribute__((sycl_global_var)) static int StaticMember; @@ -28,6 +22,11 @@ __attribute__((sycl_global_var)) S GlobalStruct; __attribute__((sycl_global_var)) static S StaticGlobal; +static union { + // expected-warning@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int AnonymousStaticUnionInstanceMember; +}; + // expected-error@+1 {{attribute takes no arguments}} __attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; @@ -44,11 +43,11 @@ __attribute__((sycl_global_var)) void F() { (void)GlobalWithAttribute; (void)ExternGlobalWithAttribute; (void)NS::NSGlobalWithAttribute; - (void)GlobalUnion.InstanceMember; (void)S::StaticMember; (void)GlobalStruct.InstanceMember; (void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)AnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} }); } From 6bb98b9bb55c1767c066f01b31cd0b70ae6aa446 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Thu, 20 May 2021 16:14:41 -0700 Subject: [PATCH 11/17] Disallowed local vars, and upgraded diagnostic to an error --- clang/include/clang/Basic/Attr.td | 6 +++++- clang/test/SemaSYCL/attr-syclglobalvar.cpp | 9 +++++---- 2 files changed, 10 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fa1446d827505..fe51deb5d4b65 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1197,9 +1197,13 @@ def SYCLDevice : InheritableAttr { let Documentation = [SYCLDeviceDocs]; } +def SYCLGlobalVarSubject : SubsetSubjecthasGlobalStorage() && + !S->isLocalVarDecl()}], + "global variables">; + def SYCLGlobalVar : InheritableAttr { let Spellings = [GNU<"sycl_global_var">]; - let Subjects = SubjectList<[GlobalVar]>; + let Subjects = SubjectList<[SYCLGlobalVarSubject], ErrorDiag>; let LangOpts = [SYCLIsDevice]; let Documentation = [SYCLGlobalVarDocs]; let SimpleHandler = 1; diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp index 52bb906fdef65..49ab71106fa1b 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -13,7 +13,7 @@ namespace NS { struct S { __attribute__((sycl_global_var)) static int StaticMember; - // expected-warning@+1 {{attribute only applies to global variables}} + // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) int InstanceMember; }; int S::StaticMember = 0; @@ -23,7 +23,7 @@ __attribute__((sycl_global_var)) S GlobalStruct; __attribute__((sycl_global_var)) static S StaticGlobal; static union { - // expected-warning@+1 {{attribute only applies to global variables}} + // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) int AnonymousStaticUnionInstanceMember; }; @@ -32,11 +32,12 @@ __attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; int GlobalNoAttribute; -// expected-warning@+1 {{attribute only applies to global variables}} +// expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) void F() { + // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) static int StaticLocalVar; - // expected-warning@+1 {{attribute only applies to global variables}} + // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) int Local; cl::sycl::kernel_single_task([=] () { From 47ce16a4368247eb618a6f74b018ff09130f7e92 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Thu, 27 May 2021 10:34:07 -0700 Subject: [PATCH 12/17] Restrict attribute to system headers, or it has no effect --- clang/lib/Sema/SemaExpr.cpp | 4 +- .../Inputs-isystem/attr-syclglobalvar.hpp | 52 +++++++++++++++++++ .../attr-syclglobalvar.cpp | 17 +++--- 3 files changed, 64 insertions(+), 9 deletions(-) create mode 100644 clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp rename clang/test/SemaSYCL/{ => attr-syclglobalvar}/attr-syclglobalvar.cpp (68%) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 7603dcee28278..4e0d71e5e6bf9 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -226,7 +226,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // Non-const globals are allowed for SYCL explicit SIMD or with the // SYCLGlobalVar attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage() && !VD->hasAttr()) + VD->hasGlobalStorage() && + !(VD->hasAttr() && + getSourceManager().isInSystemHeader(*Locs.begin()))) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp b/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp new file mode 100644 index 0000000000000..7d87eeb2705e9 --- /dev/null +++ b/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp @@ -0,0 +1,52 @@ +#include "../../Inputs/sycl.hpp" + +__attribute__((sycl_global_var)) int HppGlobalWithAttribute; + +__attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute; + +namespace NS { + __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; +} + +struct HppS { + __attribute__((sycl_global_var)) static int StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; +int HppS::StaticMember = 0; + +__attribute__((sycl_global_var)) HppS HppGlobalStruct; + +__attribute__((sycl_global_var)) static HppS HppStaticGlobal; + +static union { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; +}; + +// expected-error@+1 {{attribute takes no arguments}} +__attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg; + +int HppGlobalNoAttribute; + +// expected-error@+1 {{attribute only applies to global variables}} +__attribute__((sycl_global_var)) void HppF() { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; + (void)HppExternGlobalWithAttribute; + (void)NS::HppNSGlobalWithAttribute; + (void)HppS::StaticMember; + (void)HppGlobalStruct.InstanceMember; + (void)HppStaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)HppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../../Inputs/sycl.hpp:* {{called by}} + }); +} diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp similarity index 68% rename from clang/test/SemaSYCL/attr-syclglobalvar.cpp rename to clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp index 49ab71106fa1b..963467e3b4a7b 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp @@ -1,6 +1,7 @@ -// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only -isystem %S %s -#include "Inputs/sycl.hpp" +#include +#include "../Inputs/sycl.hpp" __attribute__((sycl_global_var)) int GlobalWithAttribute; @@ -41,14 +42,14 @@ __attribute__((sycl_global_var)) void F() { __attribute__((sycl_global_var)) int Local; cl::sycl::kernel_single_task([=] () { - (void)GlobalWithAttribute; - (void)ExternGlobalWithAttribute; - (void)NS::NSGlobalWithAttribute; - (void)S::StaticMember; - (void)GlobalStruct.InstanceMember; + (void)GlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)ExternGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)NS::NSGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)S::StaticMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)GlobalStruct.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} (void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} (void)AnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../Inputs/sycl.hpp:* {{called by}} }); } From a349032bb9593f2e9b31e3b414385152395e563d Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Wed, 2 Jun 2021 00:33:35 -0700 Subject: [PATCH 13/17] Changed to non-simple handler, added templates to test --- clang/include/clang/Basic/Attr.td | 7 +-- clang/lib/Sema/SemaDeclAttr.cpp | 16 +++++ clang/lib/Sema/SemaExpr.cpp | 8 +-- .../Inputs-isystem/attr-syclglobalvar.hpp | 52 ++++++++-------- .../attr-syclglobalvar/attr-syclglobalvar.cpp | 59 +++++++++++-------- 5 files changed, 85 insertions(+), 57 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fe51deb5d4b65..5cac0a2dbfe5f 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1197,16 +1197,11 @@ def SYCLDevice : InheritableAttr { let Documentation = [SYCLDeviceDocs]; } -def SYCLGlobalVarSubject : SubsetSubjecthasGlobalStorage() && - !S->isLocalVarDecl()}], - "global variables">; - def SYCLGlobalVar : InheritableAttr { let Spellings = [GNU<"sycl_global_var">]; - let Subjects = SubjectList<[SYCLGlobalVarSubject], ErrorDiag>; + let Subjects = SubjectList<[GlobalVar], ErrorDiag>; let LangOpts = [SYCLIsDevice]; let Documentation = [SYCLGlobalVarDocs]; - let SimpleHandler = 1; } def SYCLKernel : InheritableAttr { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 548fa364d08d9..1e440c3b6beaa 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5077,6 +5077,19 @@ static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D, handleSimpleAttribute(S, D, AL); } +static void handleSYCLGlobalVarAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + if (!S.Context.getSourceManager().isInSystemHeader(D->getLocation())) { + S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) << AL << "system header global variables"; + } + + const auto *VD = cast(D); + if (VD && VD->isLocalVarDecl()) { + S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) << AL << "global variables"; + } + + handleSimpleAttribute(S, D, AL); +} + static void handleSYCLRegisterNumAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (!AL.checkExactlyNumArgs(S, 1)) return; @@ -9171,6 +9184,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLDeviceIndirectlyCallable: handleSYCLDeviceIndirectlyCallableAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLGlobalVar: + handleSYCLGlobalVarAttr(S, D, AL); + break; case ParsedAttr::AT_SYCLRegisterNum: handleSYCLRegisterNumAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 4e0d71e5e6bf9..71f3872e51bee 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -220,15 +220,15 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, ExprEvalContexts.empty() || (!isUnevaluatedContext() && !isConstantEvaluated()); bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD); - if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static) + if (IsRuntimeEvaluated && !IsConst && + VD->getStorageClass() == SC_Static && + !VD->hasAttr()) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; // Non-const globals are allowed for SYCL explicit SIMD or with the // SYCLGlobalVar attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage() && - !(VD->hasAttr() && - getSourceManager().isInSystemHeader(*Locs.begin()))) + VD->hasGlobalStorage() && !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/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp b/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp index 7d87eeb2705e9..5a9080cc73d45 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp @@ -5,14 +5,14 @@ __attribute__((sycl_global_var)) int HppGlobalWithAttribute; __attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute; namespace NS { - __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; + __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; } struct HppS { - __attribute__((sycl_global_var)) static int StaticMember; + __attribute__((sycl_global_var)) static int StaticMember; - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int InstanceMember; + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; }; int HppS::StaticMember = 0; @@ -21,32 +21,38 @@ __attribute__((sycl_global_var)) HppS HppGlobalStruct; __attribute__((sycl_global_var)) static HppS HppStaticGlobal; static union { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; }; // expected-error@+1 {{attribute takes no arguments}} __attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg; +template struct HppStructTemplate { + __attribute__((sycl_global_var)) static T StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; + int HppGlobalNoAttribute; // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) void HppF() { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) static int StaticLocalVar; - - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int Local; - - cl::sycl::kernel_single_task([=] () { - (void)HppGlobalWithAttribute; - (void)HppExternGlobalWithAttribute; - (void)NS::HppNSGlobalWithAttribute; - (void)HppS::StaticMember; - (void)HppGlobalStruct.InstanceMember; - (void)HppStaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)HppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../../Inputs/sycl.hpp:* {{called by}} - }); + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; + (void)HppExternGlobalWithAttribute; + (void)NS::HppNSGlobalWithAttribute; + (void)HppS::StaticMember; + (void)HppGlobalStruct.InstanceMember; + (void)HppStaticGlobal.InstanceMember; + (void)HppStructTemplate::StaticMember; + (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../../Inputs/sycl.hpp:* {{called by}} + }); } diff --git a/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp index 963467e3b4a7b..38b82e90958e4 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp @@ -3,53 +3,64 @@ #include #include "../Inputs/sycl.hpp" +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} __attribute__((sycl_global_var)) int GlobalWithAttribute; +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} __attribute__((sycl_global_var)) extern int ExternGlobalWithAttribute; namespace NS { - __attribute__((sycl_global_var)) int NSGlobalWithAttribute; + // expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} + __attribute__((sycl_global_var)) int NSGlobalWithAttribute; } struct S { - __attribute__((sycl_global_var)) static int StaticMember; + // expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} + __attribute__((sycl_global_var)) static int StaticMember; - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int InstanceMember; + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; }; int S::StaticMember = 0; +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} __attribute__((sycl_global_var)) S GlobalStruct; +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} __attribute__((sycl_global_var)) static S StaticGlobal; static union { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int AnonymousStaticUnionInstanceMember; + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int AnonymousStaticUnionInstanceMember; }; // expected-error@+1 {{attribute takes no arguments}} __attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} +__attribute__((sycl_global_var)) HppStructTemplate GlobalTemplateStructWithAttribute; +HppStructTemplate GlobalTemplateStructNoAttribute; + int GlobalNoAttribute; -// expected-error@+1 {{attribute only applies to global variables}} +// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} __attribute__((sycl_global_var)) void F() { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) static int StaticLocalVar; - - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int Local; - - cl::sycl::kernel_single_task([=] () { - (void)GlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)ExternGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)NS::NSGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)S::StaticMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)GlobalStruct.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)AnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} - (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../Inputs/sycl.hpp:* {{called by}} - }); + // expected-error@+2 {{'sycl_global_var' attribute only applies to system header global variables}} + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; + (void)HppExternGlobalWithAttribute; + (void)NS::HppNSGlobalWithAttribute; + (void)HppS::StaticMember; + (void)HppGlobalStruct.InstanceMember; + (void)HppStaticGlobal.InstanceMember; + (void)HppStructTemplate::StaticMember; + (void)GlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../Inputs/sycl.hpp:* {{called by}} + }); } From 79d4ed73cfadca6812321962d0f71cae717da46f Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Wed, 2 Jun 2021 00:47:57 -0700 Subject: [PATCH 14/17] clang format --- clang/lib/Sema/SemaDeclAttr.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 1e440c3b6beaa..68cacb7d28e70 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5079,12 +5079,14 @@ static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D, static void handleSYCLGlobalVarAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (!S.Context.getSourceManager().isInSystemHeader(D->getLocation())) { - S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) << AL << "system header global variables"; + S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) + << AL << "system header global variables"; } const auto *VD = cast(D); if (VD && VD->isLocalVarDecl()) { - S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) << AL << "global variables"; + S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) + << AL << "global variables"; } handleSimpleAttribute(S, D, AL); From e9b7b2a17eaef9683413eecaee51c8fe60d0016f Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Wed, 2 Jun 2021 20:16:59 -0700 Subject: [PATCH 15/17] Simulated system include, and added macro test --- .../attr-syclglobalvar.cpp | 76 ++++++++++++++++++- .../Inputs-isystem/attr-syclglobalvar.hpp | 58 -------------- 2 files changed, 72 insertions(+), 62 deletions(-) rename clang/test/SemaSYCL/{attr-syclglobalvar => }/attr-syclglobalvar.cpp (51%) delete mode 100644 clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp diff --git a/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp similarity index 51% rename from clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp rename to clang/test/SemaSYCL/attr-syclglobalvar.cpp index 38b82e90958e4..52acd309cc3c1 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -1,7 +1,72 @@ -// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only -isystem %S %s +// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s -#include -#include "../Inputs/sycl.hpp" +#include "Inputs/sycl.hpp" + +# 5 "header.hpp" 1 3 // Simulate a system #include to enter new file named header.hpp at line 5 + +#define SYCLGLOBALVAR_ATTR_MACRO __attribute__((sycl_global_var)) + +__attribute__((sycl_global_var)) int HppGlobalWithAttribute; + +__attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute; + +namespace NS { + __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; +} + +struct HppS { + __attribute__((sycl_global_var)) static int StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; +int HppS::StaticMember = 0; + +__attribute__((sycl_global_var)) HppS HppGlobalStruct; + +__attribute__((sycl_global_var)) static HppS HppStaticGlobal; + +static union { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; +}; + +// expected-error@+1 {{attribute takes no arguments}} +__attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg; + +template struct HppStructTemplate { + __attribute__((sycl_global_var)) static T StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; + +SYCLGLOBALVAR_ATTR_MACRO int HppGlobalWithAttrMacro; + +int HppGlobalNoAttribute; + +// expected-error@+1 {{attribute only applies to global variables}} +__attribute__((sycl_global_var)) void HppF() { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; + (void)HppExternGlobalWithAttribute; + (void)NS::HppNSGlobalWithAttribute; + (void)HppS::StaticMember; + (void)HppGlobalStruct.InstanceMember; + (void)HppStaticGlobal.InstanceMember; + (void)HppStructTemplate::StaticMember; + (void)HppGlobalWithAttrMacro; + (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + }); +} + +# 69 "header.hpp" 2 // Return from the simulated #include (with the last line number of the "header.hpp" file) // expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} __attribute__((sycl_global_var)) int GlobalWithAttribute; @@ -41,6 +106,9 @@ __attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; __attribute__((sycl_global_var)) HppStructTemplate GlobalTemplateStructWithAttribute; HppStructTemplate GlobalTemplateStructNoAttribute; +// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} +SYCLGLOBALVAR_ATTR_MACRO int GlobalWithAttrMacro; + int GlobalNoAttribute; // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} @@ -61,6 +129,6 @@ __attribute__((sycl_global_var)) void F() { (void)HppStaticGlobal.InstanceMember; (void)HppStructTemplate::StaticMember; (void)GlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} - (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../Inputs/sycl.hpp:* {{called by}} + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} }); } diff --git a/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp b/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp deleted file mode 100644 index 5a9080cc73d45..0000000000000 --- a/clang/test/SemaSYCL/attr-syclglobalvar/Inputs-isystem/attr-syclglobalvar.hpp +++ /dev/null @@ -1,58 +0,0 @@ -#include "../../Inputs/sycl.hpp" - -__attribute__((sycl_global_var)) int HppGlobalWithAttribute; - -__attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute; - -namespace NS { - __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; -} - -struct HppS { - __attribute__((sycl_global_var)) static int StaticMember; - - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int InstanceMember; -}; -int HppS::StaticMember = 0; - -__attribute__((sycl_global_var)) HppS HppGlobalStruct; - -__attribute__((sycl_global_var)) static HppS HppStaticGlobal; - -static union { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; -}; - -// expected-error@+1 {{attribute takes no arguments}} -__attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg; - -template struct HppStructTemplate { - __attribute__((sycl_global_var)) static T StaticMember; - - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int InstanceMember; -}; - -int HppGlobalNoAttribute; - -// expected-error@+1 {{attribute only applies to global variables}} -__attribute__((sycl_global_var)) void HppF() { - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) static int StaticLocalVar; - - // expected-error@+1 {{attribute only applies to global variables}} - __attribute__((sycl_global_var)) int Local; - - cl::sycl::kernel_single_task([=] () { - (void)HppGlobalWithAttribute; - (void)HppExternGlobalWithAttribute; - (void)NS::HppNSGlobalWithAttribute; - (void)HppS::StaticMember; - (void)HppGlobalStruct.InstanceMember; - (void)HppStaticGlobal.InstanceMember; - (void)HppStructTemplate::StaticMember; - (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@../../Inputs/sycl.hpp:* {{called by}} - }); -} From 7219bffcba59ce2cd49b4d364e17620f9b06bc79 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Thu, 3 Jun 2021 22:44:39 -0700 Subject: [PATCH 16/17] Use SubsetSubject, return after error, no documentation, add system header diagnostic message, add parameter test --- clang/include/clang/Basic/Attr.td | 10 +- clang/include/clang/Basic/AttrDocs.td | 24 ----- .../clang/Basic/DiagnosticSemaKinds.td | 2 + clang/lib/Sema/SemaDeclAttr.cpp | 10 +- clang/test/SemaSYCL/attr-syclglobalvar.cpp | 96 +++++++++++-------- 5 files changed, 68 insertions(+), 74 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 5cac0a2dbfe5f..416cdabf5de8d 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1197,11 +1197,17 @@ def SYCLDevice : InheritableAttr { let Documentation = [SYCLDeviceDocs]; } +def GlobalStorageNonLocalVar : SubsetSubjecthasGlobalStorage() && + !S->isLocalVarDeclOrParm()}], + "global variables">; + def SYCLGlobalVar : InheritableAttr { let Spellings = [GNU<"sycl_global_var">]; - let Subjects = SubjectList<[GlobalVar], ErrorDiag>; + let Subjects = SubjectList<[GlobalStorageNonLocalVar], ErrorDiag>; let LangOpts = [SYCLIsDevice]; - let Documentation = [SYCLGlobalVarDocs]; + // Only used internally by the SYCL implementation + let Documentation = [Undocumented]; } def SYCLKernel : InheritableAttr { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 1ec7872ca43db..6e033f9e083b2 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3155,30 +3155,6 @@ implicitly inherit this attribute. }]; } -def SYCLGlobalVarDocs : Documentation { - let Category = DocCatFunction; - let Content = [{ -Normally, a SYCL kernel cannot access a global variable, but there are cases -where it is desirable to use a global variable allocated and accessed on a SYCL device. This -attribute is only available to a SYCL device compiler (that is, when passing -``-fsycl-is-device``) and only applies to global variables. It affects semantic -checks to allow use of a marked global within a SYCL kernel. - -.. code-block:: c++ - - #ifdef __SYCL_DEVICE_ONLY__ - __attribute__((sycl_global_var)) int Var; - #endif - - void F1(cl::sycl::handler& CGH) { - CGH.parallel_for_impl([=] () { - Var = 42; // device code - }); - } - - }]; -} - def RISCVInterruptDocs : Documentation { let Category = DocCatFunction; let Heading = "interrupt (RISCV)"; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9d5c936c316d9..d24c88a690141 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3402,6 +3402,8 @@ def warn_attribute_wrong_decl_type_str : Warning< "%0 attribute only applies to %1">, InGroup; def err_attribute_wrong_decl_type_str : Error< warn_attribute_wrong_decl_type_str.Text>; +def err_attribute_only_system_header : Error< + "%0 attribute only supported within a system header">; def warn_attribute_wrong_decl_type : Warning< "%0 attribute only applies to %select{" "functions" diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 68cacb7d28e70..b8fc54f53da5d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5079,14 +5079,8 @@ static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D, static void handleSYCLGlobalVarAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (!S.Context.getSourceManager().isInSystemHeader(D->getLocation())) { - S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) - << AL << "system header global variables"; - } - - const auto *VD = cast(D); - if (VD && VD->isLocalVarDecl()) { - S.Diag(AL.getLoc(), diag::err_attribute_wrong_decl_type_str) - << AL << "global variables"; + S.Diag(AL.getLoc(), diag::err_attribute_only_system_header) << AL; + return; } handleSimpleAttribute(S, D, AL); diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp index 52acd309cc3c1..968e556d5a72d 100644 --- a/clang/test/SemaSYCL/attr-syclglobalvar.cpp +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -46,7 +46,10 @@ SYCLGLOBALVAR_ATTR_MACRO int HppGlobalWithAttrMacro; int HppGlobalNoAttribute; // expected-error@+1 {{attribute only applies to global variables}} -__attribute__((sycl_global_var)) void HppF() { +__attribute__((sycl_global_var)) void HppF( + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Param +) { // expected-error@+1 {{attribute only applies to global variables}} __attribute__((sycl_global_var)) static int StaticLocalVar; @@ -54,66 +57,70 @@ __attribute__((sycl_global_var)) void HppF() { __attribute__((sycl_global_var)) int Local; cl::sycl::kernel_single_task([=] () { - (void)HppGlobalWithAttribute; - (void)HppExternGlobalWithAttribute; - (void)NS::HppNSGlobalWithAttribute; - (void)HppS::StaticMember; - (void)HppGlobalStruct.InstanceMember; - (void)HppStaticGlobal.InstanceMember; - (void)HppStructTemplate::StaticMember; - (void)HppGlobalWithAttrMacro; + (void)HppGlobalWithAttribute; // ok + (void)HppExternGlobalWithAttribute; // ok + (void)NS::HppNSGlobalWithAttribute; // ok + (void)HppS::StaticMember; // ok + (void)HppGlobalStruct.InstanceMember; // ok + (void)HppStaticGlobal.InstanceMember; // ok + (void)HppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)HppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStructTemplate::StaticMember; // ok + (void)HppGlobalWithAttrMacro; // ok (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} }); } -# 69 "header.hpp" 2 // Return from the simulated #include (with the last line number of the "header.hpp" file) +# 74 "header.hpp" 2 // Return from the simulated #include (with the last line number of the "header.hpp" file) -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -__attribute__((sycl_global_var)) int GlobalWithAttribute; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) int CppGlobalWithAttribute; -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -__attribute__((sycl_global_var)) extern int ExternGlobalWithAttribute; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) extern int CppExternGlobalWithAttribute; namespace NS { - // expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} - __attribute__((sycl_global_var)) int NSGlobalWithAttribute; + // expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} + __attribute__((sycl_global_var)) int CppNSGlobalWithAttribute; } -struct S { - // expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} +struct CppS { + // expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} __attribute__((sycl_global_var)) static int StaticMember; // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} __attribute__((sycl_global_var)) int InstanceMember; }; -int S::StaticMember = 0; +int CppS::StaticMember = 0; -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -__attribute__((sycl_global_var)) S GlobalStruct; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) CppS CppGlobalStruct; -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -__attribute__((sycl_global_var)) static S StaticGlobal; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) static CppS CppStaticGlobal; static union { // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} - __attribute__((sycl_global_var)) int AnonymousStaticUnionInstanceMember; + __attribute__((sycl_global_var)) int CppAnonymousStaticUnionInstanceMember; }; // expected-error@+1 {{attribute takes no arguments}} -__attribute__((sycl_global_var(42))) int GlobalWithAttributeArg; +__attribute__((sycl_global_var(42))) int CppGlobalWithAttributeArg; -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -__attribute__((sycl_global_var)) HppStructTemplate GlobalTemplateStructWithAttribute; -HppStructTemplate GlobalTemplateStructNoAttribute; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) HppStructTemplate CppGlobalTemplateStructWithAttribute; +HppStructTemplate CppGlobalTemplateStructNoAttribute; -// expected-error@+1 {{'sycl_global_var' attribute only applies to system header global variables}} -SYCLGLOBALVAR_ATTR_MACRO int GlobalWithAttrMacro; +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +SYCLGLOBALVAR_ATTR_MACRO int CppGlobalWithAttrMacro; int GlobalNoAttribute; // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} -__attribute__((sycl_global_var)) void F() { - // expected-error@+2 {{'sycl_global_var' attribute only applies to system header global variables}} +__attribute__((sycl_global_var)) void F( + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Param +) { // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} __attribute__((sycl_global_var)) static int StaticLocalVar; @@ -121,14 +128,23 @@ __attribute__((sycl_global_var)) void F() { __attribute__((sycl_global_var)) int Local; cl::sycl::kernel_single_task([=] () { - (void)HppGlobalWithAttribute; - (void)HppExternGlobalWithAttribute; - (void)NS::HppNSGlobalWithAttribute; - (void)HppS::StaticMember; - (void)HppGlobalStruct.InstanceMember; - (void)HppStaticGlobal.InstanceMember; - (void)HppStructTemplate::StaticMember; - (void)GlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppGlobalWithAttribute; // ok + (void)CppGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppExternGlobalWithAttribute; // ok + (void)CppExternGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)NS::HppNSGlobalWithAttribute; // ok + (void)NS::CppNSGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppS::StaticMember; // ok + (void)CppS::StaticMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppGlobalStruct.InstanceMember; // ok + (void)CppGlobalStruct.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStaticGlobal.InstanceMember; // ok + (void)CppStaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)CppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)CppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStructTemplate::StaticMember; // ok + (void)CppGlobalTemplateStructWithAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)CppGlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} }); } From 337383030bef01a626f9ba3499cdbc9196b673e6 Mon Sep 17 00:00:00 2001 From: Jeffrey T Mott Date: Fri, 4 Jun 2021 14:18:58 -0700 Subject: [PATCH 17/17] I guess not documented means not on the supported list --- clang/test/Misc/pragma-attribute-supported-attributes-list.test | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index c3c38bd645b83..80f2ac881ea0b 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -152,7 +152,6 @@ // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) -// CHECK-NEXT: SYCLGlobalVar (SubjectMatchRule_variable_is_global) // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function)