diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 769d170880362..03e86a36bf475 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11346,7 +11346,7 @@ def err_conflicting_sycl_function_attributes : Error< def err_sycl_function_attribute_mismatch : Error< "SYCL kernel without %0 attribute can't call a function with this attribute">; def err_sycl_x_y_z_arguments_must_be_one : Error< - "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">; + "all %0 attribute arguments must be '1' when the %1 attribute argument is '0'">; def err_sycl_attribute_internal_function : Error<"%0 attribute cannot be applied to a " "static function or function in an anonymous namespace">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 16776b85bf7cf..b1db79c97da04 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10302,8 +10302,6 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); template - void addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E); - template void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr); void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI, @@ -10359,7 +10357,11 @@ class Sema final { SYCLIntelFPGAMaxConcurrencyAttr *MergeSYCLIntelFPGAMaxConcurrencyAttr( Decl *D, const SYCLIntelFPGAMaxConcurrencyAttr &A); - + void AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *E); + SYCLIntelMaxGlobalWorkDimAttr * + MergeSYCLIntelMaxGlobalWorkDimAttr(Decl *D, + const SYCLIntelMaxGlobalWorkDimAttr &A); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); @@ -13332,37 +13334,6 @@ class Sema final { } }; -template -void Sema::addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *E) { - assert(E && "Attribute must have an argument."); - - if (!E->isInstantiationDependent()) { - llvm::APSInt ArgVal; - ExprResult ICE = VerifyIntegerConstantExpression(E, &ArgVal); - if (ICE.isInvalid()) - return; - E = ICE.get(); - int32_t ArgInt = ArgVal.getSExtValue(); - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { - if (ArgInt < 0) { - Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << CI << /*non-negative*/ 1; - return; - } - } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { - if (ArgInt > 3) { - Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range) - << CI << 0 << 3 << E->getSourceRange(); - return; - } - } - } - - D->addAttr(::new (Context) AttrType(Context, CI, E)); -} - inline Expr *checkMaxWorkSizeAttrExpr(Sema &S, const AttributeCommonInfo &CI, Expr *E) { assert(E && "Attribute must have an argument."); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 3a84694fb62fd..77202024d674a 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2643,6 +2643,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, NewAttr = S.MergeSYCLIntelFPGAInitiationIntervalAttr(D, *A); else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeWorkGroupSizeHintAttr(D, *A); + else if (const auto *A = dyn_cast(Attr)) + NewAttr = S.MergeSYCLIntelMaxGlobalWorkDimAttr(D, *A); else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr)) NewAttr = cast(Attr->clone(S.Context)); diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 4be847f4b1aa5..f310567cb8f08 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2989,19 +2989,6 @@ static void handleWeakImportAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // they hold equal values (1, 1, 1). static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { bool Result = true; - auto checkZeroDim = [&S, &AL](auto &A, size_t X, size_t Y, size_t Z, - bool ReverseAttrs = false) -> bool { - if (X != 1 || Y != 1 || Z != 1) { - auto Diag = - S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one); - if (ReverseAttrs) - Diag << AL << A; - else - Diag << A << AL; - return false; - } - return true; - }; // Returns the unsigned constant integer value represented by // given expression. @@ -3011,30 +2998,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { ASTContext &Ctx = S.getASTContext(); - if (AL.getKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { - ArrayRef Dims; - Attr *B = nullptr; - if (const auto *B = D->getAttr()) - Dims = B->dimensions(); - else if (const auto *B = D->getAttr()) - Dims = B->dimensions(); - if (B) { - Result &= - checkZeroDim(B, getExprValue(Dims[0], Ctx), - getExprValue(Dims[1], Ctx), getExprValue(Dims[2], Ctx)); - } - return Result; - } - - if (const auto *A = D->getAttr()) { - if ((A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue()) == 0) { - Result &= checkZeroDim(A, getExprValue(AL.getArgAsExpr(0), Ctx), - getExprValue(AL.getArgAsExpr(1), Ctx), - getExprValue(AL.getArgAsExpr(2), Ctx), - /*ReverseAttrs=*/true); - } - } - if (const auto *A = D->getAttr()) { if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <= getExprValue(A->getXDim(), Ctx)) && @@ -3144,6 +3107,27 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { } } + // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or + // ReqdWorkGroupSizeAttr, check to see if they hold equal values + // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr + // equals to 0. + if (const auto *DeclAttr = D->getAttr()) { + if (const auto *DeclExpr = dyn_cast(DeclAttr->getValue())) { + // If the value is dependent, we can not test anything. + if (!DeclExpr) + return; + + // Test the attribute value. + if (DeclExpr->getResultAsAPSInt() == 0 && + (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 || + ZDimVal.getZExtValue() != 1)) { + S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one) + << AL << DeclAttr; + return; + } + } + } + if (const auto *ExistingAttr = D->getAttr()) { // Compare attribute arguments value and warn for a mismatch. if (ExistingAttr->getXDimVal(Ctx) != XDimVal || @@ -3646,23 +3630,130 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, } // Handles max_global_work_dim. -static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &A) { - if (D->isInvalidDecl()) - return; +// Returns a OneArgResult value; EqualToOne means all argument values are +// equal to one, NotEqualToOne means at least one argument value is not +// equal to one, and Unknown means that at least one of the argument values +// could not be determined. +enum class OneArgResult { Unknown, EqualToOne, NotEqualToOne }; +static OneArgResult AreAllArgsOne(const Expr *Args[], size_t Count) { + + for (size_t Idx = 0; Idx < Count; ++Idx) { + const auto *CE = dyn_cast(Args[Idx]); + if (!CE) + return OneArgResult::Unknown; + if (CE->getResultAsAPSInt() != 1) + return OneArgResult::NotEqualToOne; + } + return OneArgResult::EqualToOne; +} + +// If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or +// ReqdWorkGroupSizeAttr, check to see if they hold equal values +// (1, 1, 1). Returns true if diagnosed. +template +static bool checkWorkGroupSizeAttrExpr(Sema &S, Decl *D, + const AttributeCommonInfo &AL) { + if (const auto *A = D->getAttr()) { + const Expr *Args[3] = {A->getXDim(), A->getYDim(), A->getZDim()}; + if (OneArgResult::NotEqualToOne == AreAllArgsOne(Args, 3)) { + S.Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one) + << A << AL; + return true; + } + } + return false; +} - Expr *E = A.getArgAsExpr(0); +void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *E) { + if (!E->isValueDependent()) { + // Validate that we have an integer constant expression and then store the + // converted constant expression into the semantic attribute so that we + // don't have to evaluate it again later. + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return; + E = Res.get(); - if (!checkWorkGroupSizeValues(S, D, A)) { - D->setInvalidDecl(); - return; + // This attribute must be in the range [0, 3]. + if (ArgVal < 0 || ArgVal > 3) { + Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range) + << CI << 0 << 3 << E->getSourceRange(); + return; + } + + // Check to see if there's a duplicate attribute with different values + // already applied to the declaration. + if (const auto *DeclAttr = D->getAttr()) { + // If the other attribute argument is instantiation dependent, we won't + // have converted it to a constant expression yet and thus we test + // whether this is a null pointer. + if (const auto *DeclExpr = dyn_cast(DeclAttr->getValue())) { + if (ArgVal != DeclExpr->getResultAsAPSInt()) { + Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI; + Diag(DeclAttr->getLoc(), diag::note_previous_attribute); + } + // Drop the duplicate attribute. + return; + } + } + + // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or + // ReqdWorkGroupSizeAttr, check to see if they hold equal values + // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr + // equals to 0. + if (ArgVal == 0) { + if (checkWorkGroupSizeAttrExpr(*this, D, + CI) || + checkWorkGroupSizeAttrExpr(*this, D, CI)) + return; + } } - if (D->getAttr()) - S.Diag(A.getLoc(), diag::warn_duplicate_attribute) << A; + D->addAttr(::new (Context) SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E)); +} - S.CheckDeprecatedSYCLAttributeSpelling(A); +SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr( + Decl *D, const SYCLIntelMaxGlobalWorkDimAttr &A) { + // Check to see if there's a duplicate attribute with different values + // already applied to the declaration. + if (const auto *DeclAttr = D->getAttr()) { + if (const auto *DeclExpr = dyn_cast(DeclAttr->getValue())) { + if (const auto *MergeExpr = dyn_cast(A.getValue())) { + if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) { + Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A; + Diag(A.getLoc(), diag::note_previous_attribute); + } + // Do not add a duplicate attribute. + return nullptr; + } + } + } + + // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or + // ReqdWorkGroupSizeAttr, check to see if they hold equal values + // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr + // equals to 0. + const auto *MergeExpr = dyn_cast(A.getValue()); + if (MergeExpr->getResultAsAPSInt() == 0) { + if (checkWorkGroupSizeAttrExpr(*this, D, + A) || + checkWorkGroupSizeAttrExpr(*this, D, A)) + return nullptr; + } + + return ::new (Context) + SYCLIntelMaxGlobalWorkDimAttr(Context, A, A.getValue()); +} - S.addIntelSingleArgAttr(D, A, E); +static void handleSYCLIntelMaxGlobalWorkDimAttr(Sema &S, Decl *D, + const ParsedAttr &AL) { + S.CheckDeprecatedSYCLAttributeSpelling(AL); + + Expr *E = AL.getArgAsExpr(0); + S.AddSYCLIntelMaxGlobalWorkDimAttr(D, AL, E); } // Handles [[intel::loop_fuse]] and [[intel::loop_fuse_independent]]. @@ -9543,7 +9634,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, handleSYCLIntelSchedulerTargetFmaxMhzAttr(S, D, AL); break; case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim: - handleMaxGlobalWorkDimAttr(S, D, AL); + handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL); break; case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset: handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 8245de5877830..3516762f35e7c 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -682,15 +682,14 @@ static void instantiateSYCLIntelNoGlobalWorkOffsetAttr( S.AddSYCLIntelNoGlobalWorkOffsetAttr(New, *A, Result.getAs()); } -template -static void instantiateIntelSYCLFunctionAttr( +static void instantiateSYCLIntelMaxGlobalWorkDimAttr( Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, - const AttrName *Attr, Decl *New) { + const SYCLIntelMaxGlobalWorkDimAttr *A, Decl *New) { EnterExpressionEvaluationContext Unevaluated( S, Sema::ExpressionEvaluationContext::ConstantEvaluated); - ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs); + ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs); if (!Result.isInvalid()) - S.addIntelSingleArgAttr(New, *Attr, Result.getAs()); + S.AddSYCLIntelMaxGlobalWorkDimAttr(New, *A, Result.getAs()); } static void instantiateSYCLIntelFPGAMaxConcurrencyAttr( @@ -960,8 +959,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *SYCLIntelMaxGlobalWorkDim = dyn_cast(TmplAttr)) { - instantiateIntelSYCLFunctionAttr( - *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New); + instantiateSYCLIntelMaxGlobalWorkDimAttr(*this, TemplateArgs, + SYCLIntelMaxGlobalWorkDim, New); continue; } if (const auto *SYCLIntelLoopFuse = diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp index c07c12b5d98d8..29cc7f084e2d4 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -33,6 +33,25 @@ struct Func { [[intelfpga::max_global_work_dim(2)]] void operator()() const {} }; +// No diagnostic is emitted because the arguments match. +[[intel::max_global_work_dim(1)]] void bar(); +[[intel::max_global_work_dim(1)]] void bar() {} + +// Checking of different argument values. +[[intel::max_global_work_dim(2)]] void baz(); // expected-note {{previous attribute is here}} +[[intel::max_global_work_dim(1)]] void baz(); // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}} + +struct TRIFuncObj { + [[intel::max_global_work_dim(0)]] void operator()() const; // expected-note {{previous attribute is here}} +}; +[[intel::max_global_work_dim(1)]] void TRIFuncObj::operator()() const {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}} + +// Checks correctness of mutual usage of different work_group_size attributes: +// reqd_work_group_size, max_work_group_size, and max_global_work_dim. +// In case the value of 'max_global_work_dim' attribute equals to 0 we shall +// ensure that if max_work_group_size and reqd_work_group_size attributes exist, +// they hold equal values (1, 1, 1). + struct TRIFuncObjGood1 { [[intel::max_global_work_dim(0)]] [[intel::max_work_group_size(1, 1, 1)]] @@ -47,14 +66,169 @@ struct TRIFuncObjGood2 { operator()() const {} }; -#ifdef TRIGGER_ERROR +struct TRIFuncObjGood3 { + [[intel::reqd_work_group_size(1)]] + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjGood4 { + [[sycl::reqd_work_group_size(1, 1, 1)]] + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjGood5 { + [[intel::max_work_group_size(1, 1, 1)]] + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjGood6 { + [[intel::reqd_work_group_size(4, 1, 1)]] + [[intel::max_global_work_dim(3)]] void + operator()() const {} +}; + +struct TRIFuncObjGood7 { + [[sycl::reqd_work_group_size(4, 1, 1)]] + [[intel::max_global_work_dim(3)]] void + operator()() const {} +}; + +struct TRIFuncObjGood8 { + [[intel::max_work_group_size(8, 1, 1)]] + [[intel::max_global_work_dim(3)]] void + operator()() const {} +}; + +// FIXME: We do not have support yet for checking +// max_work_group_size and max_global_work_dim +// attributes when merging, so the test compiles without +// any diagnostic when it shouldn't. struct TRIFuncObjBad { + [[intel::max_work_group_size(4, 4, 4)]] void + operator()() const; +}; + +[[intel::max_global_work_dim(0)]] +void TRIFuncObjBad::operator()() const {} + +// FIXME: We do not have support yet for checking +// max_work_group_size and max_global_work_dim +// attributes when merging, so the test compiles without +// any diagnostic when it shouldn't. +struct TRIFuncObjBad1 { + [[intel::reqd_work_group_size(4, 4, 4)]] void + operator()() const; +}; + +[[intel::max_global_work_dim(0)]] +void TRIFuncObjBad1::operator()() const {} + +// FIXME: We do not have support yet for checking +// max_work_group_size and max_global_work_dim +// attributes when merging, so the test compiles without +// any diagnostic when it shouldn't. +struct TRIFuncObjBad2 { + [[sycl::reqd_work_group_size(4, 4, 4)]] void + operator()() const; +}; + +[[intel::max_global_work_dim(0)]] +void TRIFuncObjBad2::operator()() const {} + +#ifdef TRIGGER_ERROR +// Checks correctness of mutual usage of different work_group_size attributes: +// reqd_work_group_size, max_work_group_size and max_global_work_dim. +// In case the value of 'max_global_work_dim' attribute equals to 0 we shall +// ensure that if max_work_group_size and reqd_work_group_size attributes exist, +// they hold equal values (1, 1, 1). + +struct TRIFuncObjBad3 { [[intel::max_global_work_dim(0)]] - [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}} - [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}} + [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} void operator()() const {} }; + +struct TRIFuncObjBad4 { + [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad5 { + [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad6 { + [[intel::reqd_work_group_size(4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad7 { + [[intel::max_global_work_dim(0)]] void + operator()() const; +}; + +[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +void TRIFuncObjBad7::operator()() const {} + +struct TRIFuncObjBad8 { + [[intel::max_global_work_dim(0)]] void + operator()() const; +}; + +[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +void TRIFuncObjBad8::operator()() const {} + +struct TRIFuncObjBad9 { + [[intel::max_global_work_dim(0)]] void + operator()() const; +}; + +[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +void TRIFuncObjBad9::operator()() const {} + +// Tests for incorrect argument values for Intel FPGA function attributes: +// reqd_work_group_size, max_work_group_size and max_global_work_dim. + +struct TRIFuncObjBad10 { + // expected-error@+2{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + [[intel::reqd_work_group_size(-4, 1)]] + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad11 { + [[intel::max_work_group_size(4, 4, 4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad12 { + [[sycl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} + [[intel::max_global_work_dim(0)]] void + operator()() const {} +}; + +struct TRIFuncObjBad13 { + [[intel::reqd_work_group_size(4)]] + [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} + void operator()() const {} +}; + +struct TRIFuncObjBad14 { + [[intel::max_work_group_size(4, 4, 4)]] + [[intel::max_global_work_dim(4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + void operator()() const {} +}; #endif // TRIGGER_ERROR int main() { @@ -137,19 +311,194 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + + h.single_task(TRIFuncObjGood3()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + h.single_task(TRIFuncObjGood4()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel6 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + h.single_task(TRIFuncObjGood5()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel7 + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + h.single_task(TRIFuncObjGood6()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel8 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + + h.single_task(TRIFuncObjGood7()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + + h.single_task(TRIFuncObjGood8()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10 + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + + h.single_task(TRIFuncObjBad()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + h.single_task(TRIFuncObjBad1()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + h.single_task(TRIFuncObjBad2()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel13 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + + // Ignore duplicate attribute with same argument value. + h.single_task( + // CHECK-LABEL: FunctionDecl {{.*}}test_kernell4 'void ()' + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + []() [[intel::max_global_work_dim(3), + intel::max_global_work_dim(3)]]{}); // Ok + #ifdef TRIGGER_ERROR [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}} - h.single_task( - []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}} + h.single_task( + []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} - h.single_task( - []() [[intel::max_global_work_dim(3), + h.single_task( + []() [[intel::max_global_work_dim(3), // expected-note {{previous attribute is here}} intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}} - h.single_task(TRIFuncObjBad()); + h.single_task(TRIFuncObjBad3()); + h.single_task(TRIFuncObjBad4()); + h.single_task(TRIFuncObjBad5()); + h.single_task(TRIFuncObjBad6()); + h.single_task(TRIFuncObjBad7()); + h.single_task(TRIFuncObjBad8()); + h.single_task(TRIFuncObjBad9()); + h.single_task(TRIFuncObjBad10()); + h.single_task(TRIFuncObjBad11()); + h.single_task(TRIFuncObjBad12()); + h.single_task(TRIFuncObjBad13()); + h.single_task(TRIFuncObjBad14()); - h.single_task( + h.single_task( []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} #endif // TRIGGER_ERROR }); diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp index 95e7a13797ab7..41bb2ce80dde5 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp @@ -34,7 +34,7 @@ constexpr int bar() { return 0; } template class KernelFunctor { public: - // expected-error@+1{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}} + // expected-error@+1{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} [[intel::max_global_work_dim(SIZE)]] void operator()() {} }; @@ -58,17 +58,80 @@ int main() { // Test that checks template parameter support on function. template -// expected-error@+1{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}} +// expected-error@+1{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} [[intel::max_global_work_dim(N)]] void func3() {} +// Test that checks template instantiations for different argument values. +template +[[intel::max_global_work_dim(1)]] void func4(); // expected-note {{previous attribute is here}} + +template +[[intel::max_global_work_dim(size)]] void func4() {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}} + +// Checks correctness of mutual usage of different work_group_size attributes: +// reqd_work_group_size, max_work_group_size, and max_global_work_dim. +// In case the value of 'max_global_work_dim' attribute equals to 0 we shall +// ensure that if max_work_group_size and reqd_work_group_size attributes exist, +// they hold equal values (1, 1, 1). +template +[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +template +[[intel::max_global_work_dim(0)]] void func5(); + +template +[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +template +[[intel::max_global_work_dim(0)]] void func6(); + +template +[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +template +[[intel::max_global_work_dim(0)]] void func7(); + +template +[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +template +[[intel::max_global_work_dim(0)]] void func8(); + +template +[[intel::max_work_group_size(N, N, N)]] void func9(); +template +[[intel::max_global_work_dim(0)]] void func9(); + +template +[[intel::reqd_work_group_size(N)]] void func10(); +template +[[intel::max_global_work_dim(0)]] void func10(); + +template +[[intel::reqd_work_group_size(N, N)]] void func11(); +template +[[intel::max_global_work_dim(0)]] void func11(); + +template +[[sycl::reqd_work_group_size(N, N, N)]] void func12(); +template +[[intel::max_global_work_dim(0)]] void func12(); + int check() { - // no error expected - func3<3>(); - //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}} - func3<-1>(); + func3<3>(); // OK + func3<-1>(); // expected-note {{in instantiation of function template specialization 'func3<-1>' requested here}} + func4<2>(); // expected-note {{in instantiation of function template specialization 'func4<2>' requested here}} + func5<2>(); // expected-note {{in instantiation of function template specialization 'func5<2>' requested here}} + func6<2>(); // expected-note {{in instantiation of function template specialization 'func6<2>' requested here}} + func7<2>(); // expected-note {{in instantiation of function template specialization 'func7<2>' requested here}} + func8<2>(); // expected-note {{in instantiation of function template specialization 'func8<2>' requested here}} + func9<1>(); // OK + func10<1>(); // OK + func11<1>(); // OK + func12<1>(); // OK return 0; } +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +[[intel::max_global_work_dim(2)]] +[[intel::max_global_work_dim(2)]] void func13() {} + // CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' // CHECK: TemplateArgument integral 3 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} @@ -77,3 +140,9 @@ int check() { // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + +// CHECK: FunctionDecl {{.*}} {{.*}} func13 'void ()' +// CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} +// CHECK-NEXT: ConstantExpr {{.*}} 'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}