From 38dc2f5e17275637a6c2f72328f9cd3ef4e1f76f Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 21 Feb 2024 08:29:31 -0800 Subject: [PATCH 01/16] [FPGA][SYCL] Allow memory attributes on non-const device_global variables Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 5 +++-- clang/test/SemaSYCL/intel-fpga-local.cpp | 9 +++++++++ 2 files changed, 12 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 4a50fb51a4a63..2336337aaa391 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2648,10 +2648,11 @@ def SYCLIntelConstVar : SubsetSubjectgetKind() != Decl::ImplicitParam && S->getKind() != Decl::ParmVar && S->getKind() != Decl::NonTypeTemplateParm && - (S->getType().isConstQualified() || + ((!S->hasAttr() || + S->getType().isConstQualified()) || S->getType().getAddressSpace() == LangAS::opencl_constant)}], - "constant variables">; + "constant variables, non-const device_global variables">; def SYCLIntelLocalStaticAgentMemVar : SubsetSubjectgetKind() != Decl::ImplicitParam && diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index d6b48db59bd9e..be55c0c01a470 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -587,3 +587,12 @@ int main() { //expected-note@+1{{conflicting attribute is here}} [[intel::bankwidth(8)]] extern const int var_bankwidth_2; [[intel::fpga_register]] const int var_bankwidth_2 =0; + +using namespace sycl::ext::oneapi; // for properties; +[[intel::numbanks(2)]] /*const*/ device_global const_glob; // OK + +[[intel::numbanks(8)]] const device_global const_glob2; // OK + +struct Foo { + [[intel::numbanks(2)]] /*const*/ device_global const_glob3; // OK +}; From db826f5c6f2f3e221e10a72b29f3b51a7f86b15b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 25 Feb 2024 22:31:39 -0800 Subject: [PATCH 02/16] Add support for device_global variables on FPGA attributes --- clang/include/clang/Basic/Attr.td | 8 +--- .../clang/Basic/DiagnosticSemaKinds.td | 4 ++ clang/lib/Sema/SemaDeclAttr.cpp | 10 +++++ clang/test/CodeGenSYCL/device_global.cpp | 7 +++ .../SemaSYCL/intel-fpga-device-global.cpp | 44 +++++++++++++++++++ clang/test/SemaSYCL/intel-fpga-local.cpp | 9 ---- 6 files changed, 67 insertions(+), 15 deletions(-) create mode 100644 clang/test/SemaSYCL/intel-fpga-device-global.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 2336337aaa391..46ba00475c5e1 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2648,11 +2648,10 @@ def SYCLIntelConstVar : SubsetSubjectgetKind() != Decl::ImplicitParam && S->getKind() != Decl::ParmVar && S->getKind() != Decl::NonTypeTemplateParm && - ((!S->hasAttr() || - S->getType().isConstQualified()) || + (S->getType().isConstQualified() || S->getType().getAddressSpace() == LangAS::opencl_constant)}], - "constant variables, non-const device_global variables">; + "constant variables">; def SYCLIntelLocalStaticAgentMemVar : SubsetSubjectgetKind() != Decl::ImplicitParam && @@ -2731,9 +2730,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelBankWidth]>; def SYCLIntelNumBanks : InheritableAttr { let Spellings = [CXX11<"intel", "numbanks">]; let Args = [ExprArgument<"Value">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelNumBanksAttrDocs]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 108ac8272057d..7c2404187158a 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12215,6 +12215,10 @@ def err_sycl_attribute_internal_decl "in an anonymous namespace">; def err_sycl_attribute_not_device_global : Error<"%0 attribute can only be applied to 'device_global' variables">; +def err_fpga_attribute_incorrrect_variable + : Error<"%0 attribute only applies to constant variables, local variables, " + "static variables, agent memory arguments, non-static data " + "members, and device_global variables">; def err_sycl_compiletime_property_duplication : Error< "can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 68949acfcd5fb..7cafa4f0b646b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7745,6 +7745,16 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, } } + if (auto *VD = dyn_cast(D)) { + if (!(VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || + (isTypeDecoratedWithDeclAttribute(VD->getType()) + || VD->getType().isConstQualified())))){ + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrrect_variable) << CI; + } + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 88b3e8f607a96..9353e9dc6378f 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -11,6 +11,9 @@ using namespace sycl; queue q; device_global A; + +[[intel::numbanks(2)]] /*const*/ device_global A1; + #ifdef SYCL_EXTERNAL SYCL_EXTERNAL device_global AExt; #endif @@ -20,8 +23,10 @@ struct Foo { static device_global C; }; device_global Foo::C; + // CHECK-RDC: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]] // CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]] +// CHECK: @A1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A1_ATTRS:[0-9]+]] // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]] // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]] @@ -70,6 +75,7 @@ void foo() { q.submit([&](handler &h) { h.single_task([=]() { (void)A; + (void)A1; (void)B; (void)Foo::C; (void)same_name; @@ -103,6 +109,7 @@ void bar() { // CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" } // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } +// CHECK: attributes #[[A1_ATTRS]] = { "sycl-unique-id"="_Z2A1" } // 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" } diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp new file mode 100644 index 0000000000000..cd4c36bfc5592 --- /dev/null +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s +#include "Inputs/sycl.hpp" + +// Tests that [[intel::numbanks()]] can be applied to device_global variables. +using namespace sycl::ext::oneapi; + +[[intel::numbanks(2)]] device_global dev_glob; // OK +[[intel::numbanks(4)]] static device_global static_dev_glob; // OK + +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +[[intel::numbanks(2)]] int K; + +struct bar { + [[intel::numbanks(2)]] /*const*/ device_global const_glob3; // OK + [[intel::numbanks(2)]] const device_global const_glob4; // OK +}; + +void foo() { + [[intel::numbanks(2)]] int A1; // OK +} + +struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { +}; + +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +[[intel::numbanks(2)]] GlobAllowedVarOnly GAVO; + +[[intel::numbanks(4)]] /*const*/ device_global Good; +[[intel::numbanks(4)]] extern device_global Bad; + +int main() { + sycl::kernel_single_task([=]() { + Good.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad.get(); + (void)GAVO; + }); + return 0; +} + +//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +[[intel::numbanks(2)]] +__attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; + diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index be55c0c01a470..d6b48db59bd9e 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -587,12 +587,3 @@ int main() { //expected-note@+1{{conflicting attribute is here}} [[intel::bankwidth(8)]] extern const int var_bankwidth_2; [[intel::fpga_register]] const int var_bankwidth_2 =0; - -using namespace sycl::ext::oneapi; // for properties; -[[intel::numbanks(2)]] /*const*/ device_global const_glob; // OK - -[[intel::numbanks(8)]] const device_global const_glob2; // OK - -struct Foo { - [[intel::numbanks(2)]] /*const*/ device_global const_glob3; // OK -}; From 5e573c19c8c631fd1d92e29ac46f7a30ec592cb2 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 25 Feb 2024 22:38:51 -0800 Subject: [PATCH 03/16] Update patch --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/test/SemaSYCL/intel-fpga-device-global.cpp | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 7c2404187158a..b48b08db0c1b7 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12218,7 +12218,7 @@ def err_sycl_attribute_not_device_global def err_fpga_attribute_incorrrect_variable : Error<"%0 attribute only applies to constant variables, local variables, " "static variables, agent memory arguments, non-static data " - "members, and device_global variables">; + "members, and non-constant device_global variables">; def err_sycl_compiletime_property_duplication : Error< "can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index cd4c36bfc5592..ee66d634d3807 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -7,7 +7,7 @@ using namespace sycl::ext::oneapi; [[intel::numbanks(2)]] device_global dev_glob; // OK [[intel::numbanks(4)]] static device_global static_dev_glob; // OK -// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} [[intel::numbanks(2)]] int K; struct bar { @@ -22,7 +22,7 @@ void foo() { struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; -// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} [[intel::numbanks(2)]] GlobAllowedVarOnly GAVO; [[intel::numbanks(4)]] /*const*/ device_global Good; @@ -38,7 +38,7 @@ int main() { return 0; } -//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and device_global variables}} +//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} [[intel::numbanks(2)]] __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; From e45c0b52b39fe32ad109a7696e3c384aed840904 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 26 Feb 2024 21:35:41 -0800 Subject: [PATCH 04/16] Address Review Comments --- .../clang/Basic/DiagnosticSemaKinds.td | 4 +-- clang/lib/Sema/SemaDeclAttr.cpp | 14 ++++++--- clang/test/CodeGenSYCL/device_global.cpp | 2 +- .../SemaSYCL/intel-fpga-device-global.cpp | 29 ++++++++++++++----- 4 files changed, 35 insertions(+), 14 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index b48b08db0c1b7..7777b7e1906e0 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12215,10 +12215,10 @@ def err_sycl_attribute_internal_decl "in an anonymous namespace">; def err_sycl_attribute_not_device_global : Error<"%0 attribute can only be applied to 'device_global' variables">; -def err_fpga_attribute_incorrrect_variable +def err_fpga_attribute_incorrect_variable : Error<"%0 attribute only applies to constant variables, local variables, " "static variables, agent memory arguments, non-static data " - "members, and non-constant device_global variables">; + "members and device_global variables">; def err_sycl_compiletime_property_duplication : Error< "can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 7cafa4f0b646b..9e34d8e560d82 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7745,13 +7745,19 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, } } - if (auto *VD = dyn_cast(D)) { + // Check attribute only applies to constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (const auto *VD = dyn_cast(D)){ if (!(VD->getKind() != Decl::ImplicitParam && VD->getKind() != Decl::NonTypeTemplateParm && ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || - (isTypeDecoratedWithDeclAttribute(VD->getType()) - || VD->getType().isConstQualified())))){ - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrrect_variable) << CI; + (VD->getKind() != Decl::ParmVar && + (isTypeDecoratedWithDeclAttribute(VD->getType()) + || VD->getType().isConstQualified() + || VD->getType().getAddressSpace() == + LangAS::opencl_constant))))){ + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; } } diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 9353e9dc6378f..790d4936dc44c 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -12,7 +12,7 @@ queue q; device_global A; -[[intel::numbanks(2)]] /*const*/ device_global A1; +[[intel::numbanks(2)]] device_global A1; #ifdef SYCL_EXTERNAL SYCL_EXTERNAL device_global AExt; diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index ee66d634d3807..296fb5710d3aa 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -1,31 +1,44 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Tests that [[intel::numbanks()]] can be applied to device_global variables. +// Tests that [[intel::numbanks()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. + using namespace sycl::ext::oneapi; -[[intel::numbanks(2)]] device_global dev_glob; // OK +[[intel::numbanks(2)]] device_global nonconst_dev_glob; // OK +[[intel::numbanks(8)]] const device_global constdev_glob; // OK [[intel::numbanks(4)]] static device_global static_dev_glob; // OK -// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; struct bar { - [[intel::numbanks(2)]] /*const*/ device_global const_glob3; // OK + [[intel::numbanks(2)]] device_global nonconst_glob3; // OK [[intel::numbanks(2)]] const device_global const_glob4; // OK + [[intel::numbanks(8)]] unsigned int numbanks[64]; }; void foo() { [[intel::numbanks(2)]] int A1; // OK + [[intel::numbanks(4)]] static unsigned int ext_five[64]; // OK +} + +void attr_on_const_no_error() +{ + //expected-no-error@+1 + [[intel::numbanks(16)]] const int const_var[64] = {0, 1}; } +//expected-no-error@+1 +void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} + struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; -// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] GlobAllowedVarOnly GAVO; -[[intel::numbanks(4)]] /*const*/ device_global Good; +[[intel::numbanks(4)]] device_global Good; [[intel::numbanks(4)]] extern device_global Bad; int main() { @@ -38,7 +51,9 @@ int main() { return 0; } -//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}} +//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; +[[intel::numbanks(8)]] +__attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; // OK From 6eb22a03daf5a0074c36596688149d3eb44df379 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 26 Feb 2024 21:39:49 -0800 Subject: [PATCH 05/16] Update tests --- clang/test/SemaSYCL/intel-fpga-device-global.cpp | 14 +++++--------- 1 file changed, 5 insertions(+), 9 deletions(-) diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 296fb5710d3aa..fbba0bd806e32 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -5,31 +5,27 @@ using namespace sycl::ext::oneapi; -[[intel::numbanks(2)]] device_global nonconst_dev_glob; // OK -[[intel::numbanks(8)]] const device_global constdev_glob; // OK [[intel::numbanks(4)]] static device_global static_dev_glob; // OK // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; struct bar { - [[intel::numbanks(2)]] device_global nonconst_glob3; // OK - [[intel::numbanks(2)]] const device_global const_glob4; // OK + [[intel::numbanks(2)]] device_global nonconst_glob3; + [[intel::numbanks(2)]] const device_global const_glob4; [[intel::numbanks(8)]] unsigned int numbanks[64]; }; void foo() { - [[intel::numbanks(2)]] int A1; // OK - [[intel::numbanks(4)]] static unsigned int ext_five[64]; // OK + [[intel::numbanks(2)]] int A1; + [[intel::numbanks(4)]] static unsigned int ext_five[64]; } void attr_on_const_no_error() { - //expected-no-error@+1 [[intel::numbanks(16)]] const int const_var[64] = {0, 1}; } -//expected-no-error@+1 void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { @@ -56,4 +52,4 @@ int main() { __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; [[intel::numbanks(8)]] -__attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; // OK +__attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; From 21d536209cc0ea2d53c4466b5ad5d7e7149b8147 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 26 Feb 2024 23:05:33 -0800 Subject: [PATCH 06/16] Allow max_replicates attribute on device_global --- clang/include/clang/Basic/Attr.td | 3 -- clang/lib/Sema/SemaDeclAttr.cpp | 17 +++++++++ .../SemaSYCL/intel-fpga-device-global.cpp | 37 +++++++++++++++++-- 3 files changed, 50 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 46ba00475c5e1..6caaf0be025ce 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2757,9 +2757,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelMerge]>; def SYCLIntelMaxReplicates : InheritableAttr { let Spellings = [CXX11<"intel", "max_replicates">]; let Args = [ExprArgument<"Value">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelMaxReplicatesAttrDocs]; } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 9e34d8e560d82..46e107c0502a6 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7853,6 +7853,23 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, << CI << /*positive*/ 0; return; } + + // Check attribute only applies to constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (const auto *VD = dyn_cast(D)){ + if (!(VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || + (VD->getKind() != Decl::ParmVar && + (isTypeDecoratedWithDeclAttribute(VD->getType()) + || VD->getType().isConstQualified() + || VD->getType().getAddressSpace() == + LangAS::opencl_constant))))){ + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; + } + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index fbba0bd806e32..5b98674a6a176 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -1,32 +1,43 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Tests that [[intel::numbanks()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. +// Tests that [[intel::numbanks()]], [[intel::max_replicates()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. using namespace sycl::ext::oneapi; [[intel::numbanks(4)]] static device_global static_dev_glob; // OK +[[intel::max_replicates(12)]] static device_global static_dev_glob1; // OK // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; +// expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::max_replicates(10)]] int K1; struct bar { - [[intel::numbanks(2)]] device_global nonconst_glob3; - [[intel::numbanks(2)]] const device_global const_glob4; + [[intel::numbanks(2)]] device_global nonconst_glob; + [[intel::numbanks(4)]] const device_global const_glob; [[intel::numbanks(8)]] unsigned int numbanks[64]; + + [[intel::max_replicates(2)]] device_global nonconst_glob1; + [[intel::max_replicates(4)]] const device_global const_glob1; + [[intel::max_replicates(4)]] unsigned int max_rep[64]; }; void foo() { - [[intel::numbanks(2)]] int A1; + [[intel::numbanks(2)]] int A; [[intel::numbanks(4)]] static unsigned int ext_five[64]; + [[intel::max_replicates(2)]] int A1; + [[intel::max_replicates(4)]] static unsigned int ext_five1[64]; } void attr_on_const_no_error() { [[intel::numbanks(16)]] const int const_var[64] = {0, 1}; + [[intel::max_replicates(16)]] const int const_var_max[64] = {0, 1}; } void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} +void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; @@ -34,14 +45,24 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] GlobAllowedVarOnly GAVO; +// expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::max_replicates(20)]] GlobAllowedVarOnly GAVO1; + + [[intel::numbanks(4)]] device_global Good; [[intel::numbanks(4)]] extern device_global Bad; +[[intel::max_replicates(8)]] device_global Good1; +[[intel::max_replicates(10)]] extern device_global Bad1; + int main() { sycl::kernel_single_task([=]() { Good.get(); + Good1.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad1.get(); (void)GAVO; }); return 0; @@ -51,5 +72,13 @@ int main() { [[intel::numbanks(2)]] __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; +//expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::max_replicates(20)]] +__attribute__((opencl_global)) unsigned int ocl_glob_max_p2d[64] = {1, 2, 3}; + + [[intel::numbanks(8)]] __attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; + +[[intel::max_replicates(16)]] +__attribute__((opencl_constant)) unsigned int const_var_max_rep[64] = {1, 2, 3}; From 5f3ea153f1570f767f15723ca3de462f042e6f29 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 27 Feb 2024 06:37:21 -0800 Subject: [PATCH 07/16] Remove duplicate codes --- clang/lib/Sema/SemaDeclAttr.cpp | 48 ++++++++++++++++----------------- 1 file changed, 24 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 46e107c0502a6..6aa59e2accfbc 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7443,6 +7443,26 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D, return false; } +// Checks if FPGA memory attributes apply on valid variables. +// Returns true if an error occured. +static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D, + const AttributeCommonInfo &CI) { + if (const auto *VD = dyn_cast(D)){ + if (!(VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || + (VD->getKind() != Decl::ParmVar && + (S.isTypeDecoratedWithDeclAttribute(VD->getType()) + || VD->getType().isConstQualified() + || VD->getType().getAddressSpace() == + LangAS::opencl_constant))))) { + S.Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; + return true; + } + } + return false; +} + void Sema::AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E) { @@ -7748,18 +7768,8 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute only applies to constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (const auto *VD = dyn_cast(D)){ - if (!(VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || - (VD->getKind() != Decl::ParmVar && - (isTypeDecoratedWithDeclAttribute(VD->getType()) - || VD->getType().isConstQualified() - || VD->getType().getAddressSpace() == - LangAS::opencl_constant))))){ - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; - } - } + if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + return; // Check to see if there's a duplicate attribute with different values // already applied to the declaration. @@ -7857,18 +7867,8 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute only applies to constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (const auto *VD = dyn_cast(D)){ - if (!(VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || - (VD->getKind() != Decl::ParmVar && - (isTypeDecoratedWithDeclAttribute(VD->getType()) - || VD->getType().isConstQualified() - || VD->getType().getAddressSpace() == - LangAS::opencl_constant))))){ - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; - } - } + if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + return; // Check to see if there's a duplicate attribute with different values // already applied to the declaration. From 3dad7b2c30f0c5c6bffb4fc32d17d3506f1e0c8c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 27 Feb 2024 07:23:21 -0800 Subject: [PATCH 08/16] Add support for force_pow2_depth attribute on device_global variable. --- clang/include/clang/Basic/Attr.td | 3 -- clang/lib/Sema/SemaDeclAttr.cpp | 6 ++++ clang/test/CodeGenSYCL/device_global.cpp | 16 ++++++--- .../SemaSYCL/intel-fpga-device-global.cpp | 35 ++++++++++++++++--- clang/test/SemaSYCL/intel-fpga-local.cpp | 2 +- 5 files changed, 49 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 6caaf0be025ce..97acdd506d6c3 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2803,9 +2803,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelNumBanks]>; def SYCLIntelForcePow2Depth : InheritableAttr { let Spellings = [CXX11<"intel", "force_pow2_depth">]; let Args = [ExprArgument<"Value">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelForcePow2DepthAttrDocs]; } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 6aa59e2accfbc..be1ac0160d4d1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -8113,6 +8113,12 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, return; } + // Check attribute only applies to constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + return; + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 790d4936dc44c..f8862e1f10660 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -12,7 +12,9 @@ queue q; device_global A; -[[intel::numbanks(2)]] device_global A1; +[[intel::numbanks(2)]] device_global Nonconst_glob; +[[intel::max_replicates(2)]] device_global Nonconst_glob1; +[[intel::force_pow2_depth(1)]] device_global Nonconst_glob2; #ifdef SYCL_EXTERNAL SYCL_EXTERNAL device_global AExt; @@ -26,7 +28,9 @@ device_global Foo::C; // CHECK-RDC: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]] // CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]] -// CHECK: @A1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A1_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Num_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Max_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob2 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Force_ATTRS:[0-9]+]] // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]] // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]] @@ -75,7 +79,9 @@ void foo() { q.submit([&](handler &h) { h.single_task([=]() { (void)A; - (void)A1; + (void)Nonconst_glob; + (void)Nonconst_glob1; + (void)Nonconst_glob2; (void)B; (void)Foo::C; (void)same_name; @@ -109,7 +115,9 @@ void bar() { // CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" } // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } -// CHECK: attributes #[[A1_ATTRS]] = { "sycl-unique-id"="_Z2A1" } +// CHECK: attributes #[[Non_Const_Num_ATTRS]] = { "sycl-unique-id"="_Z13Nonconst_glob" } +// CHECK: attributes #[[Non_Const_Max_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob1" } +// CHECK: attributes #[[Non_Const_Force_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob2" } // 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" } diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 5b98674a6a176..b8f582b51522f 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -1,18 +1,23 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Tests that [[intel::numbanks()]], [[intel::max_replicates()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. +// Tests that [[intel::numbanks()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. using namespace sycl::ext::oneapi; -[[intel::numbanks(4)]] static device_global static_dev_glob; // OK -[[intel::max_replicates(12)]] static device_global static_dev_glob1; // OK +[[intel::numbanks(4)]] static device_global static_dev_glob; +[[intel::max_replicates(12)]] static device_global static_dev_glob1; +[[intel::force_pow2_depth(1)]] static device_global static_dev_glob2; // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; + // expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::max_replicates(10)]] int K1; +// expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::force_pow2_depth(1)]] int K2; + struct bar { [[intel::numbanks(2)]] device_global nonconst_glob; [[intel::numbanks(4)]] const device_global const_glob; @@ -20,7 +25,11 @@ struct bar { [[intel::max_replicates(2)]] device_global nonconst_glob1; [[intel::max_replicates(4)]] const device_global const_glob1; - [[intel::max_replicates(4)]] unsigned int max_rep[64]; + [[intel::max_replicates(8)]] unsigned int max_rep[64]; + + [[intel::force_pow2_depth(0)]] device_global nonconst_glob2; + [[intel::force_pow2_depth(0)]] const device_global const_glob2; + [[intel::force_pow2_depth(1)]] unsigned int force_dep[64]; }; void foo() { @@ -28,16 +37,20 @@ void foo() { [[intel::numbanks(4)]] static unsigned int ext_five[64]; [[intel::max_replicates(2)]] int A1; [[intel::max_replicates(4)]] static unsigned int ext_five1[64]; + [[intel::force_pow2_depth(0)]] int A2; + [[intel::force_pow2_depth(1)]] static unsigned int ext_five2[64]; } void attr_on_const_no_error() { [[intel::numbanks(16)]] const int const_var[64] = {0, 1}; [[intel::max_replicates(16)]] const int const_var_max[64] = {0, 1}; + [[intel::force_pow2_depth(1)]] const int const_var_force[64] = {0, 1}; } void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} +void attr_on_func_arg2([[intel::force_pow2_depth(1)]] int pc2) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; @@ -48,6 +61,8 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { // expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::max_replicates(20)]] GlobAllowedVarOnly GAVO1; +// expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::force_pow2_depth(0)]] GlobAllowedVarOnly GAVO2; [[intel::numbanks(4)]] device_global Good; [[intel::numbanks(4)]] extern device_global Bad; @@ -55,15 +70,23 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { [[intel::max_replicates(8)]] device_global Good1; [[intel::max_replicates(10)]] extern device_global Bad1; +[[intel::force_pow2_depth(0)]] device_global Good2; +[[intel::force_pow2_depth(0)]] extern device_global Bad2; + int main() { sycl::kernel_single_task([=]() { Good.get(); Good1.get(); + Good2.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad1.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad2.get(); (void)GAVO; + (void)GAVO1; + (void)GAVO2; }); return 0; } @@ -76,9 +99,11 @@ __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; [[intel::max_replicates(20)]] __attribute__((opencl_global)) unsigned int ocl_glob_max_p2d[64] = {1, 2, 3}; - [[intel::numbanks(8)]] __attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; [[intel::max_replicates(16)]] __attribute__((opencl_constant)) unsigned int const_var_max_rep[64] = {1, 2, 3}; + +[[intel::force_pow2_depth(0)]] +__attribute__((opencl_constant)) unsigned int const_force_var[64] = {1, 2, 3}; diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index d6b48db59bd9e..753c5b727a23b 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -455,7 +455,7 @@ void attr_on_const_error() //expected-error@+1{{attribute only applies to local non-const variables and non-static data members}} void attr_on_func_arg([[intel::private_copies(8)]] int pc) {} -//expected-error@+1{{attribute only applies to constant variables, local variables, static variables, agent memory arguments, and non-static data members}} +//expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::force_pow2_depth(0)]] __attribute__((opencl_global)) unsigned int ocl_glob_force_p2d[64] = {1, 2, 3}; From 362a72fd5dfca46b49457522360475e4f793834e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 27 Feb 2024 09:16:53 -0800 Subject: [PATCH 09/16] Add support for bankwidth attribute on device_global variable. --- clang/include/clang/Basic/Attr.td | 3 -- clang/lib/Sema/SemaDeclAttr.cpp | 6 +++ clang/test/CodeGenSYCL/device_global.cpp | 4 ++ .../SemaSYCL/intel-fpga-device-global.cpp | 54 ++++++++++++++++++- 4 files changed, 63 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 97acdd506d6c3..22baac7e149fa 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2719,9 +2719,6 @@ def : MutualExclusions<[SYCLIntelDoublePump, SYCLIntelSinglePump, def SYCLIntelBankWidth : InheritableAttr { let Spellings = [CXX11<"intel", "bankwidth">]; let Args = [ExprArgument<"Value">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelBankWidthAttrDocs]; } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index be1ac0160d4d1..301b2162c5aa0 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7681,6 +7681,12 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, return; } + // Check attribute only applies to constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + return; + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index f8862e1f10660..ccca5c5b92f78 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -15,6 +15,7 @@ device_global A; [[intel::numbanks(2)]] device_global Nonconst_glob; [[intel::max_replicates(2)]] device_global Nonconst_glob1; [[intel::force_pow2_depth(1)]] device_global Nonconst_glob2; +[[intel::bankwidth(2)]] device_global Nonconst_glob3; #ifdef SYCL_EXTERNAL SYCL_EXTERNAL device_global AExt; @@ -31,6 +32,7 @@ device_global Foo::C; // CHECK: @Nonconst_glob = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Num_ATTRS:[0-9]+]] // CHECK: @Nonconst_glob1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Max_ATTRS:[0-9]+]] // CHECK: @Nonconst_glob2 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Force_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob3 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Bankw_ATTRS:[0-9]+]] // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]] // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]] @@ -82,6 +84,7 @@ void foo() { (void)Nonconst_glob; (void)Nonconst_glob1; (void)Nonconst_glob2; + (void)Nonconst_glob3; (void)B; (void)Foo::C; (void)same_name; @@ -118,6 +121,7 @@ void bar() { // CHECK: attributes #[[Non_Const_Num_ATTRS]] = { "sycl-unique-id"="_Z13Nonconst_glob" } // CHECK: attributes #[[Non_Const_Max_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob1" } // CHECK: attributes #[[Non_Const_Force_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob2" } +// CHECK: attributes #[[Non_Const_Bankw_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob3" } // 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" } diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index b8f582b51522f..55b9bff1d1e87 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -1,13 +1,14 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Tests that [[intel::numbanks()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. +// Tests that [[intel::numbanks()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]], [[intel::bankwidth()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. using namespace sycl::ext::oneapi; [[intel::numbanks(4)]] static device_global static_dev_glob; [[intel::max_replicates(12)]] static device_global static_dev_glob1; [[intel::force_pow2_depth(1)]] static device_global static_dev_glob2; +[[intel::bankwidth(4)]] static device_global static_dev_glob3; // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; @@ -18,6 +19,9 @@ using namespace sycl::ext::oneapi; // expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::force_pow2_depth(1)]] int K2; +// expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bankwidth(8)]] int K3; + struct bar { [[intel::numbanks(2)]] device_global nonconst_glob; [[intel::numbanks(4)]] const device_global const_glob; @@ -30,15 +34,41 @@ struct bar { [[intel::force_pow2_depth(0)]] device_global nonconst_glob2; [[intel::force_pow2_depth(0)]] const device_global const_glob2; [[intel::force_pow2_depth(1)]] unsigned int force_dep[64]; + + [[intel::bankwidth(2)]] device_global nonconst_glob3; + [[intel::bankwidth(4)]] const device_global const_glob3; + [[intel::bankwidth(16)]] unsigned int bankw[64]; }; +struct RandomStruct { + int M; +}; + +// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::numbanks(4)]] RandomStruct S; +// expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bankwidth(4)]] RandomStruct S1; +// expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::force_pow2_depth(1)]] RandomStruct S2; +// expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::max_replicates(8)]] RandomStruct S3; + void foo() { [[intel::numbanks(2)]] int A; [[intel::numbanks(4)]] static unsigned int ext_five[64]; + [[intel::numbanks(8)]] RandomStruct S; + [[intel::max_replicates(2)]] int A1; [[intel::max_replicates(4)]] static unsigned int ext_five1[64]; + [[intel::max_replicates(24)]] RandomStruct S1; + [[intel::force_pow2_depth(0)]] int A2; [[intel::force_pow2_depth(1)]] static unsigned int ext_five2[64]; + [[intel::force_pow2_depth(0)]] RandomStruct S2; + + [[intel::bankwidth(2)]] int A3; + [[intel::bankwidth(4)]] static unsigned int ext_five3[64]; + [[intel::bankwidth(8)]] RandomStruct S3; } void attr_on_const_no_error() @@ -46,11 +76,13 @@ void attr_on_const_no_error() [[intel::numbanks(16)]] const int const_var[64] = {0, 1}; [[intel::max_replicates(16)]] const int const_var_max[64] = {0, 1}; [[intel::force_pow2_depth(1)]] const int const_var_force[64] = {0, 1}; + [[intel::bankwidth(16)]] const int const_var_bankw[64] = {0, 1}; } void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} void attr_on_func_arg2([[intel::force_pow2_depth(1)]] int pc2) {} +void attr_on_func_arg3([[intel::bankwidth(8)]] int pc2) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; @@ -64,6 +96,9 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { // expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::force_pow2_depth(0)]] GlobAllowedVarOnly GAVO2; +// expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bankwidth(16)]] GlobAllowedVarOnly GAVO3; + [[intel::numbanks(4)]] device_global Good; [[intel::numbanks(4)]] extern device_global Bad; @@ -73,20 +108,29 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { [[intel::force_pow2_depth(0)]] device_global Good2; [[intel::force_pow2_depth(0)]] extern device_global Bad2; +[[intel::bankwidth(2)]] device_global Good3; +[[intel::bankwidth(2)]] extern device_global Bad3; + int main() { sycl::kernel_single_task([=]() { Good.get(); Good1.get(); Good2.get(); + Good3.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad1.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad2.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad3.get(); + (void)GAVO; (void)GAVO1; (void)GAVO2; + (void)GAVO3; }); return 0; } @@ -99,6 +143,11 @@ __attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3}; [[intel::max_replicates(20)]] __attribute__((opencl_global)) unsigned int ocl_glob_max_p2d[64] = {1, 2, 3}; +//expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bankwidth(32)]] +__attribute__((opencl_global)) unsigned int ocl_glob_bankw_p2d[64] = {1, 2, 3}; + + [[intel::numbanks(8)]] __attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; @@ -107,3 +156,6 @@ __attribute__((opencl_constant)) unsigned int const_var_max_rep[64] = {1, 2, 3}; [[intel::force_pow2_depth(0)]] __attribute__((opencl_constant)) unsigned int const_force_var[64] = {1, 2, 3}; + +[[intel::bankwidth(32)]] +__attribute__((opencl_constant)) unsigned int const_bankw_var[64] = {1, 2, 3}; From aa9c847ae7d018d6b51ea839157cdcb11cb865a8 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 28 Feb 2024 22:08:41 -0800 Subject: [PATCH 10/16] Add support for more fpga memory attributes on device_global variables. Address review comments and add field check. --- clang/include/clang/Basic/Attr.td | 50 ---- .../clang/Basic/DiagnosticSemaKinds.td | 5 +- clang/lib/Sema/SemaDeclAttr.cpp | 123 ++++++-- clang/test/CodeGenSYCL/device_global.cpp | 32 +++ .../SemaSYCL/intel-fpga-device-global.cpp | 271 +++++++++++++++++- clang/test/SemaSYCL/intel-fpga-local.cpp | 6 +- 6 files changed, 413 insertions(+), 74 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 22baac7e149fa..02d10c86ec43f 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2636,50 +2636,14 @@ def SYCLIntelEnableLoopPipelining : StmtAttr { def : MutualExclusions<[SYCLIntelDisableLoopPipelining, SYCLIntelEnableLoopPipelining]>; -def SYCLIntelLocalNonConstVar : SubsetSubjecthasLocalStorage() && - S->getKind() != Decl::ImplicitParam && - S->getKind() != Decl::ParmVar && - S->getKind() != Decl::NonTypeTemplateParm && - !S->getType().isConstQualified()}], - "local non-const variables">; - -def SYCLIntelConstVar : SubsetSubjectgetKind() != Decl::ImplicitParam && - S->getKind() != Decl::ParmVar && - S->getKind() != Decl::NonTypeTemplateParm && - (S->getType().isConstQualified() || - S->getType().getAddressSpace() == - LangAS::opencl_constant)}], - "constant variables">; - -def SYCLIntelLocalStaticAgentMemVar : SubsetSubjectgetKind() != Decl::ImplicitParam && - S->getKind() != Decl::NonTypeTemplateParm && - (S->getStorageClass() == SC_Static || - S->hasLocalStorage())}], - "local variables, static variables, agent memory arguments">; - -def SYCLIntelLocalOrStaticVar : SubsetSubjectgetKind() != Decl::ImplicitParam && - S->getKind() != Decl::ParmVar && - S->getKind() != Decl::NonTypeTemplateParm && - (S->getStorageClass() == SC_Static || - S->hasLocalStorage())}], - "local variables, static variables">; - def SYCLIntelDoublePump : Attr { let Spellings = [CXX11<"intel", "doublepump">]; - let Subjects = SubjectList<[SYCLIntelConstVar, SYCLIntelLocalOrStaticVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelDoublePumpAttrDocs]; } def SYCLIntelSinglePump : Attr { let Spellings = [CXX11<"intel", "singlepump">]; - let Subjects = SubjectList<[SYCLIntelConstVar, SYCLIntelLocalOrStaticVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelSinglePumpAttrDocs]; } @@ -2698,17 +2662,12 @@ def SYCLIntelMemory : Attr { } } }]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelMemoryAttrDocs]; } def SYCLIntelRegister : Attr { let Spellings = [CXX11<"intel", "fpga_register">]; - let Subjects = SubjectList<[SYCLIntelConstVar, SYCLIntelLocalOrStaticVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelRegisterAttrDocs]; } @@ -2735,7 +2694,6 @@ def SYCLIntelPrivateCopies : InheritableAttr { let Spellings = [CXX11<"intel", "private_copies">]; let Args = [ExprArgument<"Value">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let Subjects = SubjectList<[SYCLIntelLocalNonConstVar, Field], ErrorDiag>; let Documentation = [SYCLIntelPrivateCopiesAttrDocs]; } def : MutualExclusions<[SYCLIntelRegister, SYCLIntelPrivateCopies]>; @@ -2744,8 +2702,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelPrivateCopies]>; def SYCLIntelMerge : Attr { let Spellings = [CXX11<"intel", "merge">]; let Args = [StringArgument<"Name">, StringArgument<"Direction">]; - let Subjects = SubjectList<[SYCLIntelConstVar, SYCLIntelLocalOrStaticVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelMergeAttrDocs]; } @@ -2761,9 +2717,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelMaxReplicates]>; def SYCLIntelSimpleDualPort : Attr { let Spellings = [CXX11<"intel", "simple_dual_port">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelSimpleDualPortAttrDocs]; } @@ -2788,9 +2741,6 @@ def SYCLIntelPipeIO : InheritableAttr { def SYCLIntelBankBits : Attr { let Spellings = [CXX11<"intel", "bank_bits">]; let Args = [VariadicExprArgument<"Args">]; - let Subjects = SubjectList<[SYCLIntelConstVar, - SYCLIntelLocalStaticAgentMemVar, - Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let Documentation = [SYCLIntelBankBitsDocs]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 7777b7e1906e0..43b346083a700 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12217,8 +12217,11 @@ def err_sycl_attribute_not_device_global : Error<"%0 attribute can only be applied to 'device_global' variables">; def err_fpga_attribute_incorrect_variable : Error<"%0 attribute only applies to constant variables, local variables, " - "static variables, agent memory arguments, non-static data " + "static variables, %select{|agent memory arguments, }1non-static data " "members and device_global variables">; +def err_fpga_attribute_invalid_decl + : Error<"%0 attribute only applies to const variables, local variables, " + "non-static data members and device_global variables">; def err_sycl_compiletime_property_duplication : Error< "can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 301b2162c5aa0..8215b42c3a775 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7445,18 +7445,16 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D, // Checks if FPGA memory attributes apply on valid variables. // Returns true if an error occured. -static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D, - const AttributeCommonInfo &CI) { +static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D) { if (const auto *VD = dyn_cast(D)){ - if (!(VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) || - (VD->getKind() != Decl::ParmVar && - (S.isTypeDecoratedWithDeclAttribute(VD->getType()) - || VD->getType().isConstQualified() - || VD->getType().getAddressSpace() == - LangAS::opencl_constant))))) { - S.Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI; + if (!(isa(D) || + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + (VD->getStorageClass() == SC_Static || VD->hasLocalStorage() || + S.isTypeDecoratedWithDeclAttribute(VD->getType()) || + VD->getType().isConstQualified() || + VD->getType().getAddressSpace() == + LangAS::opencl_constant)))) { return true; } } @@ -7541,6 +7539,14 @@ static void handleSYCLIntelSinglePumpAttr(Sema &S, Decl *D, } } + // Check attribute applies to field, constant variables, local variables, + // static variables, non-static data members, and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL + << /*agent memory arguments*/ 0; + return; + } + // If the declaration does not have an [[intel::fpga_memory]] // attribute, this creates one as an implicit attribute. if (!D->hasAttr()) @@ -7564,6 +7570,14 @@ static void handleSYCLIntelDoublePumpAttr(Sema &S, Decl *D, } } + // Check attribute applies to field, constant variables, local variables, + // static variables, non-static data members, and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL + << /*agent memory arguments*/ 0; + return; + } + // If the declaration does not have an [[intel::fpga_memory]] // attribute, this creates one as an implicit attribute. if (!D->hasAttr()) @@ -7611,6 +7625,15 @@ static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D, const ParsedAttr &AL) { D->dropAttr(); } + // Check attribute applies to field, constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL + << /*agent memory arguments*/ 1; + return; + } + D->addAttr(::new (S.Context) SYCLIntelMemoryAttr(S.Context, AL, Kind)); } @@ -7643,6 +7666,14 @@ static void handleSYCLIntelRegisterAttr(Sema &S, Decl *D, } } + // Check attribute applies to field, constant variables, local variables, + // static variables, non-static data members, and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(A.getLoc(), diag::err_fpga_attribute_incorrect_variable) << A + << /*agent memory arguments*/ 0; + return; + } + if (checkIntelFPGARegisterAttrCompatibility(S, D, A)) return; @@ -7681,11 +7712,14 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, return; } - // Check attribute only applies to constant variables, local variables, + // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI + << /*agent memory arguments*/ 1; return; + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. @@ -7771,11 +7805,14 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, } } - // Check attribute only applies to constant variables, local variables, + // Check attribute applies to constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI + << /*agent memory arguments*/ 1; return; + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. @@ -7844,6 +7881,15 @@ static void handleIntelSimpleDualPortAttr(Sema &S, Decl *D, } } + // Check attribute applies to field, constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL + << /*agent memory arguments*/ 1; + return; + } + if (!D->hasAttr()) D->addAttr(SYCLIntelMemoryAttr::CreateImplicit( S.Context, SYCLIntelMemoryAttr::Default)); @@ -7870,11 +7916,14 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, return; } - // Check attribute only applies to constant variables, local variables, + // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI + << /*agent memory arguments*/ 1; return; + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. @@ -7959,6 +8008,14 @@ static void handleSYCLIntelMergeAttr(Sema &S, Decl *D, const ParsedAttr &AL) { return; } + // Check attribute applies to field, constant variables, local variables, + // static variables, non-static data members, and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL + << /*agent memory arguments*/ 0; + return; + } + if (!D->hasAttr()) D->addAttr(SYCLIntelMemoryAttr::CreateImplicit( S.Context, SYCLIntelMemoryAttr::Default)); @@ -8044,6 +8101,15 @@ void Sema::AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI, D->addAttr(SYCLIntelNumBanksAttr::CreateImplicit(Context, NBE)); } + // Check attribute applies to field, constant variables, local variables, + // static variables, agent memory arguments, non-static data members, + // and device_global variables. + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI + << /*agent memory arguments*/ 1; + return; + } + if (!D->hasAttr()) D->addAttr(SYCLIntelMemoryAttr::CreateImplicit( Context, SYCLIntelMemoryAttr::Default)); @@ -8069,6 +8135,22 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI, << CI << /*non-negative*/ 1; return; } + + // Check attribute applies to field as well as const variables, local + // variables, non-static data members, and device_global variables. + if (const auto *VD = dyn_cast(D)){ + if (!(isa(D) || + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + VD->getKind() != Decl::ParmVar && + (VD->hasLocalStorage() || + isTypeDecoratedWithDeclAttribute(VD->getType()) + || VD->getType().isConstQualified())))) { + Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI; + return; + } + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { @@ -8119,11 +8201,14 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, return; } - // Check attribute only applies to constant variables, local variables, + // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D, CI)) + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI + << /*agent memory arguments*/ 1; return; + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index ccca5c5b92f78..6a828c63719cc 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -16,6 +16,14 @@ device_global A; [[intel::max_replicates(2)]] device_global Nonconst_glob1; [[intel::force_pow2_depth(1)]] device_global Nonconst_glob2; [[intel::bankwidth(2)]] device_global Nonconst_glob3; +[[intel::simple_dual_port]] device_global Nonconst_glob4; +[[intel::fpga_memory]] device_global Nonconst_glob5; +[[intel::bank_bits(3, 4)]] device_global Nonconst_glob6; +[[intel::fpga_register]] device_global Nonconst_glob7; +[[intel::doublepump]] device_globalNonconst_glob8; +[[intel::singlepump]] device_global Nonconst_glob9; +[[intel::merge("mrg5", "width")]] device_global Nonconst_glob10; +[[intel::private_copies(8)]] device_global Nonconst_glob11; #ifdef SYCL_EXTERNAL SYCL_EXTERNAL device_global AExt; @@ -33,6 +41,14 @@ device_global Foo::C; // CHECK: @Nonconst_glob1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Max_ATTRS:[0-9]+]] // CHECK: @Nonconst_glob2 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Force_ATTRS:[0-9]+]] // CHECK: @Nonconst_glob3 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Bankw_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob4 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Simple_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob5 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Mem_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob6 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Bankbits_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob7 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Reg_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob8 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Dpump_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob9 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Spump_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob10 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Merge_ATTRS:[0-9]+]] +// CHECK: @Nonconst_glob11 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[Non_Const_Pc_ATTRS:[0-9]+]] // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]] // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]] @@ -85,6 +101,14 @@ void foo() { (void)Nonconst_glob1; (void)Nonconst_glob2; (void)Nonconst_glob3; + (void)Nonconst_glob4; + (void)Nonconst_glob5; + (void)Nonconst_glob6; + (void)Nonconst_glob7; + (void)Nonconst_glob8; + (void)Nonconst_glob9; + (void)Nonconst_glob10; + (void)Nonconst_glob11; (void)B; (void)Foo::C; (void)same_name; @@ -122,6 +146,14 @@ void bar() { // CHECK: attributes #[[Non_Const_Max_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob1" } // CHECK: attributes #[[Non_Const_Force_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob2" } // CHECK: attributes #[[Non_Const_Bankw_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob3" } +// CHECK: attributes #[[Non_Const_Simple_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob4" } +// CHECK: attributes #[[Non_Const_Mem_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob5" } +// CHECK: attributes #[[Non_Const_Bankbits_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob6" } +// CHECK: attributes #[[Non_Const_Reg_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob7" } +// CHECK: attributes #[[Non_Const_Dpump_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob8" } +// CHECK: attributes #[[Non_Const_Spump_ATTRS]] = { "sycl-unique-id"="_Z14Nonconst_glob9" } +// CHECK: attributes #[[Non_Const_Merge_ATTRS]] = { "sycl-unique-id"="_Z15Nonconst_glob10" } +// CHECK: attributes #[[Non_Const_Pc_ATTRS]] = { "sycl-unique-id"="_Z15Nonconst_glob11" } // 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" } diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 55b9bff1d1e87..728a25b7b4227 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Tests that [[intel::numbanks()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]], [[intel::bankwidth()]] only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables. +// Tests that [[intel::numbanks()]], [[intel::fpga_register]], [[intel::private_copies()]], [[intel::doublepump]], [[intel::singlepump]], [[intel::merge()]], [[intel::fpga_memory()]], [[intel::bank_bits()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]], [[intel::bankwidth()]], [[intel::simple_dual_port]] can be applied to device_global variables as well as constant variables, local variables, static variables, agent memory arguments, non-static data members. using namespace sycl::ext::oneapi; @@ -9,6 +9,13 @@ using namespace sycl::ext::oneapi; [[intel::max_replicates(12)]] static device_global static_dev_glob1; [[intel::force_pow2_depth(1)]] static device_global static_dev_glob2; [[intel::bankwidth(4)]] static device_global static_dev_glob3; +[[intel::simple_dual_port]] static device_global static_dev_glob4; +[[intel::fpga_memory]] static device_global static_dev_glob5; +[[intel::bank_bits(3, 4)]] static device_global static_dev_glob6; +[[intel::fpga_register]] static device_global static_dev_glob7; +[[intel::doublepump]] static device_global static_dev_glob8; +[[intel::singlepump]] static device_global static_dev_glob9; +[[intel::merge("mrg5", "width")]] static device_global static_dev_glob10; // expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::numbanks(2)]] int K; @@ -22,6 +29,30 @@ using namespace sycl::ext::oneapi; // expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::bankwidth(8)]] int K3; +// expected-error@+1{{'simple_dual_port' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::simple_dual_port]] int K4; + +// expected-error@+1{{'fpga_memory' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::fpga_memory]] int K5; + +// expected-error@+1{{'bank_bits' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bank_bits(3, 4)]] int K6; + +// expected-error@+1{{'fpga_register' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::fpga_register]] int K7; + +// expected-error@+1{{'doublepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::doublepump]] int K8; + +// expected-error@+1{{'singlepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::singlepump]] int K9; + +// expected-error@+1{{'merge' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::merge("mrg3", "width")]] int K10; + +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} +[[intel::private_copies(16)]] int K12; + struct bar { [[intel::numbanks(2)]] device_global nonconst_glob; [[intel::numbanks(4)]] const device_global const_glob; @@ -38,6 +69,38 @@ struct bar { [[intel::bankwidth(2)]] device_global nonconst_glob3; [[intel::bankwidth(4)]] const device_global const_glob3; [[intel::bankwidth(16)]] unsigned int bankw[64]; + + [[intel::simple_dual_port]] device_global nonconst_glob4; + [[intel::simple_dual_port]] const device_global const_glob4; + [[intel::simple_dual_port]] unsigned int simple[64]; + + [[intel::fpga_memory]] device_global nonconst_glob5; + [[intel::fpga_memory("MLAB")]] const device_global const_glob5; + [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32]; + + [[intel::bank_bits(3, 4)]] device_global nonconst_glob6; + [[intel::bank_bits(4, 5)]] const device_global const_glob6; + [[intel::bank_bits(3, 4)]] unsigned int mem_block_bits[32]; + + [[intel::fpga_register]] device_global nonconst_glob7; + [[intel::fpga_register]] const device_global const_glob7; + [[intel::fpga_register]] unsigned int reg; + + [[intel::singlepump]] device_global nonconst_glob8; + [[intel::singlepump]] const device_global const_glob8; + [[intel::singlepump]] unsigned int spump; + + [[intel::doublepump]] device_global nonconst_glob9; + [[intel::doublepump]] const device_global const_glob9; + [[intel::doublepump]] unsigned int dpump; + + [[intel::merge("mrg6", "depth")]] device_global nonconst_glob10; + [[intel::merge("mrg6", "depth")]] const device_global const_glob10; + [[intel::merge("mrg6", "width")]] unsigned int mergewidth; + + [[intel::private_copies(32)]] device_global nonconst_glob11; + [[intel::private_copies(8)]] const device_global const_glob11; + [[intel::private_copies(8)]] unsigned int pc; }; struct RandomStruct { @@ -52,6 +115,29 @@ struct RandomStruct { [[intel::force_pow2_depth(1)]] RandomStruct S2; // expected-error@+1{{'max_replicates' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::max_replicates(8)]] RandomStruct S3; +// expected-error@+1{{'simple_dual_port' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::simple_dual_port]] RandomStruct S4; + +// expected-error@+1{{'fpga_memory' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::fpga_memory]] RandomStruct S5; + +// expected-error@+1{{'bank_bits' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bank_bits(4, 5)]] RandomStruct S6; + +// expected-error@+1{{'fpga_register' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::fpga_register]] RandomStruct S7; + +// expected-error@+1{{'singlepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::singlepump]] RandomStruct S8; + +// expected-error@+1{{'doublepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::doublepump]] RandomStruct S9; + +// expected-error@+1{{'merge' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::merge("mrg1", "width")]] RandomStruct S10; + +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} +[[intel::private_copies(32)]] RandomStruct S11; void foo() { [[intel::numbanks(2)]] int A; @@ -69,6 +155,39 @@ void foo() { [[intel::bankwidth(2)]] int A3; [[intel::bankwidth(4)]] static unsigned int ext_five3[64]; [[intel::bankwidth(8)]] RandomStruct S3; + + [[intel::simple_dual_port]] int A4; + [[intel::simple_dual_port]] static unsigned int ext_five4[64]; + [[intel::simple_dual_port]] RandomStruct S4; + + [[intel::fpga_memory("BLOCK_RAM")]] int A5; + [[intel::fpga_memory("MLAB")]] static unsigned int ext_five5[64]; + [[intel::fpga_memory]] RandomStruct S5; + + [[intel::bank_bits(6, 7)]] int A6; + [[intel::bank_bits(9, 10)]] static unsigned int ext_five6[64]; + [[intel::bank_bits(4, 5)]] RandomStruct S6; + + [[intel::fpga_register]] int A7; + [[intel::fpga_register]] static unsigned int ext_five7[64]; + [[intel::fpga_register]] RandomStruct S7; + + [[intel::singlepump]] int A8; + [[intel::singlepump]] static unsigned int ext_five8[64]; + [[intel::singlepump]] RandomStruct S8; + + [[intel::doublepump]] int A9; + [[intel::doublepump]] static unsigned int ext_five9[64]; + [[intel::doublepump]] RandomStruct S9; + + [[intel::merge("mrg1", "depth")]] int A10; + [[intel::merge("mrg1", "width")]] static unsigned int ext_five10[64]; + [[intel::merge("mrg1", "width")]] RandomStruct S10; + + [[intel::private_copies(8)]] int A11; + //expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} + [[intel::private_copies(16)]] static unsigned int ext_five11[64]; + [[intel::private_copies(32)]] RandomStruct S11; } void attr_on_const_no_error() @@ -77,12 +196,26 @@ void attr_on_const_no_error() [[intel::max_replicates(16)]] const int const_var_max[64] = {0, 1}; [[intel::force_pow2_depth(1)]] const int const_var_force[64] = {0, 1}; [[intel::bankwidth(16)]] const int const_var_bankw[64] = {0, 1}; + [[intel::simple_dual_port]] const int const_var_simple_dual[64] = {0, 1}; + [[intel::fpga_memory]] const int const_var_mem[64] = {0, 1}; + [[intel::bank_bits(6, 7)]] const int const_var_bits[64] = {0, 1}; + [[intel::fpga_register]] const int const_var_regis[64] = {0, 1}; + [[intel::singlepump]] const int const_var_spump[64] = {0, 1}; + [[intel::doublepump]] const int const_var_dpump[64] = {0, 1}; + [[intel::merge("mrg6", "width")]] const int const_var_mergewid[64] = {0, 1}; } void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} void attr_on_func_arg2([[intel::force_pow2_depth(1)]] int pc2) {} void attr_on_func_arg3([[intel::bankwidth(8)]] int pc2) {} +void attr_on_func_arg4([[intel::simple_dual_port]] int pc3) {} +void attr_on_func_arg5([[intel::fpga_memory]] int pc4) {} +void attr_on_func_arg6([[intel::bank_bits(7, 8)]] int pc5) {} +void attr_on_func_arg7([[intel::fpga_register]] int pc6) {} +void attr_on_func_arg8([[intel::singlepump]] int pc7) {} +void attr_on_func_arg9([[intel::doublepump]] int pc8) {} +void attr_on_func_arg10([[intel::merge("mrg1", "width")]] int pc9) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; @@ -99,6 +232,30 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { // expected-error@+1{{'bankwidth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} [[intel::bankwidth(16)]] GlobAllowedVarOnly GAVO3; +// expected-error@+1{{'simple_dual_port' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::simple_dual_port]] GlobAllowedVarOnly GAVO4; + +// expected-error@+1{{'fpga_memory' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::fpga_memory]] GlobAllowedVarOnly GAVO5; + +// expected-error@+1{{'bank_bits' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bank_bits(6, 7)]] GlobAllowedVarOnly GAVO6; + +// expected-error@+1{{'fpga_register' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::fpga_register]] GlobAllowedVarOnly GAVO7; + +// expected-error@+1{{'singlepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::singlepump]] GlobAllowedVarOnly GAVO8; + +// expected-error@+1{{'doublepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::doublepump]] GlobAllowedVarOnly GAVO9; + +// expected-error@+1{{'merge' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::merge("mrg5", "width")]] GlobAllowedVarOnly GAVO10; + +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} +[[intel::private_copies(16)]] GlobAllowedVarOnly GAVO11; + [[intel::numbanks(4)]] device_global Good; [[intel::numbanks(4)]] extern device_global Bad; @@ -111,12 +268,44 @@ struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { [[intel::bankwidth(2)]] device_global Good3; [[intel::bankwidth(2)]] extern device_global Bad3; +[[intel::simple_dual_port]] device_global Good4; +[[intel::simple_dual_port]] extern device_global Bad4; + +[[intel::fpga_memory("MLAB")]] device_global Good5; +[[intel::fpga_memory("BLOCK_RAM")]] extern device_global Bad5; + +[[intel::bank_bits(6, 7)]] device_global Good6; +[[intel::bank_bits(7, 8)]] extern device_global Bad6; + +[[intel::fpga_register]] device_global Good7; +[[intel::fpga_register]] extern device_global Bad7; + +[[intel::doublepump]] device_global Good8; +[[intel::doublepump]] extern device_global Bad8; + +[[intel::singlepump]] device_global Good9; +[[intel::singlepump]] extern device_global Bad9; + +[[intel::merge("mrg1", "depth")]] device_global Good10; +[[intel::merge("mrg1", "depth")]] extern device_global Bad10; + +[[intel::private_copies(16)]] device_global Good11; +[[intel::private_copies(16)]] extern device_global Bad11; + int main() { sycl::kernel_single_task([=]() { Good.get(); Good1.get(); Good2.get(); Good3.get(); + Good4.get(); + Good5.get(); + Good6.get(); + Good7.get(); + Good8.get(); + Good9.get(); + Good10.get(); + Good11.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad.get(); @@ -126,11 +315,35 @@ int main() { Bad2.get(); // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} Bad3.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad4.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad5.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad6.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad7.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad8.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad9.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad10.get(); + // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}} + Bad11.get(); (void)GAVO; (void)GAVO1; (void)GAVO2; (void)GAVO3; + (void)GAVO4; + (void)GAVO5; + (void)GAVO6; + (void)GAVO7; + (void)GAVO8; + (void)GAVO9; + (void)GAVO10; + (void)GAVO11; }); return 0; } @@ -147,6 +360,53 @@ __attribute__((opencl_global)) unsigned int ocl_glob_max_p2d[64] = {1, 2, 3}; [[intel::bankwidth(32)]] __attribute__((opencl_global)) unsigned int ocl_glob_bankw_p2d[64] = {1, 2, 3}; +//expected-error@+1{{'simple_dual_port' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::simple_dual_port]] +__attribute__((opencl_global)) unsigned int ocl_glob_simple_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'fpga_memory' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::fpga_memory("MLAB")]] +__attribute__((opencl_global)) unsigned int ocl_glob_memory_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'bank_bits' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} +[[intel::bank_bits(7, 8)]] +__attribute__((opencl_global)) unsigned int ocl_glob_bank_bits_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'fpga_register' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::fpga_register]] +__attribute__((opencl_global)) unsigned int ocl_glob_reg_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'doublepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::doublepump]] +__attribute__((opencl_global)) unsigned int ocl_glob_dpump_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'singlepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::singlepump]] +__attribute__((opencl_global)) unsigned int ocl_glob_spump_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'merge' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} +[[intel::merge("mrg4", "depth")]] +__attribute__((opencl_global)) unsigned int ocl_glob_mer_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} +[[intel::private_copies(8)]] +__attribute__((opencl_global)) unsigned int ocl_glob_pc_p2d[64] = {1, 2, 3}; + +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} +[[intel::private_copies(8)]] +__attribute__((opencl_constant)) unsigned int const_var_private_copies[64] = {1, 2, 3}; + +[[intel::merge("mrg5", "width")]] +__attribute__((opencl_constant)) unsigned int const_var_merge[64] = {1, 2, 3}; + +[[intel::fpga_register]] +__attribute__((opencl_constant)) unsigned int const_var_fpga_register[64] = {1, 2, 3}; + +[[intel::fpga_memory]] +__attribute__((opencl_constant)) unsigned int const_var_fpga_memory[64] = {1, 2, 3}; + +[[intel::bank_bits(2, 3)]] +__attribute__((opencl_constant)) unsigned int const_var_bank_bits[64] = {1, 2, 3}; [[intel::numbanks(8)]] __attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; @@ -159,3 +419,12 @@ __attribute__((opencl_constant)) unsigned int const_force_var[64] = {1, 2, 3}; [[intel::bankwidth(32)]] __attribute__((opencl_constant)) unsigned int const_bankw_var[64] = {1, 2, 3}; + +[[intel::simple_dual_port]] +__attribute__((opencl_constant)) unsigned int const_simple_var[64] = {1, 2, 3}; + +[[intel::doublepump]] +__attribute__((opencl_constant)) unsigned int const_dpump_var[64] = {1, 2, 3}; + +[[intel::singlepump]] +__attribute__((opencl_constant)) unsigned int const_spump_var[64] = {1, 2, 3}; diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index 753c5b727a23b..b171c5078d9f2 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -442,17 +442,17 @@ void check_gnu_style() { int __attribute__((force_pow2_depth(0))) force_p2d; } -//expected-error@+1{{attribute only applies to local non-const variables and non-static data members}} +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} [[intel::private_copies(8)]] __attribute__((opencl_constant)) unsigned int const_var[64] = {1, 2, 3}; void attr_on_const_error() { - //expected-error@+1{{attribute only applies to local non-const variables and non-static data members}} + //expected-no-error@+1 [[intel::private_copies(8)]] const int const_var[64] = {0, 1}; } -//expected-error@+1{{attribute only applies to local non-const variables and non-static data members}} +//expected-error@+1{{'private_copies' attribute only applies to const variables, local variables, non-static data members and device_global variables}} void attr_on_func_arg([[intel::private_copies(8)]] int pc) {} //expected-error@+1{{'force_pow2_depth' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members and device_global variables}} From cf17e1ea06b01fa07dfad064d629cf5c66de4a8b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 29 Feb 2024 06:18:22 -0800 Subject: [PATCH 11/16] Remove duplicate cases --- clang/test/SemaSYCL/intel-fpga-device-global.cpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 728a25b7b4227..3d9ec1d6c2a38 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -205,18 +205,6 @@ void attr_on_const_no_error() [[intel::merge("mrg6", "width")]] const int const_var_mergewid[64] = {0, 1}; } -void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} -void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} -void attr_on_func_arg2([[intel::force_pow2_depth(1)]] int pc2) {} -void attr_on_func_arg3([[intel::bankwidth(8)]] int pc2) {} -void attr_on_func_arg4([[intel::simple_dual_port]] int pc3) {} -void attr_on_func_arg5([[intel::fpga_memory]] int pc4) {} -void attr_on_func_arg6([[intel::bank_bits(7, 8)]] int pc5) {} -void attr_on_func_arg7([[intel::fpga_register]] int pc6) {} -void attr_on_func_arg8([[intel::singlepump]] int pc7) {} -void attr_on_func_arg9([[intel::doublepump]] int pc8) {} -void attr_on_func_arg10([[intel::merge("mrg1", "width")]] int pc9) {} - struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; From 5fad45db5f5537e6924a516256a771f9acff1337 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 29 Feb 2024 06:56:53 -0800 Subject: [PATCH 12/16] Add more test case --- clang/test/SemaSYCL/intel-fpga-device-global.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 3d9ec1d6c2a38..bf1a881e5752e 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -205,6 +205,16 @@ void attr_on_const_no_error() [[intel::merge("mrg6", "width")]] const int const_var_mergewid[64] = {0, 1}; } +void attr_on_func_arg([[intel::numbanks(8)]] int pc) {} +void attr_on_func_arg1([[intel::max_replicates(8)]] int pc1) {} +void attr_on_func_arg2([[intel::force_pow2_depth(1)]] int pc2) {} +void attr_on_func_arg3([[intel::bankwidth(8)]] int pc3) {} +void attr_on_func_arg4([[intel::simple_dual_port]] int pc4) {} +void attr_on_func_arg5([[intel::fpga_memory]] int pc5) {} +void attr_on_func_arg6([[intel::bank_bits(7, 8)]] int pc6) {} +void attr_on_func_arg8([[intel::singlepump]] int pc7) {} +void attr_on_func_arg9([[intel::doublepump]] int pc8) {} + struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; From 41b79a419f7ea46a5e3e6a2a6023da6298da5214 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 29 Feb 2024 07:56:23 -0800 Subject: [PATCH 13/16] Added missing test case --- clang/test/SemaSYCL/intel-fpga-device-global.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index bf1a881e5752e..5a0852414037f 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -212,8 +212,10 @@ void attr_on_func_arg3([[intel::bankwidth(8)]] int pc3) {} void attr_on_func_arg4([[intel::simple_dual_port]] int pc4) {} void attr_on_func_arg5([[intel::fpga_memory]] int pc5) {} void attr_on_func_arg6([[intel::bank_bits(7, 8)]] int pc6) {} -void attr_on_func_arg8([[intel::singlepump]] int pc7) {} -void attr_on_func_arg9([[intel::doublepump]] int pc8) {} +void attr_on_func_arg7([[intel::singlepump]] int pc7) {} +void attr_on_func_arg8([[intel::doublepump]] int pc8) {} +void attr_on_func_arg9([[intel::fpga_register]] int pc9) {} +void attr_on_func_arg10([[intel::merge("mrg1", "width")]] int pc10) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { }; From a3144828043a826c2e45826f7936b4e38d4b957a Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 29 Feb 2024 10:18:03 -0800 Subject: [PATCH 14/16] Fix existing attributes restrictions --- clang/lib/Sema/SemaDeclAttr.cpp | 62 ++++++++++++------- .../SemaSYCL/intel-fpga-device-global.cpp | 4 ++ 2 files changed, 45 insertions(+), 21 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 8215b42c3a775..24cce31b9e29e 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7445,16 +7445,36 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D, // Checks if FPGA memory attributes apply on valid variables. // Returns true if an error occured. -static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D) { +static bool SYCLIntelConstLocalStaticAgentMemVar(Sema &S, Decl *D) { if (const auto *VD = dyn_cast(D)){ if (!(isa(D) || - (VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - (VD->getStorageClass() == SC_Static || VD->hasLocalStorage() || - S.isTypeDecoratedWithDeclAttribute(VD->getType()) || - VD->getType().isConstQualified() || - VD->getType().getAddressSpace() == - LangAS::opencl_constant)))) { + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + (VD->getStorageClass() == SC_Static || + VD->hasLocalStorage() || + (VD->getKind() != Decl::ParmVar && + (S.isTypeDecoratedWithDeclAttribute(VD->getType()) || + VD->getType().isConstQualified() || + VD->getType().getAddressSpace() == + LangAS::opencl_constant)))))) { + return true; + } + } + return false; +} + +static bool SYCLIntelConstOrLocalOrStaticVar(Sema &S, Decl *D) { + if (const auto *VD = dyn_cast(D)){ + if (!(isa(D) || + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + VD->getKind() != Decl::ParmVar && + (S.isTypeDecoratedWithDeclAttribute(VD->getType()) || + VD->getType().isConstQualified() || + VD->getType().getAddressSpace() == + LangAS::opencl_constant || + VD->getStorageClass() == SC_Static || + VD->hasLocalStorage())))) { return true; } } @@ -7541,7 +7561,7 @@ static void handleSYCLIntelSinglePumpAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL << /*agent memory arguments*/ 0; return; @@ -7572,7 +7592,7 @@ static void handleSYCLIntelDoublePumpAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL << /*agent memory arguments*/ 0; return; @@ -7628,7 +7648,7 @@ static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(S, D)) { S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL << /*agent memory arguments*/ 1; return; @@ -7668,7 +7688,7 @@ static void handleSYCLIntelRegisterAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { S.Diag(A.getLoc(), diag::err_fpga_attribute_incorrect_variable) << A << /*agent memory arguments*/ 0; return; @@ -7715,7 +7735,7 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI << /*agent memory arguments*/ 1; return; @@ -7808,7 +7828,7 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI << /*agent memory arguments*/ 1; return; @@ -7884,7 +7904,7 @@ static void handleIntelSimpleDualPortAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(S, D)) { S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL << /*agent memory arguments*/ 1; return; @@ -7919,9 +7939,9 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + << /*agent memory arguments*/ 1; return; } @@ -8010,7 +8030,7 @@ static void handleSYCLIntelMergeAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(S, D)) { + if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL << /*agent memory arguments*/ 0; return; @@ -8104,7 +8124,7 @@ void Sema::AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI << /*agent memory arguments*/ 1; return; @@ -8204,9 +8224,9 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + << /*agent memory arguments*/ 1; return; } diff --git a/clang/test/SemaSYCL/intel-fpga-device-global.cpp b/clang/test/SemaSYCL/intel-fpga-device-global.cpp index 5a0852414037f..86b31a2efa429 100644 --- a/clang/test/SemaSYCL/intel-fpga-device-global.cpp +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -212,9 +212,13 @@ void attr_on_func_arg3([[intel::bankwidth(8)]] int pc3) {} void attr_on_func_arg4([[intel::simple_dual_port]] int pc4) {} void attr_on_func_arg5([[intel::fpga_memory]] int pc5) {} void attr_on_func_arg6([[intel::bank_bits(7, 8)]] int pc6) {} +// expected-error@+1{{'singlepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} void attr_on_func_arg7([[intel::singlepump]] int pc7) {} +// expected-error@+1{{'doublepump' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} void attr_on_func_arg8([[intel::doublepump]] int pc8) {} +// expected-error@+1{{'fpga_register' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} void attr_on_func_arg9([[intel::fpga_register]] int pc9) {} +// expected-error@+1{{'merge' attribute only applies to constant variables, local variables, static variables, non-static data members and device_global variables}} void attr_on_func_arg10([[intel::merge("mrg1", "width")]] int pc10) {} struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly { From 32ae57391b3c1c4857cd21d7ba63b3b66bde9c0f Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 1 Mar 2024 14:09:17 -0800 Subject: [PATCH 15/16] Address review comments --- clang/lib/Sema/SemaDeclAttr.cpp | 120 ++++++++++++++------------------ 1 file changed, 52 insertions(+), 68 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 24cce31b9e29e..334ec82483e3c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7445,36 +7445,16 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D, // Checks if FPGA memory attributes apply on valid variables. // Returns true if an error occured. -static bool SYCLIntelConstLocalStaticAgentMemVar(Sema &S, Decl *D) { - if (const auto *VD = dyn_cast(D)){ - if (!(isa(D) || - (VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - (VD->getStorageClass() == SC_Static || - VD->hasLocalStorage() || - (VD->getKind() != Decl::ParmVar && - (S.isTypeDecoratedWithDeclAttribute(VD->getType()) || - VD->getType().isConstQualified() || - VD->getType().getAddressSpace() == - LangAS::opencl_constant)))))) { - return true; - } - } - return false; -} - -static bool SYCLIntelConstOrLocalOrStaticVar(Sema &S, Decl *D) { - if (const auto *VD = dyn_cast(D)){ +static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D) { + if (const auto *VD = dyn_cast(D)) { if (!(isa(D) || (VD->getKind() != Decl::ImplicitParam && VD->getKind() != Decl::NonTypeTemplateParm && - VD->getKind() != Decl::ParmVar && - (S.isTypeDecoratedWithDeclAttribute(VD->getType()) || + (S.isTypeDecoratedWithDeclAttribute( + VD->getType()) || VD->getType().isConstQualified() || - VD->getType().getAddressSpace() == - LangAS::opencl_constant || - VD->getStorageClass() == SC_Static || - VD->hasLocalStorage())))) { + VD->getType().getAddressSpace() == LangAS::opencl_constant || + VD->getStorageClass() == SC_Static || VD->hasLocalStorage())))) { return true; } } @@ -7561,11 +7541,12 @@ static void handleSYCLIntelSinglePumpAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { - S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL - << /*agent memory arguments*/ 0; + if ((D->getKind() == Decl::ParmVar) || + CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << AL << /*agent memory arguments*/ 0; return; - } + } // If the declaration does not have an [[intel::fpga_memory]] // attribute, this creates one as an implicit attribute. @@ -7592,11 +7573,12 @@ static void handleSYCLIntelDoublePumpAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { - S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL - << /*agent memory arguments*/ 0; + if ((D->getKind() == Decl::ParmVar) || + CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << AL << /*agent memory arguments*/ 0; return; - } + } // If the declaration does not have an [[intel::fpga_memory]] // attribute, this creates one as an implicit attribute. @@ -7648,11 +7630,11 @@ static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(S, D)) { - S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << AL << /*agent memory arguments*/ 1; return; - } + } D->addAttr(::new (S.Context) SYCLIntelMemoryAttr(S.Context, AL, Kind)); } @@ -7688,11 +7670,12 @@ static void handleSYCLIntelRegisterAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { - S.Diag(A.getLoc(), diag::err_fpga_attribute_incorrect_variable) << A - << /*agent memory arguments*/ 0; + if ((D->getKind() == Decl::ParmVar) || + CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(A.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << A << /*agent memory arguments*/ 0; return; - } + } if (checkIntelFPGARegisterAttrCompatibility(S, D, A)) return; @@ -7735,9 +7718,9 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << CI << /*agent memory arguments*/ 1; return; } @@ -7828,9 +7811,9 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << CI << /*agent memory arguments*/ 1; return; } @@ -7904,9 +7887,9 @@ static void handleIntelSimpleDualPortAttr(Sema &S, Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(S, D)) { - S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << AL << /*agent memory arguments*/ 1; return; } @@ -7939,9 +7922,9 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << CI << /*agent memory arguments*/ 1; return; } @@ -8030,11 +8013,12 @@ static void handleSYCLIntelMergeAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Check attribute applies to field, constant variables, local variables, // static variables, non-static data members, and device_global variables. - if (SYCLIntelConstOrLocalOrStaticVar(S, D)) { - S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) << AL - << /*agent memory arguments*/ 0; + if ((D->getKind() == Decl::ParmVar) || + CheckValidFPGAMemoryAttributesVar(S, D)) { + S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << AL << /*agent memory arguments*/ 0; return; - } + } if (!D->hasAttr()) D->addAttr(SYCLIntelMemoryAttr::CreateImplicit( @@ -8124,9 +8108,9 @@ void Sema::AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << CI << /*agent memory arguments*/ 1; return; } @@ -8156,16 +8140,16 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI, return; } - // Check attribute applies to field as well as const variables, local - // variables, non-static data members, and device_global variables. + // Check attribute applies to field as well as const variables, non-static + // local variables, non-static data members, and device_global variables. if (const auto *VD = dyn_cast(D)){ if (!(isa(D) || (VD->getKind() != Decl::ImplicitParam && VD->getKind() != Decl::NonTypeTemplateParm && VD->getKind() != Decl::ParmVar && (VD->hasLocalStorage() || - isTypeDecoratedWithDeclAttribute(VD->getType()) - || VD->getType().isConstQualified())))) { + isTypeDecoratedWithDeclAttribute( + VD->getType()))))) { Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI; return; } @@ -8224,9 +8208,9 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, // Check attribute applies to field, constant variables, local variables, // static variables, agent memory arguments, non-static data members, // and device_global variables. - if (SYCLIntelConstLocalStaticAgentMemVar(*this, D)) { - Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) << CI - << /*agent memory arguments*/ 1; + if (CheckValidFPGAMemoryAttributesVar(*this, D)) { + Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) + << CI << /*agent memory arguments*/ 1; return; } From 4a4a016500a6896ff47c8e0147a01615719be90b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 1 Mar 2024 15:43:00 -0800 Subject: [PATCH 16/16] Fix clang format errors --- clang/lib/Sema/SemaDeclAttr.cpp | 32 ++++++++++++++++---------------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 334ec82483e3c..b11f223f1ce5d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7720,7 +7720,7 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, // and device_global variables. if (CheckValidFPGAMemoryAttributesVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) - << CI << /*agent memory arguments*/ 1; + << CI << /*agent memory arguments*/ 1; return; } @@ -7813,7 +7813,7 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI, // and device_global variables. if (CheckValidFPGAMemoryAttributesVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) - << CI << /*agent memory arguments*/ 1; + << CI << /*agent memory arguments*/ 1; return; } @@ -7924,7 +7924,7 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, // and device_global variables. if (CheckValidFPGAMemoryAttributesVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) - << CI << /*agent memory arguments*/ 1; + << CI << /*agent memory arguments*/ 1; return; } @@ -8110,7 +8110,7 @@ void Sema::AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI, // and device_global variables. if (CheckValidFPGAMemoryAttributesVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) - << CI << /*agent memory arguments*/ 1; + << CI << /*agent memory arguments*/ 1; return; } @@ -8142,17 +8142,17 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI, // Check attribute applies to field as well as const variables, non-static // local variables, non-static data members, and device_global variables. - if (const auto *VD = dyn_cast(D)){ - if (!(isa(D) || - (VD->getKind() != Decl::ImplicitParam && - VD->getKind() != Decl::NonTypeTemplateParm && - VD->getKind() != Decl::ParmVar && - (VD->hasLocalStorage() || - isTypeDecoratedWithDeclAttribute( - VD->getType()))))) { - Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI; - return; - } + if (const auto *VD = dyn_cast(D)) { + if (!(isa(D) || + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + VD->getKind() != Decl::ParmVar && + (VD->hasLocalStorage() || + isTypeDecoratedWithDeclAttribute( + VD->getType()))))) { + Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI; + return; + } } // Check to see if there's a duplicate attribute with different values @@ -8210,7 +8210,7 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, // and device_global variables. if (CheckValidFPGAMemoryAttributesVar(*this, D)) { Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable) - << CI << /*agent memory arguments*/ 1; + << CI << /*agent memory arguments*/ 1; return; }