diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index cbc03a363233a..73b0e74a221c7 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2646,50 +2646,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]; } @@ -2708,17 +2672,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]; } @@ -2729,9 +2688,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]; } @@ -2740,9 +2696,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]; } @@ -2751,7 +2704,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]>; @@ -2760,8 +2712,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]; } @@ -2770,9 +2720,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]; } @@ -2780,9 +2727,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]; } @@ -2807,9 +2751,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]; } @@ -2819,9 +2760,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/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index af29649364bd2..54feb5d797ed9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12218,6 +12218,13 @@ 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_incorrect_variable + : Error<"%0 attribute only applies to constant variables, local variables, " + "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 68949acfcd5fb..b11f223f1ce5d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7443,6 +7443,24 @@ 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) { + if (const auto *VD = dyn_cast(D)) { + if (!(isa(D) || + (VD->getKind() != Decl::ImplicitParam && + VD->getKind() != Decl::NonTypeTemplateParm && + (S.isTypeDecoratedWithDeclAttribute( + VD->getType()) || + VD->getType().isConstQualified() || + VD->getType().getAddressSpace() == LangAS::opencl_constant || + VD->getStorageClass() == SC_Static || VD->hasLocalStorage())))) { + return true; + } + } + return false; +} + void Sema::AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E) { @@ -7521,6 +7539,15 @@ 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 ((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. if (!D->hasAttr()) @@ -7544,6 +7571,15 @@ 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 ((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. if (!D->hasAttr()) @@ -7591,6 +7627,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)); } @@ -7623,6 +7668,15 @@ 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 ((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; @@ -7661,6 +7715,15 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI, return; } + // 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; + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { @@ -7745,6 +7808,15 @@ 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)) { + 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. if (const auto *DeclAttr = D->getAttr()) { @@ -7812,6 +7884,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)); @@ -7837,6 +7918,16 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI, << CI << /*positive*/ 0; return; } + + // 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; + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { @@ -7920,6 +8011,15 @@ 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 ((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( S.Context, SYCLIntelMemoryAttr::Default)); @@ -8005,6 +8105,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)); @@ -8030,6 +8139,22 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI, << CI << /*non-negative*/ 1; return; } + + // 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; + } + } + // Check to see if there's a duplicate attribute with different values // already applied to the declaration. if (const auto *DeclAttr = D->getAttr()) { @@ -8080,6 +8205,15 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D, return; } + // 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; + } + // 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..6a828c63719cc 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -11,6 +11,20 @@ using namespace sycl; queue q; 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; +[[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; #endif @@ -20,8 +34,21 @@ 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: @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: @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]+]] @@ -70,6 +97,18 @@ void foo() { q.submit([&](handler &h) { h.single_task([=]() { (void)A; + (void)Nonconst_glob; + (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; @@ -103,6 +142,18 @@ void bar() { // CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" } // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" } +// 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 #[[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 new file mode 100644 index 0000000000000..86b31a2efa429 --- /dev/null +++ b/clang/test/SemaSYCL/intel-fpga-device-global.cpp @@ -0,0 +1,434 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s +#include "Inputs/sycl.hpp" + +// 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; + +[[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; +[[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; + +// 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; + +// 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; + [[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(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]; + + [[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 { + 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; +// 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; + [[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; + + [[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() +{ + [[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}; + [[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 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 { +}; + +// 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; + +// 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; + +// 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; + +[[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; + +[[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(); + // 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(); + // 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; +} + +//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}; + +//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}; + +//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}; + +//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}; + +[[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}; + +[[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 d6b48db59bd9e..b171c5078d9f2 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -442,20 +442,20 @@ 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{{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};