From 845ff48268e0c228af3b3503bcab4276fb90c155 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 1 Feb 2021 12:21:31 -0800 Subject: [PATCH 01/10] [SYCL] Remove arbitrary range Part 1 Remove arbitrary range for SYCLIntelSchedulerTargetFmaxMhz and SYCLIntelLoopFuse attributes. Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 16 ---------------- clang/include/clang/Sema/Sema.h | 7 +++++++ clang/lib/Sema/SemaDeclAttr.cpp | 12 ++++++++---- clang/test/SemaSYCL/loop_fusion.cpp | 7 ++++--- .../test/SemaSYCL/scheduler_target_fmax_mhz.cpp | 12 ++++-------- 5 files changed, 23 insertions(+), 31 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index ead6e7952b80c..a62a80518c1f6 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1269,14 +1269,6 @@ def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr { let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [SYCLIntelSchedulerTargetFmaxMhzAttrDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 0; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def SYCLIntelMaxWorkGroupSize : InheritableAttr { @@ -1331,14 +1323,6 @@ def SYCLIntelLoopFuse : InheritableAttr { let Accessors = [Accessor<"isIndependent", [CXX11<"intel", "loop_fuse_independent">]>]; let Documentation = [SYCLIntelLoopFuseDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 0; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def C11NoReturn : InheritableAttr { diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index faa6af40ebe87..7dfb77fede0ee 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12970,6 +12970,13 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, return; } } + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz) { + if (ArgInt < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI.getAttrName() << /*non-negative*/ 1; + return; + } + } } D->addAttr(::new (Context) AttrType(Context, CI, E)); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 6851053a96e75..33497e83ebc39 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3210,7 +3210,8 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.Diag(AL.getLoc(), diag::note_spelling_suggestion) << "'intel::scheduler_target_fmax_mhz'"; - S.AddOneConstantValueAttr(D, AL, E); + S.addIntelSYCLSingleArgFunctionAttr( + D, AL, E); } // Handles max_global_work_dim. @@ -3287,10 +3288,13 @@ static bool checkSYCLIntelLoopFuseArgument(Sema &S, return true; } - SYCLIntelLoopFuseAttr TmpAttr(S.Context, CI, E); - ExprResult ICE; + if (!ArgVal->isNonNegative()) { + S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; + return true; + } - return S.checkRangedIntegralArgument(E, &TmpAttr, ICE); + return false; } void Sema::addSYCLIntelLoopFuseAttr(Decl *D, const AttributeCommonInfo &CI, diff --git a/clang/test/SemaSYCL/loop_fusion.cpp b/clang/test/SemaSYCL/loop_fusion.cpp index 389b3f3d98317..e92b7c99110db 100644 --- a/clang/test/SemaSYCL/loop_fusion.cpp +++ b/clang/test/SemaSYCL/loop_fusion.cpp @@ -4,8 +4,8 @@ [[intel::loop_fuse("foo")]] void func() {} // expected-error{{'loop_fuse' attribute requires an integer constant}} -[[intel::loop_fuse(1048577)]] void func1() {} // expected-error{{'loop_fuse' attribute requires integer constant between 0 and 1048576 inclusive}} -[[intel::loop_fuse_independent(-1)]] void func2() {} // expected-error{{'loop_fuse_independent' attribute requires integer constant between 0 and 1048576 inclusive}} +[[intel::loop_fuse(1048577)]] void func1() {} // OK +[[intel::loop_fuse_independent(-1)]] void func2() {} // expected-error{{'loop_fuse_independent' attribute requires a non-negative integral compile time constant expression}} [[intel::loop_fuse(0, 1)]] void func3() {} // expected-error{{'loop_fuse' attribute takes no more than 1 argument}} [[intel::loop_fuse_independent(2, 3)]] void func4() {} // expected-error{{'loop_fuse_independent' attribute takes no more than 1 argument}} @@ -48,13 +48,14 @@ [[intel::loop_fuse]] void func16(); template -[[intel::loop_fuse(N)]] void func17(); // expected-error{{'loop_fuse' attribute requires integer constant between 0 and 1048576 inclusive}} +[[intel::loop_fuse(N)]] void func17(); // expected-error{{'loop_fuse' attribute requires a non-negative integral compile time constant expression}} template [[intel::loop_fuse(Ty{})]] void func18() {} // expected-error{{'loop_fuse' attribute requires an integer constant}} void checkTemplates() { func17<-1>(); // expected-note{{in instantiation of}} + func17<0>(); // OK func18(); // expected-note{{in instantiation of}} } diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp index d0203e41447d6..379cf57d7b23c 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -10,20 +10,16 @@ template [[intel::scheduler_target_fmax_mhz(N)]] void zoo() {} int main() { - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel1 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 5 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5 // expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}} // expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}} cl::sycl::kernel_single_task( []() [[intelfpga::scheduler_target_fmax_mhz(5)]]{}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel2 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 cl::sycl::kernel_single_task( []() { func(); }); @@ -39,10 +35,10 @@ int main() { [[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}} cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}} + []() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}} + []() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} cl::sycl::kernel_single_task( []() [[intel::scheduler_target_fmax_mhz(1), intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}} From 224b7eb3143490e013f46b3fc7796e8ddcc4fa16 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 1 Feb 2021 13:12:00 -0800 Subject: [PATCH 02/10] Rename function to be more generic Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 9 ++++----- clang/lib/Sema/SemaDeclAttr.cpp | 14 +++++-------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 20 +++++++++---------- 3 files changed, 19 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 7dfb77fede0ee..31e7f1251b6f1 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10094,8 +10094,8 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); template - void addIntelSYCLSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); + void addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E); template void addIntelSYCLTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XDimExpr, Expr *YDimExpr, @@ -12936,9 +12936,8 @@ class Sema final { }; template -void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, - const AttributeCommonInfo &CI, - Expr *E) { +void Sema::addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E) { assert(E && "Attribute must have an argument."); if (!E->isInstantiationDependent()) { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 33497e83ebc39..12de38a296d1d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3158,7 +3158,7 @@ static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (D->getAttr()) S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; - S.addIntelSYCLSingleArgFunctionAttr(D, AL, E); + S.addSingleArgFunctionAttr(D, AL, E); } // Handles num_simd_work_items. @@ -3176,8 +3176,7 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D, S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::num_simd_work_items'"; - S.addIntelSYCLSingleArgFunctionAttr(D, Attr, - E); + S.addSingleArgFunctionAttr(D, Attr, E); } // Handles use_stall_enable_clusters @@ -3210,8 +3209,7 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.Diag(AL.getLoc(), diag::note_spelling_suggestion) << "'intel::scheduler_target_fmax_mhz'"; - S.addIntelSYCLSingleArgFunctionAttr( - D, AL, E); + S.addSingleArgFunctionAttr(D, AL, E); } // Handles max_global_work_dim. @@ -3234,8 +3232,7 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::max_global_work_dim'"; - S.addIntelSYCLSingleArgFunctionAttr(D, Attr, - E); + S.addSingleArgFunctionAttr(D, Attr, E); } SYCLIntelLoopFuseAttr * @@ -5526,8 +5523,7 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, ? Attr.getArgAsExpr(0) : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, Attr.getLoc()); - S.addIntelSYCLSingleArgFunctionAttr(D, Attr, - E); + S.addSingleArgFunctionAttr(D, Attr, E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index efb3e77e2f502..481fc09707b86 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -629,15 +629,15 @@ static void instantiateSYCLIntelLoopFuseAttr( } template -static void instantiateIntelSYCLFunctionAttr( - Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, - const AttrName *Attr, Decl *New) { +static void +instantiateFunctionAttr(Sema &S, + const MultiLevelTemplateArgumentList &TemplateArgs, + const AttrName *Attr, Decl *New) { EnterExpressionEvaluationContext Unevaluated( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); if (!Result.isInvalid()) - S.addIntelSYCLSingleArgFunctionAttr(New, *Attr, - Result.getAs()); + S.addSingleArgFunctionAttr(New, *Attr, Result.getAs()); } /// Determine whether the attribute A might be relevent to the declaration D. @@ -820,25 +820,25 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *IntelReqdSubGroupSize = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( + instantiateFunctionAttr( *this, TemplateArgs, IntelReqdSubGroupSize, New); continue; } if (const auto *SYCLIntelNumSimdWorkItems = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( + instantiateFunctionAttr( *this, TemplateArgs, SYCLIntelNumSimdWorkItems, New); continue; } if (const auto *SYCLIntelSchedulerTargetFmaxMhz = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( + instantiateFunctionAttr( *this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New); continue; } if (const auto *SYCLIntelMaxGlobalWorkDim = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( + instantiateFunctionAttr( *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New); continue; } @@ -850,7 +850,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *SYCLIntelNoGlobalWorkOffset = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( + instantiateFunctionAttr( *this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New); continue; } From 9ba170f6c49e1660dc5be3b2b7d085f6c15d49ec Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 1 Feb 2021 15:49:33 -0800 Subject: [PATCH 03/10] Remove arb range part 2 Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 8 -------- clang/include/clang/Sema/Sema.h | 9 ++++++++- clang/lib/Sema/SemaDeclAttr.cpp | 4 ++-- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 5 +++++ clang/test/SemaSYCL/intel-fpga-local.cpp | 19 +++---------------- 5 files changed, 18 insertions(+), 27 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a62a80518c1f6..dc7f4f1fd7ac4 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2042,14 +2042,6 @@ def IntelFPGAPrivateCopies : InheritableAttr { let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[IntelFPGALocalNonConstVar, Field], ErrorDiag>; let Documentation = [IntelFPGAPrivateCopiesAttrDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 0; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } // Two string arguments. diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 31e7f1251b6f1..b7d47d86302d8 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12969,7 +12969,8 @@ void Sema::addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, return; } } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz) { + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz || + CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI.getAttrName() << /*non-negative*/ 1; @@ -12978,6 +12979,12 @@ void Sema::addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, } } + if (CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) { + if (!D->hasAttr()) + D->addAttr(IntelFPGAMemoryAttr::CreateImplicit( + Context, IntelFPGAMemoryAttr::Default)); + } + D->addAttr(::new (Context) AttrType(Context, CI, E)); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 12de38a296d1d..adad345f95c7d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5862,8 +5862,8 @@ static void handleIntelFPGAPrivateCopiesAttr(Sema &S, Decl *D, S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::private_copies'"; - S.AddOneConstantValueAttr( - D, Attr, Attr.getArgAsExpr(0)); + S.addSingleArgFunctionAttr(D, Attr, + Attr.getArgAsExpr(0)); } static void handleIntelFPGAForcePow2DepthAttr(Sema &S, Decl *D, diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 481fc09707b86..72c92d841be33 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -587,6 +587,11 @@ static void instantiateIntelFPGAMemoryAttr( std::is_same::value) return S.AddOneConstantPowerTwoValueAttr(New, *Attr, Result.getAs()); + + else if (std::is_same::value) + return S.addSingleArgFunctionAttr(New, *Attr, + Result.getAs()); + return S.AddOneConstantValueAttr(New, *Attr, Result.getAs()); } diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index 08f914f06a999..bdea81ac3fcec 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -48,8 +48,6 @@ void check_ast() //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} [[intel::private_copies(8)]] unsigned int private_copies[64]; @@ -375,8 +373,6 @@ void diagnostics() //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} //expected-warning@+2 {{attribute 'intelfpga::private_copies' is deprecated}} //expected-note@+1 {{did you mean to use 'intel::private_copies' instead?}} @@ -390,24 +386,19 @@ void diagnostics() //CHECK: VarDecl{{.*}}pc_pc //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}16 //CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} //expected-warning@+2{{is already applied}} [[intel::private_copies(8)]] [[intel::private_copies(16)]] unsigned int pc_pc[64]; - //expected-error@+1{{'private_copies' attribute requires integer constant between 0 and 1048576 inclusive}} + //expected-error@+1{{'private_copies' attribute requires a non-negative integral compile time constant expression}} [[intel::private_copies(-4)]] unsigned int pc_negative[64]; - int i_private_copies = 32; // expected-note {{declared here}} - //expected-error@+1{{expression is not an integral constant expression}} + int i_private_copies = 32; + //expected-error@+1{{'private_copies' attribute requires an integer constant}} [[intel::private_copies(i_private_copies)]] - //expected-note@-1{{read of non-const variable 'i_private_copies' is not allowed in a constant expression}} unsigned int pc_nonconst[64]; //expected-error@+1{{'private_copies' attribute takes one argument}} @@ -704,8 +695,6 @@ struct foo { //CHECK: FieldDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}4 //CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} [[intel::private_copies(4)]] unsigned int private_copies[64]; @@ -755,8 +744,6 @@ void check_template_parameters() { //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr - //CHECK-NEXT: ConstantExpr - //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: SubstNonTypeTemplateParmExpr //CHECK-NEXT: NonTypeTemplateParmDecl //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} From 64b469832daa9028ba9a2ccc15840e59090a6864 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 01:36:35 -0800 Subject: [PATCH 04/10] Revert "Rename function to be more generic" This reverts commit 224b7eb3143490e013f46b3fc7796e8ddcc4fa16. --- clang/include/clang/Sema/Sema.h | 9 +++++---- clang/lib/Sema/SemaDeclAttr.cpp | 14 ++++++++----- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 20 +++++++++---------- 3 files changed, 24 insertions(+), 19 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index b7d47d86302d8..3b54afe5f6973 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10094,8 +10094,8 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); template - void addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); + void addIntelSYCLSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E); template void addIntelSYCLTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XDimExpr, Expr *YDimExpr, @@ -12936,8 +12936,9 @@ class Sema final { }; template -void Sema::addSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E) { +void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *E) { assert(E && "Attribute must have an argument."); if (!E->isInstantiationDependent()) { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index adad345f95c7d..72ef71e4263c9 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3158,7 +3158,7 @@ static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (D->getAttr()) S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; - S.addSingleArgFunctionAttr(D, AL, E); + S.addIntelSYCLSingleArgFunctionAttr(D, AL, E); } // Handles num_simd_work_items. @@ -3176,7 +3176,8 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D, S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::num_simd_work_items'"; - S.addSingleArgFunctionAttr(D, Attr, E); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } // Handles use_stall_enable_clusters @@ -3209,7 +3210,8 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.Diag(AL.getLoc(), diag::note_spelling_suggestion) << "'intel::scheduler_target_fmax_mhz'"; - S.addSingleArgFunctionAttr(D, AL, E); + S.addIntelSYCLSingleArgFunctionAttr( + D, AL, E); } // Handles max_global_work_dim. @@ -3232,7 +3234,8 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::max_global_work_dim'"; - S.addSingleArgFunctionAttr(D, Attr, E); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } SYCLIntelLoopFuseAttr * @@ -5523,7 +5526,8 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, ? Attr.getArgAsExpr(0) : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, Attr.getLoc()); - S.addSingleArgFunctionAttr(D, Attr, E); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 72c92d841be33..b645018638cff 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -634,15 +634,15 @@ static void instantiateSYCLIntelLoopFuseAttr( } template -static void -instantiateFunctionAttr(Sema &S, - const MultiLevelTemplateArgumentList &TemplateArgs, - const AttrName *Attr, Decl *New) { +static void instantiateIntelSYCLFunctionAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const AttrName *Attr, Decl *New) { EnterExpressionEvaluationContext Unevaluated( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); if (!Result.isInvalid()) - S.addSingleArgFunctionAttr(New, *Attr, Result.getAs()); + S.addIntelSYCLSingleArgFunctionAttr(New, *Attr, + Result.getAs()); } /// Determine whether the attribute A might be relevent to the declaration D. @@ -825,25 +825,25 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *IntelReqdSubGroupSize = dyn_cast(TmplAttr)) { - instantiateFunctionAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, IntelReqdSubGroupSize, New); continue; } if (const auto *SYCLIntelNumSimdWorkItems = dyn_cast(TmplAttr)) { - instantiateFunctionAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, SYCLIntelNumSimdWorkItems, New); continue; } if (const auto *SYCLIntelSchedulerTargetFmaxMhz = dyn_cast(TmplAttr)) { - instantiateFunctionAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New); continue; } if (const auto *SYCLIntelMaxGlobalWorkDim = dyn_cast(TmplAttr)) { - instantiateFunctionAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New); continue; } @@ -855,7 +855,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *SYCLIntelNoGlobalWorkOffset = dyn_cast(TmplAttr)) { - instantiateFunctionAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New); continue; } From 7f8e23bdd902bc9bdef710e2b824bd349a1bdd81 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 03:30:36 -0800 Subject: [PATCH 05/10] Fix tests after merge and revert to less generic function names Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 22 +++++++++---------- clang/lib/Sema/SemaDeclAttr.cpp | 20 ++++++++--------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 11 +++++----- clang/test/SemaSYCL/intel-fpga-local.cpp | 17 ++++++++++++-- .../SemaSYCL/scheduler_target_fmax_mhz.cpp | 4 ++++ 5 files changed, 44 insertions(+), 30 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 398083c5c46ee..fe584533ba9a7 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10206,12 +10206,12 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); template - void addIntelSYCLSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); + void addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E); template - void addIntelSYCLTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, - Expr *ZDimExpr); + void addIntelTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDimExpr, Expr *YDimExpr, + Expr *ZDimExpr); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); @@ -13059,9 +13059,8 @@ class Sema final { }; template -void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, - const AttributeCommonInfo &CI, - Expr *E) { +void Sema::addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E) { assert(E && "Attribute must have an argument."); if (!E->isInstantiationDependent()) { @@ -13142,10 +13141,9 @@ static Expr *checkMaxWorkSizeAttrExpr(Sema &S, const AttributeCommonInfo &CI, } template -void Sema::addIntelSYCLTripleArgFunctionAttr(Decl *D, - const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, - Expr *ZDimExpr) { +void Sema::addIntelTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDimExpr, Expr *YDimExpr, + Expr *ZDimExpr) { assert((XDimExpr && YDimExpr && ZDimExpr) && "argument has unexpected null value"); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 4d4b52b2e8d9f..0524b182d178b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3150,8 +3150,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (!checkWorkGroupSizeValues(S, D, AL)) return; - S.addIntelSYCLTripleArgFunctionAttr(D, AL, XDimExpr, YDimExpr, - ZDimExpr); + S.addIntelTripleArgFunctionAttr(D, AL, XDimExpr, YDimExpr, + ZDimExpr); } // Handles work_group_size_hint. @@ -3193,7 +3193,7 @@ static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (D->getAttr()) S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; - S.addIntelSYCLSingleArgFunctionAttr(D, AL, E); + S.addIntelSingleArgFunctionAttr(D, AL, E); } // Handles num_simd_work_items. @@ -3208,7 +3208,7 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgFunctionAttr(D, A, E); } // Handles use_stall_enable_clusters @@ -3239,8 +3239,8 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(AL); - S.addIntelSYCLSingleArgFunctionAttr( - D, AL, E); + S.addIntelSingleArgFunctionAttr(D, AL, + E); } // Handles max_global_work_dim. @@ -3260,7 +3260,7 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgFunctionAttr(D, A, E); } SYCLIntelLoopFuseAttr * @@ -5587,7 +5587,7 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, ? A.getArgAsExpr(0) : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, A.getLoc()); - S.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgFunctionAttr(D, A, E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. @@ -5897,8 +5897,8 @@ static void handleIntelFPGAPrivateCopiesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addSingleArgFunctionAttr(D, Attr, - Attr.getArgAsExpr(0)); + S.addIntelSingleArgFunctionAttr( + D, A, A.getArgAsExpr(0)); } static void handleIntelFPGAForcePow2DepthAttr(Sema &S, Decl *D, diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index a8dd24826c9bd..9c888889dc6b6 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -571,8 +571,8 @@ static void instantiateIntelSYCTripleLFunctionAttr( return; Expr *ZDimExpr = Result.getAs(); - S.addIntelSYCLTripleArgFunctionAttr(New, *Attr, XDimExpr, YDimExpr, - ZDimExpr); + S.addIntelTripleArgFunctionAttr(New, *Attr, XDimExpr, YDimExpr, + ZDimExpr); } template @@ -589,8 +589,8 @@ static void instantiateIntelFPGAMemoryAttr( Result.getAs()); else if (std::is_same::value) - return S.addSingleArgFunctionAttr(New, *Attr, - Result.getAs()); + return S.addIntelSingleArgFunctionAttr(New, *Attr, + Result.getAs()); return S.AddOneConstantValueAttr(New, *Attr, Result.getAs()); @@ -641,8 +641,7 @@ static void instantiateIntelSYCLFunctionAttr( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); if (!Result.isInvalid()) - S.addIntelSYCLSingleArgFunctionAttr(New, *Attr, - Result.getAs()); + S.addIntelSingleArgFunctionAttr(New, *Attr, Result.getAs()); } /// Determine whether the attribute A might be relevent to the declaration D. diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index bdea81ac3fcec..49878559d012a 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -48,6 +48,8 @@ void check_ast() //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} [[intel::private_copies(8)]] unsigned int private_copies[64]; @@ -373,6 +375,8 @@ void diagnostics() //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} //expected-warning@+2 {{attribute 'intelfpga::private_copies' is deprecated}} //expected-note@+1 {{did you mean to use 'intel::private_copies' instead?}} @@ -386,8 +390,12 @@ void diagnostics() //CHECK: VarDecl{{.*}}pc_pc //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}16 //CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} //expected-warning@+2{{is already applied}} [[intel::private_copies(8)]] @@ -396,9 +404,10 @@ void diagnostics() //expected-error@+1{{'private_copies' attribute requires a non-negative integral compile time constant expression}} [[intel::private_copies(-4)]] unsigned int pc_negative[64]; - int i_private_copies = 32; - //expected-error@+1{{'private_copies' attribute requires an integer constant}} + int i_private_copies = 32; // expected-note {{declared here}} + //expected-error@+1{{expression is not an integral constant expression}} [[intel::private_copies(i_private_copies)]] + //expected-note@-1{{read of non-const variable 'i_private_copies' is not allowed in a constant expression}} unsigned int pc_nonconst[64]; //expected-error@+1{{'private_copies' attribute takes one argument}} @@ -695,6 +704,8 @@ struct foo { //CHECK: FieldDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}4 //CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} [[intel::private_copies(4)]] unsigned int private_copies[64]; @@ -744,6 +755,8 @@ void check_template_parameters() { //CHECK: VarDecl{{.*}}private_copies //CHECK: IntelFPGAMemoryAttr{{.*}}Implicit //CHECK: IntelFPGAPrivateCopiesAttr + //CHECK-NEXT: ConstantExpr + //CHECK-NEXT: value:{{.*}}8 //CHECK-NEXT: SubstNonTypeTemplateParmExpr //CHECK-NEXT: NonTypeTemplateParmDecl //CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp index 38618d369339f..c2492a2f96c67 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -12,6 +12,8 @@ template int main() { // CHECK: FunctionDecl {{.*}}test_kernel1 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 5 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5 // expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}} // expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}} @@ -20,6 +22,8 @@ int main() { // CHECK: FunctionDecl {{.*}}test_kernel2 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 2 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 cl::sycl::kernel_single_task( []() { func(); }); From 22753fee8ec18d8996121002a6b01e769406e492 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 03:44:55 -0800 Subject: [PATCH 06/10] Remove dead code Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 7 ------- 1 file changed, 7 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index fe584533ba9a7..76ff86428fc97 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13176,13 +13176,6 @@ void Sema::AddOneConstantValueAttr(Decl *D, const AttributeCommonInfo &CI, return; E = ICE.get(); } - - if (IntelFPGAPrivateCopiesAttr::classof(&TmpAttr)) { - if (!D->hasAttr()) - D->addAttr(IntelFPGAMemoryAttr::CreateImplicit( - Context, IntelFPGAMemoryAttr::Default)); - } - D->addAttr(::new (Context) AttrType(Context, CI, E)); } From 6e4e9247dd867795be3f439e63fdbd44d35517ef Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 04:03:35 -0800 Subject: [PATCH 07/10] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp | 2 +- clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 9c888889dc6b6..ae3e7da18f014 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -588,7 +588,7 @@ static void instantiateIntelFPGAMemoryAttr( return S.AddOneConstantPowerTwoValueAttr(New, *Attr, Result.getAs()); - else if (std::is_same::value) + if (std::is_same::value) return S.addIntelSingleArgFunctionAttr(New, *Attr, Result.getAs()); diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp index c2492a2f96c67..cf5911c1b35a4 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -10,7 +10,7 @@ template [[intel::scheduler_target_fmax_mhz(N)]] void zoo() {} int main() { - // CHECK: FunctionDecl {{.*}}test_kernel1 'void ()' + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 5 @@ -20,7 +20,7 @@ int main() { cl::sycl::kernel_single_task( []() [[intelfpga::scheduler_target_fmax_mhz(5)]]{}); - // CHECK: FunctionDecl {{.*}}test_kernel2 'void ()' + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()' // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 2 From 65e683c9f6cac63df68ac3fb0b2534ae18aea597 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 06:44:52 -0800 Subject: [PATCH 08/10] Remove arbitrary range for remaining attributes Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 32 ------------------- clang/include/clang/Sema/Sema.h | 24 ++++++++------ clang/lib/Sema/SemaDeclAttr.cpp | 19 +++++------ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 3 +- clang/test/SemaSYCL/intel-fpga-local.cpp | 16 +++++----- 5 files changed, 34 insertions(+), 60 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 8d8a92a94d723..5935897980b2d 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1998,14 +1998,6 @@ def IntelFPGABankWidth : Attr { Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Documentation = [IntelFPGABankWidthAttrDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 1; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def IntelFPGANumBanks : Attr { @@ -2016,14 +2008,6 @@ def IntelFPGANumBanks : Attr { Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Documentation = [IntelFPGANumBanksAttrDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 1; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def IntelFPGAPrivateCopies : InheritableAttr { @@ -2054,14 +2038,6 @@ def IntelFPGAMaxReplicates : Attr { Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Documentation = [IntelFPGAMaxReplicatesAttrDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 1; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def IntelFPGASimpleDualPort : Attr { @@ -2097,14 +2073,6 @@ def IntelFPGABankBits : Attr { Field], ErrorDiag>; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Documentation = [IntelFPGABankBitsDocs]; - let AdditionalMembers = [{ - static unsigned getMinValue() { - return 0; - } - static unsigned getMaxValue() { - return 1024*1024; - } - }]; } def IntelFPGAForcePow2Depth : Attr { diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 76ff86428fc97..cc5f386fac26a 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13071,22 +13071,23 @@ void Sema::addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, E = ICE.get(); int32_t ArgInt = ArgVal.getSExtValue(); if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNumSimdWorkItems || - CI.getParsedKind() == ParsedAttr::AT_IntelReqdSubGroupSize) { + CI.getParsedKind() == ParsedAttr::AT_IntelReqdSubGroupSize || + CI.getParsedKind() == ParsedAttr::AT_IntelFPGAMaxReplicates) { if (ArgInt <= 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << CI.getAttrName() << /*positive*/ 0; + << CI << /*positive*/ 0; return; } } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << CI.getAttrName() << /*non-negative*/ 1; + << CI << /*non-negative*/ 1; return; } if (ArgInt > 3) { Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range) - << CI.getAttrName() << 0 << 3 << E->getSourceRange(); + << CI << 0 << 3 << E->getSourceRange(); return; } } @@ -13094,7 +13095,7 @@ void Sema::addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << CI.getAttrName() << /*non-negative*/ 1; + << CI << /*non-negative*/ 1; return; } } @@ -13186,12 +13187,15 @@ void Sema::AddOneConstantPowerTwoValueAttr(Decl *D, AttrType TmpAttr(Context, CI, E); if (!E->isValueDependent()) { - ExprResult ICE; - if (checkRangedIntegralArgument(E, &TmpAttr, ICE)) + llvm::APSInt Value; + ExprResult ICE = VerifyIntegerConstantExpression(E, &Value); + if (ICE.isInvalid()) return; - Expr::EvalResult Result; - E->EvaluateAsInt(Result, Context); - llvm::APSInt Value = Result.Val.getInt(); + if (!Value.isStrictlyPositive()) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*positive*/ 0; + return; + } if (!Value.isPowerOf2()) { Diag(CI.getLoc(), diag::err_attribute_argument_not_power_of_two) << &TmpAttr; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 0524b182d178b..74e84e7cc3400 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5720,9 +5720,7 @@ static void handleOneConstantPowerTwoValueAttr(Sema &S, Decl *D, if (checkAttrMutualExclusion(S, D, A)) return; - if (A.getKind() == ParsedAttr::AT_IntelFPGABankWidth || - A.getKind() == ParsedAttr::AT_IntelFPGANumBanks) - S.CheckDeprecatedSYCLAttributeSpelling(A); + S.CheckDeprecatedSYCLAttributeSpelling(A); S.AddOneConstantPowerTwoValueAttr(D, A, A.getArgAsExpr(0)); } @@ -5759,8 +5757,8 @@ static void handleIntelFPGAMaxReplicatesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.AddOneConstantValueAttr(D, A, - A.getArgAsExpr(0)); + S.addIntelSingleArgFunctionAttr( + D, A, A.getArgAsExpr(0)); } /// Handle the merge attribute. @@ -5837,11 +5835,14 @@ void Sema::AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr::EvalResult Result; ListIsValueDep = ListIsValueDep || E->isValueDependent(); if (!E->isValueDependent()) { - ExprResult ICE; - if (checkRangedIntegralArgument(E, &TmpAttr, ICE)) + ExprResult ICE = VerifyIntegerConstantExpression(E, &Value); + if (ICE.isInvalid()) return; - if (E->EvaluateAsInt(Result, Context)) - Value = Result.Val.getInt(); + if (!Value.isNonNegative()) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; + return; + } E = ICE.get(); } Args.push_back(E); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index ae3e7da18f014..cba4f97e31e6f 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -588,7 +588,8 @@ static void instantiateIntelFPGAMemoryAttr( return S.AddOneConstantPowerTwoValueAttr(New, *Attr, Result.getAs()); - if (std::is_same::value) + if (std::is_same::value || + std::is_same::value) return S.addIntelSingleArgFunctionAttr(New, *Attr, Result.getAs()); diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index 49878559d012a..796b789c5ab55 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -314,9 +314,9 @@ void diagnostics() //expected-note@+1 {{did you mean to use 'intel::max_replicates' instead?}} [[intelfpga::max_replicates(2)]] unsigned int max_replicates[64]; - //expected-error@+1{{'max_replicates' attribute requires integer constant between 1 and 1048576 inclusive}} + //expected-error@+1{{'max_replicates' attribute requires a positive integral compile time constant expression}} [[intel::max_replicates(0)]] unsigned int maxrepl_zero[64]; - //expected-error@+1{{'max_replicates' attribute requires integer constant between 1 and 1048576 inclusive}} + //expected-error@+1{{'max_replicates' attribute requires a positive integral compile time constant expression}} [[intel::max_replicates(-1)]] unsigned int maxrepl_negative[64]; //expected-error@+3{{'max_replicates' and 'fpga_register' attributes are not compatible}} @@ -356,7 +356,7 @@ void diagnostics() //expected-error@+1{{must be a constant power of two greater than zero}} [[intel::bankwidth(3)]] unsigned int bw_invalid_value[64]; - //expected-error@+1{{requires integer constant between 1 and 1048576}} + //expected-error@+1{{requires a positive integral compile time constant expression}} [[intel::bankwidth(-4)]] unsigned int bw_negative[64]; int i_bankwidth = 32; // expected-note {{declared here}} @@ -368,7 +368,7 @@ void diagnostics() //expected-error@+1{{'bankwidth' attribute takes one argument}} [[intel::bankwidth(4, 8)]] unsigned int bw_two_args[64]; - //expected-error@+1{{requires integer constant between 1 and 1048576}} + //expected-error@+1{{requires a positive integral compile time constant expression}} [[intel::bankwidth(0)]] unsigned int bw_zero[64]; // private_copies_ @@ -446,7 +446,7 @@ void diagnostics() //expected-error@+1{{must be a constant power of two greater than zero}} [[intel::numbanks(15)]] unsigned int nb_invalid_arg[64]; - //expected-error@+1{{requires integer constant between 1 and 1048576}} + //expected-error@+1{{requires a positive integral compile time constant expression}} [[intel::numbanks(-4)]] unsigned int nb_negative[64]; int i_numbanks = 32; // expected-note {{declared here}} @@ -458,7 +458,7 @@ void diagnostics() //expected-error@+1{{'numbanks' attribute takes one argument}} [[intel::numbanks(4, 8)]] unsigned int nb_two_args[64]; - //expected-error@+1{{requires integer constant between 1 and 1048576}} + //expected-error@+1{{requires a positive integral compile time constant expression}} [[intel::numbanks(0)]] unsigned int nb_zero[64]; // merge @@ -554,7 +554,7 @@ void diagnostics() //expected-error@+1{{attribute takes at least 1 argument}} [[intel::bank_bits]] unsigned int bb_no_arg[4]; - //expected-error@+1{{requires integer constant between 0 and 1048576}} + //expected-error@+1{{'bank_bits' attribute requires a non-negative integral compile time constant expression}} [[intel::bank_bits(-1)]] unsigned int bb_negative_arg[4]; // force_pow2_depth @@ -796,7 +796,7 @@ void check_template_parameters() { //expected-error@+1{{'numbanks' attribute takes one argument}} [[intel::numbanks(A, B)]] int numbanks_negative; - //expected-error@+1{{'max_replicates' attribute requires integer constant between 1 and 1048576}} + //expected-error@+1{{'max_replicates' attribute requires a positive integral compile time constant expression}} [[intel::max_replicates(D)]] [[intel::max_replicates(C)]] //expected-warning@-1{{attribute 'max_replicates' is already applied}} From d4e99833a2a6d0e2143fde683efd05b48b41bf83 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 08:12:21 -0800 Subject: [PATCH 09/10] Windows fail fix and implement review comments Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 18 ++++++++--------- clang/lib/Sema/SemaDeclAttr.cpp | 20 ++++++++----------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 14 ++++++------- 3 files changed, 23 insertions(+), 29 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index cc5f386fac26a..8a5f289d43724 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10206,12 +10206,10 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); template - void addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); + void addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E); template - void addIntelTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, - Expr *ZDimExpr); + void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); @@ -13059,8 +13057,8 @@ class Sema final { }; template -void Sema::addIntelSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E) { +void Sema::addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E) { assert(E && "Attribute must have an argument."); if (!E->isInstantiationDependent()) { @@ -13142,9 +13140,9 @@ static Expr *checkMaxWorkSizeAttrExpr(Sema &S, const AttributeCommonInfo &CI, } template -void Sema::addIntelTripleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, - Expr *ZDimExpr) { +void Sema::addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDimExpr, Expr *YDimExpr, + Expr *ZDimExpr) { assert((XDimExpr && YDimExpr && ZDimExpr) && "argument has unexpected null value"); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 74e84e7cc3400..7cf0344ba9153 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3150,8 +3150,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (!checkWorkGroupSizeValues(S, D, AL)) return; - S.addIntelTripleArgFunctionAttr(D, AL, XDimExpr, YDimExpr, - ZDimExpr); + S.addIntelTripleArgAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr); } // Handles work_group_size_hint. @@ -3193,7 +3192,7 @@ static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { if (D->getAttr()) S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; - S.addIntelSingleArgFunctionAttr(D, AL, E); + S.addIntelSingleArgAttr(D, AL, E); } // Handles num_simd_work_items. @@ -3208,7 +3207,7 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } // Handles use_stall_enable_clusters @@ -3239,8 +3238,7 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(AL); - S.addIntelSingleArgFunctionAttr(D, AL, - E); + S.addIntelSingleArgAttr(D, AL, E); } // Handles max_global_work_dim. @@ -3260,7 +3258,7 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } SYCLIntelLoopFuseAttr * @@ -5587,7 +5585,7 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, ? A.getArgAsExpr(0) : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, A.getLoc()); - S.addIntelSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. @@ -5757,8 +5755,7 @@ static void handleIntelFPGAMaxReplicatesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSingleArgFunctionAttr( - D, A, A.getArgAsExpr(0)); + S.addIntelSingleArgAttr(D, A, A.getArgAsExpr(0)); } /// Handle the merge attribute. @@ -5898,8 +5895,7 @@ static void handleIntelFPGAPrivateCopiesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSingleArgFunctionAttr( - D, A, A.getArgAsExpr(0)); + S.addIntelSingleArgAttr(D, A, A.getArgAsExpr(0)); } static void handleIntelFPGAForcePow2DepthAttr(Sema &S, Decl *D, diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index cba4f97e31e6f..9abf868021c3c 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -571,8 +571,7 @@ static void instantiateIntelSYCTripleLFunctionAttr( return; Expr *ZDimExpr = Result.getAs(); - S.addIntelTripleArgFunctionAttr(New, *Attr, XDimExpr, YDimExpr, - ZDimExpr); + S.addIntelTripleArgAttr(New, *Attr, XDimExpr, YDimExpr, ZDimExpr); } template @@ -590,11 +589,12 @@ static void instantiateIntelFPGAMemoryAttr( if (std::is_same::value || std::is_same::value) - return S.addIntelSingleArgFunctionAttr(New, *Attr, - Result.getAs()); - - return S.AddOneConstantValueAttr(New, *Attr, + return S.addIntelSingleArgAttr(New, *Attr, Result.getAs()); + + if (std::is_same::value) + return S.AddOneConstantValueAttr(New, *Attr, + Result.getAs()); } } @@ -642,7 +642,7 @@ static void instantiateIntelSYCLFunctionAttr( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); if (!Result.isInvalid()) - S.addIntelSingleArgFunctionAttr(New, *Attr, Result.getAs()); + S.addIntelSingleArgAttr(New, *Attr, Result.getAs()); } /// Determine whether the attribute A might be relevent to the declaration D. From 68570c2eeb928cfaa3afe89eee321baffa830539 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 12 Feb 2021 09:48:18 -0800 Subject: [PATCH 10/10] Fix MSVC build fail Signed-off-by: Elizabeth Andrews --- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 57 +++++++++++-------- 1 file changed, 33 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 9abf868021c3c..fb17e6ac68888 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -574,28 +574,37 @@ static void instantiateIntelSYCTripleLFunctionAttr( S.addIntelTripleArgAttr(New, *Attr, XDimExpr, YDimExpr, ZDimExpr); } -template -static void instantiateIntelFPGAMemoryAttr( +static void instantiateIntelFPGAForcePow2DepthAttr( Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, - const AttrName *Attr, Decl *New) { + const IntelFPGAForcePow2DepthAttr *Attr, Decl *New) { EnterExpressionEvaluationContext Unevaluated( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); - if (!Result.isInvalid()) { - if (std::is_same::value || - std::is_same::value) - return S.AddOneConstantPowerTwoValueAttr(New, *Attr, - Result.getAs()); + if (!Result.isInvalid()) + return S.AddOneConstantValueAttr( + New, *Attr, Result.getAs()); +} - if (std::is_same::value || - std::is_same::value) - return S.addIntelSingleArgAttr(New, *Attr, - Result.getAs()); +static void instantiateIntelFPGABankWidthAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const IntelFPGABankWidthAttr *Attr, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); + if (!Result.isInvalid()) + S.AddOneConstantPowerTwoValueAttr( + New, *Attr, Result.getAs()); +} - if (std::is_same::value) - return S.AddOneConstantValueAttr(New, *Attr, - Result.getAs()); - } +static void instantiateIntelFPGANumBanksAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const IntelFPGANumBanksAttr *Attr, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); + if (!Result.isInvalid()) + S.AddOneConstantPowerTwoValueAttr( + New, *Attr, Result.getAs()); } static void instantiateIntelFPGABankBitsAttr( @@ -790,23 +799,23 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, if (const auto *IntelFPGABankWidth = dyn_cast(TmplAttr)) { - instantiateIntelFPGAMemoryAttr( - *this, TemplateArgs, IntelFPGABankWidth, New); + instantiateIntelFPGABankWidthAttr(*this, TemplateArgs, IntelFPGABankWidth, + New); } if (const auto *IntelFPGANumBanks = dyn_cast(TmplAttr)) { - instantiateIntelFPGAMemoryAttr( - *this, TemplateArgs, IntelFPGANumBanks, New); + instantiateIntelFPGANumBanksAttr(*this, TemplateArgs, IntelFPGANumBanks, + New); } if (const auto *IntelFPGAPrivateCopies = dyn_cast(TmplAttr)) { - instantiateIntelFPGAMemoryAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, IntelFPGAPrivateCopies, New); } if (const auto *IntelFPGAMaxReplicates = dyn_cast(TmplAttr)) { - instantiateIntelFPGAMemoryAttr( + instantiateIntelSYCLFunctionAttr( *this, TemplateArgs, IntelFPGAMaxReplicates, New); } if (const auto *IntelFPGABankBits = @@ -816,8 +825,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *IntelFPGAForcePow2Depth = dyn_cast(TmplAttr)) { - instantiateIntelFPGAMemoryAttr( - *this, TemplateArgs, IntelFPGAForcePow2Depth, New); + instantiateIntelFPGAForcePow2DepthAttr(*this, TemplateArgs, + IntelFPGAForcePow2Depth, New); } if (const auto *SYCLIntelPipeIO = dyn_cast(TmplAttr)) { instantiateSYCLIntelPipeIOAttr(*this, TemplateArgs, SYCLIntelPipeIO, New);