diff --git a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp index 521126f990e97..87bf2dd622880 100644 --- a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp +++ b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp @@ -57,9 +57,8 @@ void SingleWorkItemBarrierCheck::check(const MatchFinder::MatchResult &Result) { bool IsNDRange = false; if (MatchedDecl->hasAttr()) { const auto *Attribute = MatchedDecl->getAttr(); - if (*Attribute->getXDimVal(*Result.Context) > 1 || - *Attribute->getYDimVal(*Result.Context) > 1 || - *Attribute->getZDimVal(*Result.Context) > 1) + if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 || + *Attribute->getZDimVal() > 1) IsNDRange = true; } if (IsNDRange) // No warning if kernel is treated as an NDRange. diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 9eeba021fe2a9..ac14b6dd4248a 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -3251,17 +3251,20 @@ def ReqdWorkGroupSize : InheritableAttr { ExprArgument<"ZDim", /*optional*/1>]; let Subjects = SubjectList<[Function], ErrorDiag>; let AdditionalMembers = [{ - ArrayRef dimensions() const { - return {getXDim(), getYDim(), getZDim()}; - } - Optional getXDimVal(ASTContext &Ctx) const { - return getXDim()->getIntegerConstantExpr(Ctx); + Optional getXDimVal() const { + if (const auto *CE = dyn_cast(getXDim())) + return CE->getResultAsAPSInt(); + return None; } - Optional getYDimVal(ASTContext &Ctx) const { - return getYDim()->getIntegerConstantExpr(Ctx); + Optional getYDimVal() const { + if (const auto *CE = dyn_cast(getYDim())) + return CE->getResultAsAPSInt(); + return None; } - Optional getZDimVal(ASTContext &Ctx) const { - return getZDim()->getIntegerConstantExpr(Ctx); + Optional getZDimVal() const { + if (const auto *CE = dyn_cast(getZDim())) + return CE->getResultAsAPSInt(); + return None; } }]; let Documentation = [ReqdWorkGroupSizeAttrDocs]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index e0116feddc392..3c2e0a863d5ad 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10541,9 +10541,6 @@ class Sema final { void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); - template - void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr); void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XDim, Expr *YDim, Expr *ZDim); WorkGroupSizeHintAttr * @@ -10640,6 +10637,10 @@ class Sema final { void AddSYCLAddIRAttributesGlobalVariableAttr(Decl *D, const AttributeCommonInfo &CI, MutableArrayRef Args); + void AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDim, Expr *YDim, Expr *ZDim); + ReqdWorkGroupSizeAttr * + MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); @@ -13736,62 +13737,6 @@ class Sema final { } }; -inline Expr *checkMaxWorkSizeAttrExpr(Sema &S, const AttributeCommonInfo &CI, - Expr *E) { - assert(E && "Attribute must have an argument."); - - if (!E->isInstantiationDependent()) { - llvm::APSInt ArgVal; - ExprResult ICE = S.VerifyIntegerConstantExpression(E, &ArgVal); - - if (ICE.isInvalid()) - return nullptr; - - E = ICE.get(); - - if (ArgVal.isNegative()) { - S.Diag(E->getExprLoc(), - diag::warn_attribute_requires_non_negative_integer_argument) - << E->getType() << S.Context.UnsignedLongLongTy - << E->getSourceRange(); - return E; - } - - unsigned Val = ArgVal.getZExtValue(); - if (Val == 0) { - S.Diag(E->getExprLoc(), diag::err_attribute_argument_is_zero) - << CI << E->getSourceRange(); - return nullptr; - } - } - return E; -} - -template -void Sema::addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI, - Expr *XDimExpr, Expr *YDimExpr, - Expr *ZDimExpr) { - - assert((XDimExpr && YDimExpr && ZDimExpr) && - "argument has unexpected null value"); - - // Accept template arguments for now as they depend on something else. - // We'll get to check them when they eventually get instantiated. - if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && - !ZDimExpr->isValueDependent()) { - - // Save ConstantExpr in semantic attribute - XDimExpr = checkMaxWorkSizeAttrExpr(*this, CI, XDimExpr); - YDimExpr = checkMaxWorkSizeAttrExpr(*this, CI, YDimExpr); - ZDimExpr = checkMaxWorkSizeAttrExpr(*this, CI, ZDimExpr); - - if (!XDimExpr || !YDimExpr || !ZDimExpr) - return; - } - D->addAttr(::new (Context) - WorkGroupAttrType(Context, CI, XDimExpr, YDimExpr, ZDimExpr)); -} - /// RAII object that enters a new expression evaluation context. class EnterExpressionEvaluationContext { Sema &Actions; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index de6c39f72272e..900c485173bce 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -636,22 +636,13 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const ReqdWorkGroupSizeAttr *A = FD->getAttr()) { - ASTContext &ClangCtx = FD->getASTContext(); - Optional XDimVal = A->getXDimVal(ClangCtx); - Optional YDimVal = A->getYDimVal(ClangCtx); - Optional ZDimVal = A->getZDimVal(ClangCtx); - - // For a SYCLDevice ReqdWorkGroupSizeAttr arguments are reversed. - if (getLangOpts().SYCLIsDevice) - std::swap(XDimVal, ZDimVal); - + // Attributes arguments (first and third) are reversed on SYCLDevice. llvm::Metadata *AttrMDArgs[] = { - llvm::ConstantAsMetadata::get( - Builder.getInt32(XDimVal->getZExtValue())), - llvm::ConstantAsMetadata::get( - Builder.getInt32(YDimVal->getZExtValue())), - llvm::ConstantAsMetadata::get( - Builder.getInt32(ZDimVal->getZExtValue()))}; + llvm::ConstantAsMetadata::get(Builder.getInt( + getLangOpts().SYCLIsDevice ? *A->getZDimVal() : *A->getXDimVal())), + llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())), + llvm::ConstantAsMetadata::get(Builder.getInt( + getLangOpts().SYCLIsDevice ? *A->getXDimVal() : *A->getZDimVal()))}; Fn->setMetadata("reqd_work_group_size", llvm::MDNode::get(Context, AttrMDArgs)); } @@ -715,7 +706,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, const auto *CE = cast(A->getValue()); Optional ArgVal = CE->getResultAsAPSInt(); llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get( - Builder.getInt32(ArgVal->getSExtValue()))}; + Builder.getInt32(ArgVal->getZExtValue()))}; Fn->setMetadata("num_simd_work_items", llvm::MDNode::get(Context, AttrMDArgs)); } diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 2f7713d83769b..7bfda66e64db6 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -8372,10 +8372,9 @@ void TCETargetCodeGenInfo::setTargetAttributes( SmallVector Operands; Operands.push_back(llvm::ConstantAsMetadata::get(F)); - ASTContext &Ctx = M.getContext(); - unsigned XDim = Attr->getXDimVal(Ctx)->getZExtValue(); - unsigned YDim = Attr->getYDimVal(Ctx)->getZExtValue(); - unsigned ZDim = Attr->getZDimVal(Ctx)->getZExtValue(); + unsigned XDim = Attr->getXDimVal()->getZExtValue(); + unsigned YDim = Attr->getYDimVal()->getZExtValue(); + unsigned ZDim = Attr->getZDimVal()->getZExtValue(); Operands.push_back(llvm::ConstantAsMetadata::get( llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, XDim)))); @@ -9255,9 +9254,9 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( Max = FlatWGS->getMax()->EvaluateKnownConstInt(Ctx).getExtValue(); } if (ReqdWGS) { - XDim = ReqdWGS->getXDimVal(Ctx)->getZExtValue(); - YDim = ReqdWGS->getYDimVal(Ctx)->getZExtValue(); - ZDim = ReqdWGS->getZDimVal(Ctx)->getZExtValue(); + XDim = ReqdWGS->getXDimVal()->getZExtValue(); + YDim = ReqdWGS->getYDimVal()->getZExtValue(); + ZDim = ReqdWGS->getZDimVal()->getZExtValue(); } if (ReqdWGS && Min == 0 && Max == 0) Min = Max = XDim * YDim * ZDim; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index ba40ad380ff8b..f765ad84ca7c0 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2805,6 +2805,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeSYCLAddIRAttributesGlobalVariableAttr(D, *A); + else if (const auto *A = dyn_cast(Attr)) + NewAttr = S.MergeReqdWorkGroupSizeAttr(D, *A); else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr)) NewAttr = cast(Attr->clone(S.Context)); @@ -3395,27 +3397,6 @@ static void adjustDeclContextForDeclaratorDecl(DeclaratorDecl *NewD, FixSemaDC(VD->getDescribedVarTemplate()); } -template -static void checkDimensionsAndSetDiagnostics(Sema &S, FunctionDecl *New, - FunctionDecl *Old) { - const auto *NewDeclAttr = New->getAttr(); - const auto *OldDeclAttr = Old->getAttr(); - - if (!NewDeclAttr || !OldDeclAttr) - return; - - ASTContext &Ctx = S.getASTContext(); - if (NewDeclAttr->getXDimVal(Ctx) != OldDeclAttr->getXDimVal(Ctx) || - NewDeclAttr->getYDimVal(Ctx) != OldDeclAttr->getYDimVal(Ctx) || - NewDeclAttr->getZDimVal(Ctx) != OldDeclAttr->getZDimVal(Ctx)) { - S.Diag(New->getLocation(), diag::err_conflicting_sycl_function_attributes) - << OldDeclAttr << NewDeclAttr; - S.Diag(New->getLocation(), diag::warn_duplicate_attribute) << OldDeclAttr; - S.Diag(OldDeclAttr->getLocation(), diag::note_conflicting_attribute); - S.Diag(NewDeclAttr->getLocation(), diag::note_conflicting_attribute); - } -} - /// MergeFunctionDecl - We just parsed a function 'New' from /// declarator D which has the same name and scope as a previous /// declaration 'Old'. Figure out how to resolve this situation, @@ -3504,8 +3485,6 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, } } - checkDimensionsAndSetDiagnostics(*this, New, Old); - if (const auto *ILA = New->getAttr()) if (!Old->hasAttr()) { Diag(New->getLocation(), diag::err_attribute_missing_on_first_decl) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f75c9a963d186..93706678dc8ca 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3250,181 +3250,6 @@ static void handleWeakImportAttr(Sema &S, Decl *D, const ParsedAttr &AL) { D->addAttr(::new (S.Context) WeakImportAttr(S.Context, AL)); } -// Checks correctness of mutual usage of different work_group_size attributes: -// reqd_work_group_size, max_work_group_size and max_global_work_dim. -// Values of reqd_work_group_size arguments shall be equal or less than values -// coming from max_work_group_size. -// 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). -static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { - bool Result = true; - - // Returns the unsigned constant integer value represented by - // given expression. - auto getExprValue = [](const Expr *E, ASTContext &Ctx) { - return E->getIntegerConstantExpr(Ctx)->getZExtValue(); - }; - - ASTContext &Ctx = S.getASTContext(); - - // The arguments to reqd_work_group_size are ordered based on which index - // increments the fastest. In OpenCL, the first argument is the index that - // increments the fastest, and in SYCL, the last argument is the index that - // increments the fastest. - // - // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are - // available in SYCL modes and follow the SYCL rules. - // __attribute__((reqd_work_group_size)) is only available in OpenCL mode - // and follows the OpenCL rules. - if (const auto *A = D->getAttr()) { - bool CheckFirstArgument = - S.getLangOpts().OpenCL - ? getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getZDimVal() - : getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getXDimVal(); - bool CheckSecondArgument = - getExprValue(AL.getArgAsExpr(1), Ctx) > *A->getYDimVal(); - bool CheckThirdArgument = - S.getLangOpts().OpenCL - ? getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getXDimVal() - : getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getZDimVal(); - - if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) { - S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes) - << AL << A; - S.Diag(A->getLocation(), diag::note_conflicting_attribute); - Result &= false; - } - } - - if (const auto *A = D->getAttr()) { - if (!((getExprValue(AL.getArgAsExpr(0), Ctx) >= - getExprValue(A->getXDim(), Ctx)) && - (getExprValue(AL.getArgAsExpr(1), Ctx) >= - getExprValue(A->getYDim(), Ctx)) && - (getExprValue(AL.getArgAsExpr(2), Ctx) >= - getExprValue(A->getZDim(), Ctx)))) { - S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes) - << AL << A; - S.Diag(A->getLocation(), diag::note_conflicting_attribute); - Result &= false; - } - } - return Result; -} - -// Handles reqd_work_group_size. -template -static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { - if (D->isInvalidDecl()) - return; - - S.CheckDeprecatedSYCLAttributeSpelling(AL); - // __attribute__((reqd_work_group_size)) and [[cl::reqd_work_group_size]] - // all require exactly three arguments. - if ((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && - AL.getAttributeSpellingListIndex() == - ReqdWorkGroupSizeAttr::CXX11_cl_reqd_work_group_size) || - AL.getSyntax() == ParsedAttr::AS_GNU) { - if (!AL.checkExactlyNumArgs(S, 3)) - return; - } - - Expr *XDimExpr = AL.getArgAsExpr(0); - - // If no attribute argument is specified, set the second and third argument - // to the default value 1, but only if the sycl::reqd_work_group_size - // spelling was used. - auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) { - assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && - AL.getScopeName()->isStr("sycl")); - return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc()); - }; - - Expr *YDimExpr = - AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL); - - Expr *ZDimExpr = - AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL); - - ASTContext &Ctx = S.getASTContext(); - - if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && - !ZDimExpr->isValueDependent()) { - llvm::APSInt XDimVal, YDimVal, ZDimVal; - ExprResult XDim = S.VerifyIntegerConstantExpression(XDimExpr, &XDimVal); - ExprResult YDim = S.VerifyIntegerConstantExpression(YDimExpr, &YDimVal); - ExprResult ZDim = S.VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); - - if (XDim.isInvalid()) - return; - XDimExpr = XDim.get(); - - if (YDim.isInvalid()) - return; - YDimExpr = YDim.get(); - - if (ZDim.isInvalid()) - return; - ZDimExpr = ZDim.get(); - - // If the num_simd_work_items attribute is specified on a declaration it - // must evenly divide the index that increments fastest in the - // reqd_work_group_size attribute. In OpenCL, the first argument increments - // the fastest, and in SYCL, the last argument increments the fastest. - if (const auto *A = D->getAttr()) { - int64_t NumSimdWorkItems = - A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue(); - - unsigned WorkGroupSize = S.getLangOpts().OpenCL ? XDimVal.getZExtValue() - : ZDimVal.getZExtValue(); - - if (WorkGroupSize % NumSimdWorkItems != 0) { - S.Diag(A->getLocation(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) - << A << AL; - S.Diag(AL.getLoc(), diag::note_conflicting_attribute); - return; - } - } - - // If the declaration has a 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 || - ExistingAttr->getYDimVal(Ctx) != YDimVal || - ExistingAttr->getZDimVal(Ctx) != ZDimVal) { - S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; - S.Diag(ExistingAttr->getLocation(), diag::note_conflicting_attribute); - } - } - if (!checkWorkGroupSizeValues(S, D, AL)) - return; - } - - S.addIntelTripleArgAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr); -} - // Returns a DupArgResult value; Same means the args have the same value, // Different means the args do not have the same value, and Unknown means that // the args cannot (yet) be compared. @@ -3558,12 +3383,17 @@ static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) { S.AddWorkGroupSizeHintAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr); } -// Handles max_work_group_size attribute. -// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on a -// declaration along with [[intel::max_global_work_dim()]] attribute, -// check to see if all arguments of [[intel::max_work_group_size(X, Y, Z)]] -// attribute hold value 1 in case the argument of -// [[intel::max_global_work_dim()]] attribute equals to 0. +// Checks correctness of mutual usage of different work_group_size attributes: +// reqd_work_group_size, max_work_group_size, and max_global_work_dim. +// +// If [[intel::max_work_group_size(X, Y, Z)]] or +// [[sycl::reqd_work_group_size(X, Y, Z)]] or +// [[cl::reqd_work_group_size(X, Y, Z)]] +// or __attribute__((reqd_work_group_size)) attribute is specified on a +// declaration along with [[intel::max_global_work_dim()]] attribute, check to +// see if all arguments of 'max_work_group_size' or different spellings of +// 'reqd_work_group_size' attribute hold value 1 in case the argument of +// [[intel::max_global_work_dim()]] attribute value equals to 0. static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, const Expr *YDim, const Expr *ZDim) { // If any of the operand is still value dependent, we can't test anything. @@ -3582,21 +3412,57 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, ZDimExpr->getResultAsAPSInt() != 1)); } -// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on -// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] -// attribute, check to see if values of reqd_work_group_size arguments are -// equal or less than values of max_work_group_size attribute arguments. -static bool checkWorkGroupSizeAttrValues(const Expr *RWGS, const Expr *MWGS) { +// Checks correctness of mutual usage of different work_group_size attributes: +// reqd_work_group_size and max_work_group_size. +// +// If the 'reqd_work_group_size' attribute is specified on a declaration along +// with 'max_work_group_size' attribute, check to see if values of +// 'reqd_work_group_size' attribute arguments are equal to or less than values +// of 'max_work_group_size' attribute arguments. +// +// The arguments to reqd_work_group_size are ordered based on which index +// increments the fastest. In OpenCL, the first argument is the index that +// increments the fastest, and in SYCL, the last argument is the index that +// increments the fastest. +// +// __attribute__((reqd_work_group_size)) follows the OpenCL rules in OpenCL +// mode. All spellings of reqd_work_group_size attribute (regardless of +// syntax used) follow the SYCL rules when in SYCL mode. +static bool checkMaxAllowedWorkGroupSize( + Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, + const Expr *MWGSXDim, const Expr *MWGSYDim, const Expr *MWGSZDim) { // If any of the operand is still value dependent, we can't test anything. - const auto *RWGSCE = dyn_cast(RWGS); - const auto *MWGSCE = dyn_cast(MWGS); - - if (!RWGSCE || !MWGSCE) + const auto *RWGSXDimExpr = dyn_cast(RWGSXDim); + const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); + const auto *RWGSZDimExpr = dyn_cast(RWGSZDim); + const auto *MWGSXDimExpr = dyn_cast(MWGSXDim); + const auto *MWGSYDimExpr = dyn_cast(MWGSYDim); + const auto *MWGSZDimExpr = dyn_cast(MWGSZDim); + + if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !MWGSXDimExpr || + !MWGSYDimExpr || !MWGSZDimExpr) return false; - // Otherwise, check if value of reqd_work_group_size argument is - // greater than value of max_work_group_size attribute argument. - return RWGSCE->getResultAsAPSInt() > MWGSCE->getResultAsAPSInt(); + // Otherwise, check if values of 'reqd_work_group_size' attribute arguments + // are greater than values of 'max_work_group_size' attribute arguments. + bool CheckFirstArgument = + S.getLangOpts().OpenCL + ? RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSZDimExpr->getResultAsAPSInt().getZExtValue() + : RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSXDimExpr->getResultAsAPSInt().getZExtValue(); + + bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSYDimExpr->getResultAsAPSInt().getZExtValue(); + + bool CheckThirdArgument = + S.getLangOpts().OpenCL + ? RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSXDimExpr->getResultAsAPSInt().getZExtValue() + : RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSZDimExpr->getResultAsAPSInt().getZExtValue(); + + return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, @@ -3632,33 +3498,17 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, if (!XDim || !YDim || !ZDim) return; - // If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on - // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] - // attribute, check to see if values of reqd_work_group_size arguments are - // equal or less than values of max_work_group_size attribute arguments. - // - // The arguments to reqd_work_group_size are ordered based on which index - // increments the fastest. In OpenCL, the first argument is the index that - // increments the fastest, and in SYCL, the last argument is the index that - // increments the fastest. + // If the 'max_work_group_size' attribute is specified on a declaration along + // with 'reqd_work_group_size' attribute, check to see if values of + // 'reqd_work_group_size' attribute arguments are equal to or less than values + // of 'max_work_group_size' attribute arguments. // - // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are - // available in SYCL modes and follow the SYCL rules. - // __attribute__((reqd_work_group_size)) is only available in OpenCL mode - // and follows the OpenCL rules. + // We emit diagnostic if values of 'reqd_work_group_size' attribute arguments + // are greater than values of 'max_work_group_size' attribute arguments. if (const auto *DeclAttr = D->getAttr()) { - bool CheckFirstArgument = - getLangOpts().OpenCL - ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim) - : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim); - bool CheckSecondArgument = - checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim); - bool CheckThirdArgument = - getLangOpts().OpenCL - ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim) - : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim); - - if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) { + if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), + DeclAttr->getYDim(), DeclAttr->getZDim(), + XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) << CI << DeclAttr; Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); @@ -3667,7 +3517,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, } // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr, check to see if - // the attribute holds equal values to (1, 1, 1) in case the value of + // the attribute holds values equal to (1, 1, 1) in case the value of // SYCLIntelMaxGlobalWorkDimAttr equals to 0. if (const auto *DeclAttr = D->getAttr()) { if (InvalidWorkGroupSizeAttrs(DeclAttr->getValue(), XDim, YDim, ZDim)) { @@ -3726,33 +3576,17 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( return nullptr; } - // If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on - // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] - // attribute, check to see if values of reqd_work_group_size arguments are - // equal or less than values of max_work_group_size attribute arguments. + // If the 'max_work_group_size' attribute is specified on a declaration along + // with 'reqd_work_group_size' attribute, check to see if values of + // 'reqd_work_group_size' attribute arguments are equal to or less than values + // of 'max_work_group_size' attribute arguments. // - // The arguments to reqd_work_group_size are ordered based on which index - // increments the fastest. In OpenCL, the first argument is the index that - // increments the fastest, and in SYCL, the last argument is the index that - // increments the fastest. - // - // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are - // available in SYCL modes and follow the SYCL rules. - // __attribute__((reqd_work_group_size)) is only available in OpenCL mode - // and follows the OpenCL rules. + // We emit diagnostic if values of 'reqd_work_group_size' attribute arguments + // are greater than values of 'max_work_group_size' attribute arguments. if (const auto *DeclAttr = D->getAttr()) { - bool CheckFirstArgument = - getLangOpts().OpenCL - ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim()) - : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim()); - bool CheckSecondArgument = - checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim()); - bool CheckThirdArgument = - getLangOpts().OpenCL - ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim()) - : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim()); - - if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) { + if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), + DeclAttr->getYDim(), DeclAttr->getZDim(), + A.getXDim(), A.getYDim(), A.getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) << DeclAttr << &A; Diag(A.getLoc(), diag::note_conflicting_attribute); @@ -3760,10 +3594,9 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( } } - // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr, - // check to see if the attribute holds equal values to - // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr - // equals to 0. + // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr, check to see if + // the attribute holds values equal to (1, 1, 1) in case the value of + // SYCLIntelMaxGlobalWorkDimAttr equals to 0. if (const auto *DeclAttr = D->getAttr()) { if (InvalidWorkGroupSizeAttrs(DeclAttr->getValue(), A.getXDim(), A.getYDim(), A.getZDim())) { @@ -3777,12 +3610,255 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( Context, A, A.getXDim(), A.getYDim(), A.getZDim()); } +// Handles max_work_group_size attribute. static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { S.AddSYCLIntelMaxWorkGroupSizeAttr(D, AL, AL.getArgAsExpr(0), AL.getArgAsExpr(1), AL.getArgAsExpr(2)); } +// Handles reqd_work_group_size. +// If the 'reqd_work_group_size' attribute is specified on a declaration along +// with 'num_simd_work_items' attribute, the required work group size specified +// by 'num_simd_work_items' attribute must evenly divide the index that +// increments fastest in the 'reqd_work_group_size' attribute. +// +// The arguments to reqd_work_group_size are ordered based on which index +// increments the fastest. In OpenCL, the first argument is the index that +// increments the fastest, and in SYCL, the last argument is the index that +// increments the fastest. +// +// __attribute__((reqd_work_group_size)) follows the OpenCL rules in OpenCL +// mode. All spellings of reqd_work_group_size attribute (regardless of +// syntax used) follow the SYCL rules when in SYCL mode. +static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, + const Expr *RWGSXDim, const Expr *RWGSZDim) { + // If any of the operand is still value dependent, we can't test anything. + const auto *NSWIValueExpr = dyn_cast(NSWIValue); + const auto *RWGSXDimExpr = dyn_cast(RWGSXDim); + const auto *RWGSZDimExpr = dyn_cast(RWGSZDim); + + if (!NSWIValueExpr || !RWGSXDimExpr || !RWGSZDimExpr) + return false; + + // Otherwise, check which argument increments the fastest + // in OpenCL vs SYCL mode. + unsigned WorkGroupSize = + S.getLangOpts().OpenCL + ? (RWGSXDimExpr->getResultAsAPSInt()).getZExtValue() + : (RWGSZDimExpr->getResultAsAPSInt()).getZExtValue(); + + // Check if the required work group size specified by 'num_simd_work_items' + // attribute evenly divides the index that increments fastest in the + // 'reqd_work_group_size' attribute. + return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getZExtValue() != 0; +} + +void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDim, Expr *YDim, Expr *ZDim) { + // Returns nullptr if diagnosing, otherwise returns the original expression + // or the original expression converted to a constant expression. + auto CheckAndConvertArg = [&](Expr *E) -> Expr * { + // Check if the expression is not value dependent. + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a strictly positive value. + if (ArgVal <= 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*positive*/ 0; + return nullptr; + } + } + return E; + }; + + // Check all three argument values, and if any are bad, bail out. This will + // convert the given expressions into constant expressions when possible. + XDim = CheckAndConvertArg(XDim); + YDim = CheckAndConvertArg(YDim); + ZDim = CheckAndConvertArg(ZDim); + if (!XDim || !YDim || !ZDim) + return; + + // If the declaration has a ReqdWorkGroupSizeAttr, check to see if + // the attribute holds values equal to (1, 1, 1) in case the value of + // SYCLIntelMaxGlobalWorkDimAttr equals to 0. + if (const auto *DeclAttr = D->getAttr()) { + if (InvalidWorkGroupSizeAttrs(DeclAttr->getValue(), XDim, YDim, ZDim)) { + Diag(CI.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one) + << CI << DeclAttr; + } + } + + // If the 'max_work_group_size' attribute is specified on a declaration along + // with 'reqd_work_group_size' attribute, check to see if values of + // 'reqd_work_group_size' attribute arguments are equal to or less than values + // of 'max_work_group_size' attribute arguments. + // + // We emit diagnostic if values of 'reqd_work_group_size' attribute arguments + // are greater than values of 'max_work_group_size' attribute arguments. + if (const auto *DeclAttr = D->getAttr()) { + if (checkMaxAllowedWorkGroupSize(*this, XDim, YDim, ZDim, + DeclAttr->getXDim(), DeclAttr->getYDim(), + DeclAttr->getZDim())) { + Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) + << CI << DeclAttr; + Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); + return; + } + } + + // If the 'reqd_work_group_size' attribute is specified on a declaration + // along with 'num_simd_work_items' attribute, the required work group size + // specified by 'num_simd_work_items' attribute must evenly divide the index + // that increments fastest in the 'reqd_work_group_size' attribute. + if (const auto *DeclAttr = D->getAttr()) { + if (CheckWorkGroupSize(*this, DeclAttr->getValue(), XDim, ZDim)) { + Diag(DeclAttr->getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) + << DeclAttr << CI; + Diag(CI.getLoc(), diag::note_conflicting_attribute); + return; + } + } + + // If the attribute was already applied with different arguments, then + // diagnose the second attribute as a duplicate and don't add it. + if (const auto *Existing = D->getAttr()) { + DupArgResult Results[] = {AreArgValuesIdentical(XDim, Existing->getXDim()), + AreArgValuesIdentical(YDim, Existing->getYDim()), + AreArgValuesIdentical(ZDim, Existing->getZDim())}; + // If any of the results are known to be different, we can diagnose at this + // point and drop the attribute. + if (llvm::is_contained(Results, DupArgResult::Different)) { + Diag(CI.getLoc(), diag::err_duplicate_attribute) << CI; + Diag(Existing->getLoc(), diag::note_previous_attribute); + return; + } + + // If all of the results are known to be the same, we can silently drop the + // attribute. Otherwise, we have to add the attribute and resolve its + // differences later. + if (llvm::all_of(Results, + [](DupArgResult V) { return V == DupArgResult::Same; })) + return; + } + + D->addAttr(::new (Context) + ReqdWorkGroupSizeAttr(Context, CI, XDim, YDim, ZDim)); +} + +ReqdWorkGroupSizeAttr * +Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { + // If the declaration has a ReqdWorkGroupSizeAttr, check to see if the + // attribute holds values equal to (1, 1, 1) in case the value of + // SYCLIntelMaxGlobalWorkDimAttr equals to 0. + if (const auto *DeclAttr = D->getAttr()) { + if (InvalidWorkGroupSizeAttrs(DeclAttr->getValue(), A.getXDim(), + A.getYDim(), A.getZDim())) { + Diag(A.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one) + << &A << DeclAttr; + return nullptr; + } + } + + // If the 'max_work_group_size' attribute is specified on a declaration along + // with 'reqd_work_group_size' attribute, check to see if values of + // 'reqd_work_group_size' attribute arguments are equal or less than values + // of 'max_work_group_size' attribute arguments. + // + // We emit diagnostic if values of 'reqd_work_group_size' attribute arguments + // are greater than values of 'max_work_group_size' attribute arguments. + if (const auto *DeclAttr = D->getAttr()) { + if (checkMaxAllowedWorkGroupSize( + *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), + DeclAttr->getYDim(), DeclAttr->getZDim())) { + Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) + << DeclAttr << &A; + Diag(A.getLoc(), diag::note_conflicting_attribute); + return nullptr; + } + } + + // If the 'reqd_work_group_size' attribute is specified on a declaration + // along with 'num_simd_work_items' attribute, the required work group size + // specified by 'num_simd_work_items' attribute must evenly divide the index + // that increments fastest in the 'reqd_work_group_size' attribute. + if (const auto *DeclAttr = D->getAttr()) { + if (CheckWorkGroupSize(*this, DeclAttr->getValue(), A.getXDim(), + A.getZDim())) { + Diag(DeclAttr->getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) + << DeclAttr << &A; + Diag(A.getLoc(), diag::note_conflicting_attribute); + return nullptr; + } + } + + // Check to see if there's a duplicate attribute already applied. + if (const auto *DeclAttr = D->getAttr()) { + DupArgResult Results[] = { + AreArgValuesIdentical(DeclAttr->getXDim(), A.getXDim()), + AreArgValuesIdentical(DeclAttr->getYDim(), A.getYDim()), + AreArgValuesIdentical(DeclAttr->getZDim(), A.getZDim())}; + + // If any of the results are known to be different, we can diagnose at this + // point and drop the attribute. + if (llvm::is_contained(Results, DupArgResult::Different)) { + Diag(DeclAttr->getLoc(), diag::err_duplicate_attribute) << &A; + Diag(A.getLoc(), diag::note_previous_attribute); + return nullptr; + } + + // If all of the results are known to be the same, we can silently drop the + // attribute. Otherwise, we have to add the attribute and resolve its + // differences later. + if (llvm::all_of(Results, + [](DupArgResult V) { return V == DupArgResult::Same; })) + return nullptr; + } + + return ::new (Context) + ReqdWorkGroupSizeAttr(Context, A, A.getXDim(), A.getYDim(), A.getZDim()); +} + +static void handleReqdWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { + S.CheckDeprecatedSYCLAttributeSpelling(AL); + + // __attribute__((reqd_work_group_size)) and [[cl::reqd_work_group_size]] + // all require exactly three arguments. + if ((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && + AL.getAttributeSpellingListIndex() == + ReqdWorkGroupSizeAttr::CXX11_cl_reqd_work_group_size) || + AL.getSyntax() == ParsedAttr::AS_GNU) { + if (!AL.checkExactlyNumArgs(S, 3)) + return; + } + + Expr *XDimExpr = AL.getArgAsExpr(0); + + // If no attribute argument is specified, set the second and third argument + // to the default value 1, but only if the sycl::reqd_work_group_size + // spelling was used. + auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) { + assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && + AL.getScopeName()->isStr("sycl")); + return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc()); + }; + + Expr *YDimExpr = + AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL); + + Expr *ZDimExpr = + AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL); + + S.AddReqdWorkGroupSizeAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr); +} + void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, Expr *E) { if (!E->isValueDependent()) { @@ -3923,38 +3999,17 @@ void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, } } - // If the reqd_work_group_size attribute is specified on a declaration - // along with num_simd_work_items, the required work group size specified - // by num_simd_work_items attribute must evenly divide the index that - // increments fastest in the reqd_work_group_size attribute. - // - // The arguments to reqd_work_group_size are ordered based on which index - // increments the fastest. In OpenCL, the first argument is the index that - // increments the fastest, and in SYCL, the last argument is the index that - // increments the fastest. + // If the 'reqd_work_group_size' attribute is specified on a declaration + // along with 'num_simd_work_items' attribute, the required work group size + // specified by 'num_simd_work_items' attribute must evenly divide the index + // that increments fastest in the 'reqd_work_group_size' attribute. if (const auto *DeclAttr = D->getAttr()) { - Expr *XDimExpr = DeclAttr->getXDim(); - Expr *YDimExpr = DeclAttr->getYDim(); - Expr *ZDimExpr = DeclAttr->getZDim(); - - if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && - !ZDimExpr->isValueDependent()) { - llvm::APSInt XDimVal, ZDimVal; - ExprResult XDim = VerifyIntegerConstantExpression(XDimExpr, &XDimVal); - ExprResult ZDim = VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); - - if (XDim.isInvalid() || ZDim.isInvalid()) - return; - - unsigned WorkGroupSize = getLangOpts().OpenCL ? XDimVal.getZExtValue() - : ZDimVal.getZExtValue(); - - if (WorkGroupSize % ArgVal.getSExtValue() != 0) { - Diag(CI.getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) - << CI << DeclAttr; - Diag(DeclAttr->getLocation(), diag::note_conflicting_attribute); - return; - } + if (CheckWorkGroupSize(*this, E, DeclAttr->getXDim(), + DeclAttr->getZDim())) { + Diag(CI.getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) + << CI << DeclAttr; + Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); + return; } } } @@ -3978,6 +4033,21 @@ SYCLIntelNumSimdWorkItemsAttr *Sema::MergeSYCLIntelNumSimdWorkItemsAttr( } } } + + // If the 'reqd_work_group_size' attribute is specified on a declaration + // along with 'num_simd_work_items' attribute, the required work group size + // specified by 'num_simd_work_items' attribute must evenly divide the index + // that increments fastest in the 'reqd_work_group_size' attribute. + if (const auto *DeclAttr = D->getAttr()) { + if (CheckWorkGroupSize(*this, A.getValue(), DeclAttr->getXDim(), + DeclAttr->getZDim())) { + Diag(A.getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) + << &A << DeclAttr; + Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); + return nullptr; + } + } + return ::new (Context) SYCLIntelNumSimdWorkItemsAttr(Context, A, A.getValue()); } @@ -4211,9 +4281,9 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D, } // 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. + // ReqdWorkGroupSizeAttr, check to see if the attribute holds values equal + // to (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr equals + // to 0. if (ArgVal == 0) { if (checkWorkGroupSizeAttrExpr(*this, D, CI) || @@ -4243,9 +4313,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr( } // 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. + // ReqdWorkGroupSizeAttr, check to see if the attribute holds values equal to + // (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, @@ -10978,7 +11047,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, handleWorkGroupSizeHint(S, D, AL); break; case ParsedAttr::AT_ReqdWorkGroupSize: - handleWorkGroupSize(S, D, AL); + handleReqdWorkGroupSize(S, D, AL); break; case ParsedAttr::AT_SYCLIntelMaxWorkGroupSize: handleSYCLIntelMaxWorkGroupSize(S, D, AL); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8174a24bbc27c..fac53fcfbf94a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3857,10 +3857,9 @@ static void PropagateAndDiagnoseDeviceAttr( case attr::Kind::ReqdWorkGroupSize: { auto *RWGSA = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { - ASTContext &Ctx = S.getASTContext(); - if (Existing->getXDimVal(Ctx) != RWGSA->getXDimVal(Ctx) || - Existing->getYDimVal(Ctx) != RWGSA->getYDimVal(Ctx) || - Existing->getZDimVal(Ctx) != RWGSA->getZDimVal(Ctx)) { + if (*Existing->getXDimVal() != *RWGSA->getXDimVal() || + *Existing->getYDimVal() != *RWGSA->getYDimVal() || + *Existing->getZDimVal() != *RWGSA->getZDimVal()) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); @@ -3869,10 +3868,9 @@ static void PropagateAndDiagnoseDeviceAttr( } } else if (auto *Existing = SYCLKernel->getAttr()) { - ASTContext &Ctx = S.getASTContext(); - if (*Existing->getXDimVal() < RWGSA->getXDimVal(Ctx) || - *Existing->getYDimVal() < RWGSA->getYDimVal(Ctx) || - *Existing->getZDimVal() < RWGSA->getZDimVal(Ctx)) { + if (*Existing->getXDimVal() < *RWGSA->getXDimVal() || + *Existing->getYDimVal() < *RWGSA->getYDimVal() || + *Existing->getZDimVal() < *RWGSA->getZDimVal()) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); @@ -3905,10 +3903,9 @@ static void PropagateAndDiagnoseDeviceAttr( case attr::Kind::SYCLIntelMaxWorkGroupSize: { auto *SIMWGSA = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { - ASTContext &Ctx = S.getASTContext(); - if (Existing->getXDimVal(Ctx) > *SIMWGSA->getXDimVal() || - Existing->getYDimVal(Ctx) > *SIMWGSA->getYDimVal() || - Existing->getZDimVal(Ctx) > *SIMWGSA->getZDimVal()) { + if (*Existing->getXDimVal() > *SIMWGSA->getXDimVal() || + *Existing->getYDimVal() > *SIMWGSA->getYDimVal() || + *Existing->getZDimVal() > *SIMWGSA->getZDimVal()) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 1994219e08ecc..edd5da5030608 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -595,31 +595,6 @@ static void instantiateDependentAMDGPUWavesPerEUAttr( S.addAMDGPUWavesPerEUAttr(New, Attr, MinExpr, MaxExpr); } -template -static void instantiateIntelSYCTripleLFunctionAttr( - Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, - const AttrName *Attr, Decl *New) { - EnterExpressionEvaluationContext Unevaluated( - S, Sema::ExpressionEvaluationContext::ConstantEvaluated); - - ExprResult Result = S.SubstExpr(Attr->getXDim(), TemplateArgs); - if (Result.isInvalid()) - return; - Expr *XDimExpr = Result.getAs(); - - Result = S.SubstExpr(Attr->getYDim(), TemplateArgs); - if (Result.isInvalid()) - return; - Expr *YDimExpr = Result.getAs(); - - Result = S.SubstExpr(Attr->getZDim(), TemplateArgs); - if (Result.isInvalid()) - return; - Expr *ZDimExpr = Result.getAs(); - - S.addIntelTripleArgAttr(New, *Attr, XDimExpr, YDimExpr, ZDimExpr); -} - static void instantiateIntelFPGAForcePow2DepthAttr( Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, const IntelFPGAForcePow2DepthAttr *Attr, Decl *New) { @@ -891,6 +866,25 @@ static void instantiateSYCLIntelMaxWorkGroupSizeAttr( ZResult.get()); } +static void instantiateReqdWorkGroupSizeAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const ReqdWorkGroupSizeAttr *A, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + ExprResult XResult = S.SubstExpr(A->getXDim(), TemplateArgs); + if (XResult.isInvalid()) + return; + ExprResult YResult = S.SubstExpr(A->getYDim(), TemplateArgs); + if (YResult.isInvalid()) + return; + ExprResult ZResult = S.SubstExpr(A->getZDim(), TemplateArgs); + if (ZResult.isInvalid()) + return; + + S.AddReqdWorkGroupSizeAttr(New, *A, XResult.get(), YResult.get(), + ZResult.get()); +} + // This doesn't take any template parameters, but we have a custom action that // needs to happen when the kernel itself is instantiated. We need to run the // ItaniumMangler to mark the names required to name this kernel. @@ -1116,8 +1110,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *ReqdWorkGroupSize = dyn_cast(TmplAttr)) { - instantiateIntelSYCTripleLFunctionAttr( - *this, TemplateArgs, ReqdWorkGroupSize, New); + instantiateReqdWorkGroupSizeAttr(*this, TemplateArgs, ReqdWorkGroupSize, + New); continue; } if (const auto *SYCLIntelMaxWorkGroupSize = diff --git a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl index 4ac0a55541ab3..30142ed4ec2cd 100644 --- a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl +++ b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl @@ -32,9 +32,9 @@ void f_kernel_image2d_t( kernel image2d_t image ) { // expected-error {{'kernel' int __kernel x; // expected-error {{'__kernel' attribute only applies to functions}} } -kernel __attribute__((reqd_work_group_size(1,2,0))) void kernel11(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} -kernel __attribute__((reqd_work_group_size(1,0,2))) void kernel12(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} -kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} +kernel __attribute__((reqd_work_group_size(1,2,0))) void kernel11(){} // expected-error {{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} +kernel __attribute__((reqd_work_group_size(1,0,2))) void kernel12(){} // expected-error {{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} +kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expected-error {{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} __attribute__((intel_reqd_sub_group_size(8))) void kernel14(){} // expected-error {{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel}} kernel __attribute__((intel_reqd_sub_group_size(0))) void kernel15() {} // expected-error {{'intel_reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} @@ -45,7 +45,6 @@ kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_s // expected-note {{previous attribute is here}} __kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a positive integral compile time constant expression}} -__kernel __attribute__((reqd_work_group_size(8, 16, -32))) void neg2() {} //expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} // 4294967294 is a negative integer if treated as signed. // Should compile successfully, since we expect an unsigned. 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 2a9d375a010b0..1dc5ccca9a53c 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -110,38 +110,22 @@ struct TRIFuncObjGood9 { [[intel::max_global_work_dim(1)]] void TRIFuncObjGood9::operator()() const {} -// FIXME: We do not have support yet for checking -// reqd_work_group_size and max_global_work_dim -// attributes when merging, so the test compiles without -// any diagnostic when it shouldn't. +#ifdef TRIGGER_ERROR struct TRIFuncObjBad1 { - [[sycl::reqd_work_group_size(4, 4, 4)]] void + [[sycl::reqd_work_group_size(4, 4, 4)]] void // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} operator()() const; }; [[intel::max_global_work_dim(0)]] void TRIFuncObjBad1::operator()() const {} -// FIXME: We do not have support yet for checking -// reqd_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 { +struct TRIFuncObjBad2 { [[intel::max_global_work_dim(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'}} @@ -149,89 +133,91 @@ struct TRIFuncObjBad3 { operator()() const {} }; -struct TRIFuncObjBad4 { +struct TRIFuncObjBad3 { [[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 { +struct TRIFuncObjBad4 { [[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 { +struct TRIFuncObjBad5 { [[sycl::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 { +struct TRIFuncObjBad6 { [[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 {} +void +TRIFuncObjBad6::operator()() const {} -struct TRIFuncObjBad8 { +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 TRIFuncObjBad8::operator()() const {} +void +TRIFuncObjBad7::operator()() const {} -struct TRIFuncObjBad9 { +struct TRIFuncObjBad8 { [[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 {} +void +TRIFuncObjBad8::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'}} +struct TRIFuncObjBad9 { + // expected-error@+1{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} [[sycl::reqd_work_group_size(-4, 1)]] [[intel::max_global_work_dim(0)]] void operator()() const {} }; -struct TRIFuncObjBad11 { +struct TRIFuncObjBad10 { [[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}} +struct TRIFuncObjBad11 { + [[sycl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} [[intel::max_global_work_dim(0)]] void operator()() const {} }; -struct TRIFuncObjBad13 { +struct TRIFuncObjBad12 { [[sycl::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 { +struct TRIFuncObjBad13 { [[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 {} }; -struct TRIFuncObjBad15 { +struct TRIFuncObjBad14 { [[intel::max_work_group_size(4, 4, 4)]] void // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} operator()() const; }; -[[intel::max_global_work_dim(0)]] void TRIFuncObjBad15::operator()() const {} +[[intel::max_global_work_dim(0)]] void TRIFuncObjBad14::operator()() const {} #endif // TRIGGER_ERROR int main() { @@ -432,39 +418,10 @@ int main() { // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +#ifdef TRIGGER_ERROR 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{{$}} +#endif // TRIGGER_ERROR // Ignore duplicate attribute with same argument value. h.single_task( @@ -498,10 +455,9 @@ int main() { h.single_task(TRIFuncObjBad12()); h.single_task(TRIFuncObjBad13()); h.single_task(TRIFuncObjBad14()); - h.single_task(TRIFuncObjBad15()); - h.single_task( - []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} + 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 }); return 0; diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp index 524884b00675c..baf7133861c3f 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp @@ -79,12 +79,8 @@ f9() {} [[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK -// FIXME: We do not have support yet for checking -// reqd_work_group_size and max_work_group_size -// attributes when merging, so the test compiles without -// any diagnostic when it shouldn't. -[[sycl::reqd_work_group_size(64, 64, 64)]] void f12(); -[[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected error but now OK. +[[sycl::reqd_work_group_size(64, 64, 64)]] void f12(); // expected-note {{conflicting attribute is here}} +[[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} [[intel::max_work_group_size(16, 16, 16)]] // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(16, 64, 16)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp index 1b3903849ce87..87d10ac0c0867 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp @@ -32,18 +32,6 @@ class Functor32 { }; #endif // TRIGGER_ERROR -class Functor33 { -public: - // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} -}; - -class Functor30 { -public: - // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {} -}; - class Functor16 { public: [[sycl::reqd_work_group_size(16)]] void operator()() const {} @@ -112,32 +100,16 @@ int main() { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 32 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 32 + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -4 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - Functor33 f33; - h.single_task(f33); - - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 - // CHECK: ReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 30 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 30 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -30 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 30 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -30 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - Functor30 f30; - h.single_task(f30); + Functor64 f64; + h.single_task(f64); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp index a9ce1857d779b..1dea68a817583 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -32,30 +32,15 @@ void bar() { [[sycl::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} -class Functor33 { -public: - // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} -}; - +#ifdef TRIGGER_ERROR [[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} - -class Functor30 { -public: - // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {} -}; +#endif // TRIGGER_ERROR class Functor16 { public: [[sycl::reqd_work_group_size(16)]] void operator()() const {} }; -class Functor64 { -public: - [[sycl::reqd_work_group_size(64, 64)]] void operator()() const {} -}; - class Functor16x16x16 { public: [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} @@ -94,36 +79,30 @@ int main() { FunctorAttr fattr; h.single_task(fattr); - Functor33 f33; - h.single_task(f33); - - Functor30 f30; - h.single_task(f30); - - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { + h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { f32x32x32(); }); #ifdef TRIGGER_ERROR Functor8 f8; - h.single_task(f8); + h.single_task(f8); - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} f4x1x1(); f32x1x1(); }); - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} f16x1x1(); f16x16x1(); }); - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} f32x32x32(); f32x32x1(); }); // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { + h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); @@ -182,31 +161,6 @@ int main() { // CHECK-NEXT: value: Int 32 // CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} // CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -4 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name6 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 30 -// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -30 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' -// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -30 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' -// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name7 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 32 // CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} // CHECK-NEXT: ConstantExpr{{.*}}'int' diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 14bc2eeabd70e..19b7508bd297b 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -174,41 +174,54 @@ struct TRIFuncObjBad12 { }; struct TRIFuncObjBad13 { - [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} - [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - void operator()() const {} + [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} + void + operator()() const {} }; struct TRIFuncObjBad14 { - [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} - void operator()() const {} -}; - -struct TRIFuncObjBad15 { [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} [[sycl::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; -struct TRIFuncObjBad16 { +struct TRIFuncObjBad15 { [[sycl::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; -struct TRIFuncObjBad17 { +struct TRIFuncObjBad16 { [[intel::num_simd_work_items(3)]] [[sycl::reqd_work_group_size(3, 3, 3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; +struct TRIFuncObjBad17 { + [[intel::num_simd_work_items(-1)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} + [[sycl::reqd_work_group_size(-1)]] // expected-error{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + void + operator()() const {} +}; + struct TRIFuncObjBad18 { - [[intel::num_simd_work_items(-1)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - [[sycl::reqd_work_group_size(-1)]] // expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - void operator()() const {} + [[intel::num_simd_work_items(5)]] void // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} + operator()() const; +}; + +[[sycl::reqd_work_group_size(10, 5, 9)]] // expected-note{{conflicting attribute is here}} +void +TRIFuncObjBad18::operator()() const {} + +struct TRIFuncObjBad19 { + [[sycl::reqd_work_group_size(10, 5, 9)]] void // expected-note{{conflicting attribute is here}} + operator()() const; }; +[[intel::num_simd_work_items(5)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +void +TRIFuncObjBad19::operator()() const {} #endif // TRIGGER_ERROR // If the declaration has a [[sycl::reqd_work_group_size()]] // or [[cl::reqd_work_group_size()]] or @@ -239,6 +252,20 @@ struct TRIFuncObjGood4 { operator()() const {} }; +struct TRIFuncObjGood5 { + [[intel::num_simd_work_items(5)]] void + operator()() const; +}; + +[[sycl::reqd_work_group_size(3, 10, 5)]] void TRIFuncObjGood5::operator()() const {} + +struct TRIFuncObjGood6 { + [[sycl::reqd_work_group_size(3, 10, 5)]] void + operator()() const; +}; + +[[intel::num_simd_work_items(5)]] void TRIFuncObjGood6::operator()() const {} + [[intel::num_simd_work_items(2)]] __attribute__((reqd_work_group_size(3, 2, 6))) void func6(); // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \ expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} @@ -382,9 +409,44 @@ int main() { h.single_task(TRIFuncObjBad18()); h.single_task( - []() [[intel::num_simd_work_items(1), intel::num_simd_work_items(2)]]{}); // expected-warning{{attribute 'num_simd_work_items' is already applied with different arguments}} \ - // expected-note {{previous attribute is here}} + []() [[intel::num_simd_work_items(1), intel::num_simd_work_items(2)]] {}); // expected-warning{{attribute 'num_simd_work_items' is already applied with different arguments}} // expected-note {{previous attribute is here}} + + h.single_task(TRIFuncObjBad19()); + #endif // TRIGGER_ERROR + h.single_task(TRIFuncObjGood5()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel30 + // CHECK: SYCLIntelNumSimdWorkItemsAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: IntegerLiteral{{.*}}5{{$}} + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 10 + // CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: IntegerLiteral{{.*}}5{{$}} + + h.single_task(TRIFuncObjGood6()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel31 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 10 + // CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: IntegerLiteral{{.*}}5{{$}} + // CHECK: SYCLIntelNumSimdWorkItemsAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: IntegerLiteral{{.*}}5{{$}} }); return 0; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index af8730ea98700..1a5357a9f0ab2 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -24,14 +24,8 @@ func1(); [[sycl::reqd_work_group_size(8, 8, 8)]] void func2() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} //third case - expect error -[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}} -void -func3(); - -[[sycl::reqd_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}} -void -// expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} -func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(4, 4, 4)]] void func3(); // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} // fourth case - expect warning. [[intel::max_work_group_size(4, 4, 4)]] void func4(); // expected-note {{previous attribute is here}} diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp index 8bc681b5b690d..bbfa3136cfb2d 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -22,14 +22,11 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro [[cl::reqd_work_group_size(4)]] void four_with_more_feeling(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \ // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} class Functor32 { public: - // expected-note@+3{{conflicting attribute is here}} - // expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}} - // expected-error@+2{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}} - [[sycl::reqd_work_group_size(1, 1, 32)]] void + [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} + [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error{{attribute 'reqd_work_group_size' is already applied with different arguments}} operator()() const {} }; #endif // TRIGGER_ERROR @@ -45,6 +42,13 @@ class FunctorAttr { // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} }; +struct TRIFuncObjGood { + [[sycl::reqd_work_group_size(1, 2, 3)]] void + operator()() const; +}; + +[[sycl::reqd_work_group_size(1, 2, 3)]] void TRIFuncObjGood::operator()() const {} + int main() { q.submit([&](handler &h) { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 @@ -74,6 +78,19 @@ int main() { // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 FunctorAttr fattr; h.single_task(fattr); + + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} + h.single_task(TRIFuncObjGood()); }); return 0; } diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index 634257a4aebe1..8f0c02dcdd1f0 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -48,11 +48,8 @@ class Functor16 { #ifdef TRIGGER_ERROR class Functor32 { public: - // expected-note@+3{{conflicting attribute is here}} - // expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}} - // expected-error@+2 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}} - [[sycl::reqd_work_group_size(1, 1, 32)]] void + [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} + [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} operator()() const {} }; #endif @@ -118,6 +115,10 @@ int main() { }); #endif + // Ignore duplicate attribute. + h.single_task( + []() [[sycl::reqd_work_group_size(2, 2, 2), + sycl::reqd_work_group_size(2, 2, 2)]] {}); }); return 0; } @@ -166,3 +167,16 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 32 // CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} +// +// CHECK: FunctionDecl {{.*}}test_kernel11 +// CHECK: ReqdWorkGroupSizeAttr +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} +// CHECK-NOT: ReqdWorkGroupSizeAttr diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp new file mode 100644 index 0000000000000..318cbf4efeeb0 --- /dev/null +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -0,0 +1,130 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// Check the basics. +[[sycl::reqd_work_group_size]] void f(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} +[[sycl::reqd_work_group_size(12, 12, 12, 12)]] void f0(); // expected-error {{'reqd_work_group_size' attribute takes no more than 3 arguments}} +[[sycl::reqd_work_group_size("derp", 1, 2)]] void f1(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char[5]'}} +[[sycl::reqd_work_group_size(1, 1, 1)]] int i; // expected-error {{'reqd_work_group_size' attribute only applies to functions}} + +class Functor33 { +public: + // expected-error@+1{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} +}; + +class Functor30 { +public: + // expected-error@+1 2{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {} +}; + +// Tests for 'reqd_work_group_size' attribute duplication. +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +[[sycl::reqd_work_group_size(6, 6, 6)]] [[sycl::reqd_work_group_size(6, 6, 6)]] void f2() {} + +// No diagnostic is emitted because the arguments match. +[[sycl::reqd_work_group_size(32, 32, 32)]] void f3(); +[[sycl::reqd_work_group_size(32, 32, 32)]] void f3(); // OK + +// Produce a conflicting attribute warning when the args are different. +[[sycl::reqd_work_group_size(6, 6, 6)]] // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(16, 16, 16)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} +f4() {} + +// Catch the easy case where the attributes are all specified at once with +// different arguments. +struct TRIFuncObjGood1 { + // expected-note@+2 {{previous attribute is here}} + // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + [[sycl::reqd_work_group_size(64)]] [[sycl::reqd_work_group_size(128)]] void operator()() const {} +}; + +struct TRIFuncObjGood2 { + // expected-note@+2 {{previous attribute is here}} + // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + [[sycl::reqd_work_group_size(64, 64)]] [[sycl::reqd_work_group_size(128, 128)]] void operator()() const {} +}; + +struct TRIFuncObjGood3 { + [[sycl::reqd_work_group_size(8, 8)]] void // expected-note {{previous attribute is here}} + operator()() const; +}; + +[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} +void +TRIFuncObjGood3::operator()() const {} + +// Show that the attribute works on member functions. +class Functor { +public: + [[sycl::reqd_work_group_size(16, 16, 16)]] [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const; + [[sycl::reqd_work_group_size(16, 16, 16)]] [[sycl::reqd_work_group_size(32, 32, 32)]] void operator()(int) const; // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} +}; + +class FunctorC { +public: + [[intel::max_work_group_size(64, 64, 64)]] [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()() const; + [[intel::max_work_group_size(16, 16, 16)]] // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(64, 64, 64)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} + operator()(int) const; +}; + +// Ensure that template arguments behave appropriately based on instantiations. +template +[[sycl::reqd_work_group_size(N, 1, 1)]] void f6(); // #f6 + +// Test that template redeclarations also get diagnosed properly. +template +[[sycl::reqd_work_group_size(1, 1, 1)]] void f7(); // #f7prev + +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void f7() {} // #f7 + +// Test that a template redeclaration where the difference is known up front is +// diagnosed immediately, even without instantiation. +template +[[sycl::reqd_work_group_size(X, 1, Z)]] void f8(); // expected-note {{previous attribute is here}} +template +[[sycl::reqd_work_group_size(X, 2, Z)]] void f8(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} + +void instantiate() { + f6<1>(); // OK + // expected-error@#f6 {{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + f6<-1>(); // expected-note {{in instantiation}} + // expected-error@#f6 {{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} + f6<0>(); // expected-note {{in instantiation}} + f7<1, 1, 1>(); // OK, args are the same on the redecl. + // expected-error@#f7 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-note@#f7prev {{previous attribute is here}} + f7<2, 2, 2>(); // expected-note {{in instantiation}} +} + +// Tests for 'reqd_work_group_size' attribute duplication. + +[[sycl::reqd_work_group_size(8)]] // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} +f8(){}; + +[[sycl::reqd_work_group_size(32, 32, 1)]] [[sycl::reqd_work_group_size(32, 32)]] void f9() {} // OK + +// Test that template redeclarations also get diagnosed properly. +template +[[sycl::reqd_work_group_size(64, 1, 1)]] void f10(); // #f10prev +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void f10() {} // #f10err + +void test() { + f10<64, 1, 1>(); // OK, args are the same on the redecl. + // expected-error@#f10err {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-note@#f10prev {{previous attribute is here}} + f10<1, 1, 64>(); // expected-note {{in instantiation}} +} + +struct TRIFuncObjBad { + [[sycl::reqd_work_group_size(32, 1, 1)]] void // expected-note {{previous attribute is here}} + operator()() const; +}; + +[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} +void +TRIFuncObjBad::operator()() const {} 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 42f907356d3cf..5cee0bacfeb71 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 @@ -123,6 +123,16 @@ template template [[intel::max_work_group_size(N, N, N)]] void func14(); +template +[[intel::max_global_work_dim(0)]] void func15(); +template +[[sycl::reqd_work_group_size(N, N, N)]] void func15(); // 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 func16(); +template +[[sycl::reqd_work_group_size(N, N, N)]] void func16(); + int check() { func3<3>(); // OK func3<-1>(); // expected-note {{in instantiation of function template specialization 'func3<-1>' requested here}} @@ -137,11 +147,13 @@ int check() { func12<1>(); // OK func13<6>(); // expected-note {{in instantiation of function template specialization 'func13<6>' requested here}} func14<1>(); // OK + func15<6>(); // expected-note {{in instantiation of function template specialization 'func15<6>' requested here}} + func16<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 func15() {} +[[intel::max_global_work_dim(2)]] [[intel::max_global_work_dim(2)]] void func17() {} // CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' // CHECK: TemplateArgument integral 3 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} @@ -151,7 +163,7 @@ int check() { // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}} func15 'void ()' +// CHECK: FunctionDecl {{.*}} {{.*}} func17 'void ()' // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 2 diff --git a/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp index 12ffca8ab74f2..19ac44e439bcd 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp @@ -84,3 +84,18 @@ int check() { // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +[[sycl::reqd_work_group_size(4, 4, 4)]] [[sycl::reqd_work_group_size(4, 4, 4)]] void func4() {} +// CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()' +// 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-NOT: ReqdWorkGroupSizeAttr diff --git a/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp b/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp index 77698887ec8df..4704ac849537e 100644 --- a/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp @@ -130,6 +130,41 @@ template template [[intel::num_simd_work_items(2)]] void func12(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[intel::num_simd_work_items(N)]] void func13(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +__attribute__((reqd_work_group_size(8, 6, 3))) void func13(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'reqd_work_group_size' is deprecated}} expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} + +template +[[intel::num_simd_work_items(N)]] void func14(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[cl::reqd_work_group_size(8, 4, 5)]] void func14(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + +template +[[intel::num_simd_work_items(3)]] void func15(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[sycl::reqd_work_group_size(N, N, N)]] void func15(); // expected-note{{conflicting attribute is here}} + +template +[[intel::num_simd_work_items(N)]] void func16(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void func16(); // expected-note{{conflicting attribute is here}} + +template +[[intel::num_simd_work_items(3)]] void func17(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void func17(); // expected-note{{conflicting attribute is here}} + +template +[[intel::num_simd_work_items(2)]] void func18(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void func18(); // expected-note{{conflicting attribute is here}} + +template +[[intel::num_simd_work_items(2)]] void func19(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +template +[[sycl::reqd_work_group_size(N, N, N)]] void func19(); // expected-note{{conflicting attribute is here}} + int check1() { func6<3>(); // OK func6<2>(); // expected-note {{in instantiation of function template specialization 'func6<2>' requested here}} @@ -145,5 +180,19 @@ int check1() { func11<8, 6, 2>(); // OK func12<3>(); // expected-note {{in instantiation of function template specialization 'func12<3>' requested here}} func12<2>(); // OK + func13<3>(); // OK + func13<2>(); // expected-note {{in instantiation of function template specialization 'func13<2>' requested here}} + func14<4>(); // expected-note {{in instantiation of function template specialization 'func14<4>' requested here}} + func14<5>(); // OK + func15<5>(); // expected-note {{in instantiation of function template specialization 'func15<5>' requested here}} + func15<3>(); // OK + func16<6, 3, 5, 3>(); // expected-note {{in instantiation of function template specialization 'func16<6, 3, 5, 3>' requested here}} + func16<9, 6, 3, 3>(); // OK + func17<6, 3, 5>(); // expected-note {{in instantiation of function template specialization 'func17<6, 3, 5>' requested here}} + func17<9, 6, 3>(); // OK + func18<6, 4, 5>(); // expected-note {{in instantiation of function template specialization 'func18<6, 4, 5>' requested here}} + func18<8, 6, 2>(); // OK + func19<3>(); // expected-note {{in instantiation of function template specialization 'func19<3>' requested here}} + func19<2>(); // OK return 0; }