diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 14d78e539412f..5935897980b2d 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1282,14 +1282,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 { @@ -1344,14 +1336,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 { @@ -2014,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 { @@ -2032,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 { @@ -2049,14 +2017,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. @@ -2078,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 { @@ -2121,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 eea057e8aeda8..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 addIntelSYCLSingleArgFunctionAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E); + void addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E); template - void addIntelSYCLTripleArgFunctionAttr(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,9 +13057,8 @@ class Sema final { }; template -void Sema::addIntelSYCLSingleArgFunctionAttr(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()) { @@ -13072,27 +13069,42 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, 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; + } + } + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz || + CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) { + if (ArgInt < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; return; } } } + if (CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) { + if (!D->hasAttr()) + D->addAttr(IntelFPGAMemoryAttr::CreateImplicit( + Context, IntelFPGAMemoryAttr::Default)); + } + D->addAttr(::new (Context) AttrType(Context, CI, E)); } @@ -13128,10 +13140,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::addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDimExpr, Expr *YDimExpr, + Expr *ZDimExpr) { assert((XDimExpr && YDimExpr && ZDimExpr) && "argument has unexpected null value"); @@ -13164,13 +13175,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)); } @@ -13181,12 +13185,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 233cc149d5a71..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.addIntelSYCLTripleArgFunctionAttr(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.addIntelSYCLSingleArgFunctionAttr(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.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } // Handles use_stall_enable_clusters @@ -3239,7 +3238,7 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(AL); - S.AddOneConstantValueAttr(D, AL, E); + S.addIntelSingleArgAttr(D, AL, E); } // Handles max_global_work_dim. @@ -3259,7 +3258,7 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - S.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } SYCLIntelLoopFuseAttr * @@ -3312,10 +3311,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, @@ -5583,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.addIntelSYCLSingleArgFunctionAttr(D, A, E); + S.addIntelSingleArgAttr(D, A, E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. @@ -5716,9 +5718,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)); } @@ -5755,8 +5755,7 @@ static void handleIntelFPGAMaxReplicatesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.AddOneConstantValueAttr(D, A, - A.getArgAsExpr(0)); + S.addIntelSingleArgAttr(D, A, A.getArgAsExpr(0)); } /// Handle the merge attribute. @@ -5833,11 +5832,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 (!Value.isNonNegative()) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; return; - if (E->EvaluateAsInt(Result, Context)) - Value = Result.Val.getInt(); + } E = ICE.get(); } Args.push_back(E); @@ -5893,8 +5895,7 @@ static void handleIntelFPGAPrivateCopiesAttr(Sema &S, Decl *D, S.CheckDeprecatedSYCLAttributeSpelling(A); - S.AddOneConstantValueAttr(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 f83bda22e1717..fb17e6ac68888 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -571,25 +571,40 @@ static void instantiateIntelSYCTripleLFunctionAttr( return; Expr *ZDimExpr = Result.getAs(); - S.addIntelSYCLTripleArgFunctionAttr(New, *Attr, XDimExpr, YDimExpr, - ZDimExpr); + 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()); - return S.AddOneConstantValueAttr(New, *Attr, - Result.getAs()); - } + if (!Result.isInvalid()) + return S.AddOneConstantValueAttr( + 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()); +} + +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( @@ -636,8 +651,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.addIntelSingleArgAttr(New, *Attr, Result.getAs()); } /// Determine whether the attribute A might be relevent to the declaration D. @@ -785,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 = @@ -811,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); diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index 08f914f06a999..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_ @@ -401,7 +401,7 @@ void diagnostics() [[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}} @@ -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}} 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 bd6e2703d3d1f..cf5911c1b35a4 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -41,10 +41,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}}