From 3985bb2c4fd466027049c2ecc578060a650f632c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 09:38:29 -0800 Subject: [PATCH 01/23] [SYCL] Refactor reqd_work_group_size attribute implementation Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 21 +- clang/include/clang/Sema/Sema.h | 64 +- clang/lib/CodeGen/CodeGenFunction.cpp | 33 +- clang/lib/CodeGen/TargetInfo.cpp | 25 +- clang/lib/Sema/SemaDecl.cpp | 25 +- clang/lib/Sema/SemaDeclAttr.cpp | 647 ++++++++++++------ clang/lib/Sema/SemaSYCL.cpp | 21 +- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 48 +- clang/test/SemaOpenCL/invalid-kernel-attrs.cl | 7 +- .../intel-max-global-work-dim-device.cpp | 52 +- .../SemaSYCL/intel-max-work-group-size.cpp | 8 +- ...eqd-work-group-size-device-direct-prop.cpp | 32 +- .../intel-reqd-work-group-size-device.cpp | 33 +- .../SemaSYCL/num_simd_work_items_device.cpp | 78 ++- .../redeclaration-attribute-propagation.cpp | 10 +- ...eqd-work-group-size-device-direct-prop.cpp | 26 +- .../SemaSYCL/reqd-work-group-size-device.cpp | 22 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 118 ++++ ...ice-intel-max-global-work-dim-template.cpp | 16 +- ...ce-intel-reqd-work-group-size-template.cpp | 15 + ...cl-device-num_simd_work_items-template.cpp | 49 ++ 21 files changed, 838 insertions(+), 512 deletions(-) create mode 100644 clang/test/SemaSYCL/reqd_work_group_size.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 05f33867af877..53dd866aa1586 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -3027,17 +3027,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 923d087dc01ac..24e6c0c54e257 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10519,9 +10519,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 * @@ -10601,6 +10598,11 @@ class Sema final { SYCLIntelMaxWorkGroupSizeAttr * MergeSYCLIntelMaxWorkGroupSizeAttr(Decl *D, const SYCLIntelMaxWorkGroupSizeAttr &A); + + 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); @@ -13679,62 +13681,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 e83ce1b74f548..8bf72c2e7b61a 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -636,24 +636,25 @@ 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. + /*if (getLangOpts().SYCLIsDevice) { + llvm::Metadata *AttrMDArgs[] = { + llvm::ConstantAsMetadata::get(Builder.getInt(*A->getZDimVal())), + llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())), + llvm::ConstantAsMetadata::get(Builder.getInt(*A->getXDimVal()))}; + Fn->setMetadata("reqd_work_group_size", + llvm::MDNode::get(Context, AttrMDArgs)); + }*/ 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)); + llvm::MDNode::get(Context, AttrMDArgs)); } bool IsKernelOrDevice = diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index a0d38dcb798b8..e60ae92ab1f40 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -8372,10 +8372,15 @@ 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(); + const auto *XDimExpr = cast(Attr->getXDim()); + const auto *YDimExpr = cast(Attr->getYDim()); + const auto *ZDimExpr = cast(Attr->getZDim()); + Optional XDimVal = XDimExpr->getResultAsAPSInt(); + Optional YDimVal = YDimExpr->getResultAsAPSInt(); + Optional ZDimVal = ZDimExpr->getResultAsAPSInt(); + unsigned XDim = XDimVal->getZExtValue(); + unsigned YDim = YDimVal->getZExtValue(); + unsigned ZDim = ZDimVal->getZExtValue(); Operands.push_back(llvm::ConstantAsMetadata::get( llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, XDim)))); @@ -9255,9 +9260,15 @@ 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(); + const auto *XDimExpr = cast(ReqdWGS->getXDim()); + const auto *YDimExpr = cast(ReqdWGS->getYDim()); + const auto *ZDimExpr = cast(ReqdWGS->getZDim()); + Optional XDimVal = XDimExpr->getResultAsAPSInt(); + Optional YDimVal = YDimExpr->getResultAsAPSInt(); + Optional ZDimVal = ZDimExpr->getResultAsAPSInt(); + XDim = XDimVal->getZExtValue(); + YDim = YDimVal->getZExtValue(); + ZDim = ZDimVal->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 27d7d0d5c5900..cdddcf5c00726 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2783,6 +2783,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, NewAttr = S.MergeSYCLIntelPipeIOAttr(D, *A); else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeSYCLIntelMaxWorkGroupSizeAttr(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)); @@ -3373,27 +3375,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, @@ -3482,8 +3463,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 6eb1f0f6bdc51..a44dbc0e5aa21 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. @@ -3586,17 +3411,36 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // 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) { +static bool checkWorkGroupSizeAttrValues(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(); + bool CheckFirstArgument = + S.getLangOpts().OpenCL + ? RWGSXDimExpr->getResultAsAPSInt() > MWGSZDimExpr->getResultAsAPSInt() + : RWGSXDimExpr->getResultAsAPSInt() > MWGSXDimExpr->getResultAsAPSInt(); + bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt() > MWGSYDimExpr->getResultAsAPSInt(); + bool CheckThirdArgument = + S.getLangOpts().OpenCL + ? RWGSZDimExpr->getResultAsAPSInt() > MWGSXDimExpr->getResultAsAPSInt() + : RWGSZDimExpr->getResultAsAPSInt() > MWGSZDimExpr->getResultAsAPSInt(); + + return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, @@ -3647,18 +3491,8 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. 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 (checkWorkGroupSizeAttrValues(*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); @@ -3741,18 +3575,9 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. 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 (checkWorkGroupSizeAttrValues(*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); @@ -3783,6 +3608,362 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, 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 reqd_work_group_size attribute, +// check to see if values of reqd_work_group_size arguments are +// equal or greater than values of reqd_work_group_size attribute arguments. +static bool checkReqdWorkGroupSizeAttrValues(Sema &S, const Expr *RWGSXDim, + const Expr *RWGSYDim, + const Expr *RWGSZDim, + const Expr *R1WGSXDim, + const Expr *R1WGSYDim, + const Expr *R1WGSZDim) { + // If any of the operand is still value dependent, we can't test anything. + const auto *RWGSXDimExpr = dyn_cast(RWGSXDim); + const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); + const auto *RWGSZDimExpr = dyn_cast(RWGSZDim); + const auto *R1WGSXDimExpr = dyn_cast(R1WGSXDim); + const auto *R1WGSYDimExpr = dyn_cast(R1WGSYDim); + const auto *R1WGSZDimExpr = dyn_cast(R1WGSZDim); + if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !R1WGSXDimExpr || !R1WGSYDimExpr || + !R1WGSZDimExpr) + return false; + + // Otherwise, check if value of reqd_work_group_size argument is + // less than value of reqd_work_group_size attribute argument. + bool CheckFirstArgument = + S.getLangOpts().OpenCL + ? RWGSXDimExpr->getResultAsAPSInt() < R1WGSZDimExpr->getResultAsAPSInt() + : RWGSXDimExpr->getResultAsAPSInt() < R1WGSXDimExpr->getResultAsAPSInt(); + bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt() < R1WGSYDimExpr->getResultAsAPSInt(); + bool CheckThirdArgument = + S.getLangOpts().OpenCL + ? RWGSZDimExpr->getResultAsAPSInt() < R1WGSXDimExpr->getResultAsAPSInt() + : RWGSZDimExpr->getResultAsAPSInt() < R1WGSZDimExpr->getResultAsAPSInt(); + + return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; +} + +// 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. +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 must evenly divide the index that increments fastest in the + // reqd_work_group_size attribute. + return WorkGroupSize % (NSWIValueExpr->getResultAsAPSInt()).getSExtValue() != 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 equal values 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 + // 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. + // + // __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. + if (const auto *DeclAttr = D->getAttr()) { + if (checkWorkGroupSizeAttrValues(*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 reqd_work_group_size attribute, + // check to see if values of reqd_work_group_size arguments are + // equal or greater than values of reqd_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. + if (const auto *DeclAttr = D->getAttr()) { + if (checkReqdWorkGroupSizeAttrValues(*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, 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. + 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::warn_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 equal values 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 [[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. + // + // __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. + if (const auto *DeclAttr = D->getAttr()) { + if (checkWorkGroupSizeAttrValues(*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 reqd_work_group_size attribute, + // check to see if values of reqd_work_group_size arguments are + // equal or greater than values of reqd_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. + if (const auto *DeclAttr = D->getAttr()) { + if (checkReqdWorkGroupSizeAttrValues(*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, 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. + 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::warn_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()) { @@ -3932,29 +4113,16 @@ void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, // 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. 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 +4146,29 @@ SYCLIntelNumSimdWorkItemsAttr *Sema::MergeSYCLIntelNumSimdWorkItemsAttr( } } } + + // 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. + // + // __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. + 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()); } @@ -10719,7 +10910,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 4a94c25bbfd17..7f113e8f1a258 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3902,10 +3902,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); @@ -3914,10 +3913,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); @@ -3950,10 +3948,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 590ab0c355e79..fa61e72634d7c 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) { @@ -855,6 +830,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. @@ -1080,8 +1074,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..21b868b84f742 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -110,31 +110,23 @@ 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 + [[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 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 @@ -195,8 +187,7 @@ void TRIFuncObjBad9::operator()() const {} // 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'}} + // 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 {} @@ -209,7 +200,7 @@ struct TRIFuncObjBad11 { }; struct TRIFuncObjBad12 { - [[sycl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} + [[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 {} }; @@ -432,39 +423,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( 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..d9ab0c294d812 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 @@ -30,19 +30,19 @@ class Functor32 { // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} }; -#endif // TRIGGER_ERROR class Functor33 { public: - // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + // 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-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + // 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 {} }; +#endif // TRIGGER_ERROR class Functor16 { public: @@ -109,35 +109,13 @@ int main() { FunctorAttr fattr; h.single_task(fattr); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 - // CHECK: ReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 32 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 32 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -4 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 +#ifdef TRIGGER_ERROR 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); +#endif // TRIGGER_ERROR }); 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..f3b0a434262de 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -32,9 +32,10 @@ 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}} +#ifdef TRIGGER_ERROR class Functor33 { public: - // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + // 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 {} }; @@ -42,9 +43,10 @@ class Functor33 { class Functor30 { public: - // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + // 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 {} }; +#endif // TRIGGER_ERROR class Functor16 { public: @@ -94,11 +96,13 @@ int main() { FunctorAttr fattr; h.single_task(fattr); +#ifdef TRIGGER_ERROR Functor33 f33; h.single_task(f33); Functor30 f30; h.single_task(f30); +#endif // TRIGGER_ERROR h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { f32x32x32(); @@ -176,31 +180,6 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 128 // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// 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' diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 14bc2eeabd70e..097acce521145 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -174,14 +174,14 @@ struct TRIFuncObjBad12 { }; struct TRIFuncObjBad13 { - [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} + [[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}} + [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} void operator()() const {} }; @@ -205,10 +205,25 @@ struct TRIFuncObjBad17 { 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'}} + [[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 TRIFuncObjBad19 { + [[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 TRIFuncObjBad19::operator()() const {} + +struct TRIFuncObjBad20 { + [[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 TRIFuncObjBad20::operator()() const {} #endif // TRIGGER_ERROR // If the declaration has a [[sycl::reqd_work_group_size()]] // or [[cl::reqd_work_group_size()]] or @@ -239,6 +254,22 @@ 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 +413,46 @@ 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()); + + h.single_task(TRIFuncObjBad20()); + #endif // TRIGGER_ERROR + h.single_task(TRIFuncObjGood5()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel31 + // 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_kernel32 + // 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..6fa4bedf404aa 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-warning {{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..d4967ae0aeaaa 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 @@ -25,10 +25,9 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro // 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-note@+2{{conflicting attribute is here}} // 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(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; @@ -45,6 +44,14 @@ 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 +81,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..a690c5a989ab9 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(1, 1, 32)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} 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..892c9c40a24aa --- /dev/null +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -0,0 +1,118 @@ +// 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}} + +// 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-warning {{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-warning@+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-warning@+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-warning {{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-warning {{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)]] [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()(int) const; // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} expected-note {{conflicting attribute is here}} +}; + +// 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-warning {{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-warning@#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}} +} + +// If the reqd_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 arguments are +// equal or greater than values coming from reqd_work_group_size attribute. +[[sycl::reqd_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f10() {} // OK + +[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 8)]] void f11() {}; // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + +[[sycl::reqd_work_group_size(32, 32, 1)]] [[sycl::reqd_work_group_size(32, 32)]] void f12() {} // OK + +// Test that template redeclarations also get diagnosed properly. +template +[[sycl::reqd_work_group_size(64, 1, 1)]] void f13(); // #f13conflict + +template +[[sycl::reqd_work_group_size(X, Y, Z)]] void f13() {} // #f13 + +void test() { + f13<64, 1, 1>(); // OK, args are the same on the redecl. + // expected-error@#f13 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // expected-note@#f13conflict {{conflicting attribute is here}} + f13<1, 1, 64>(); // expected-note {{in instantiation}} +} + +struct TRIFuncObjBad { + [[sycl::reqd_work_group_size(32, 1, 1)]] void // expected-note {{conflicting attribute is here}} + operator()() const; +}; + +[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +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..3123a67646f25 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; } From 6dd0cd5dbfbced1c58b3aa7c6ba8c39f6e35f0cb Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 10:31:52 -0800 Subject: [PATCH 02/23] Fix format issues Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 22 +-- clang/lib/Sema/SemaDeclAttr.cpp | 157 +++++++++--------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 6 +- .../SemaSYCL/num_simd_work_items_device.cpp | 31 ++-- .../redeclaration-attribute-propagation.cpp | 4 +- ...eqd-work-group-size-device-direct-prop.cpp | 6 +- .../SemaSYCL/reqd-work-group-size-device.cpp | 2 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 11 +- ...cl-device-num_simd_work_items-template.cpp | 12 +- 9 files changed, 127 insertions(+), 124 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 8bf72c2e7b61a..1dd7f3b628ee0 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -637,23 +637,13 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, if (const ReqdWorkGroupSizeAttr *A = FD->getAttr()) { // Attributes arguments (first and third) are reversed on SYCLDevice. - /*if (getLangOpts().SYCLIsDevice) { - llvm::Metadata *AttrMDArgs[] = { - llvm::ConstantAsMetadata::get(Builder.getInt(*A->getZDimVal())), - llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())), - llvm::ConstantAsMetadata::get(Builder.getInt(*A->getXDimVal()))}; - Fn->setMetadata("reqd_work_group_size", - llvm::MDNode::get(Context, AttrMDArgs)); - }*/ llvm::Metadata *AttrMDArgs[] = { - 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::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)); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index a44dbc0e5aa21..a9cff77c78e94 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3411,12 +3411,9 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // 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(Sema &S, const Expr *RWGSXDim, - const Expr *RWGSYDim, - const Expr *RWGSZDim, - const Expr *MWGSXDim, - const Expr *MWGSYDim, - const Expr *MWGSZDim) { +static bool checkWorkGroupSizeAttrValues( + 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 *RWGSXDimExpr = dyn_cast(RWGSXDim); const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); @@ -3424,21 +3421,27 @@ static bool checkWorkGroupSizeAttrValues(Sema &S, const Expr *RWGSXDim, 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) + + 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. - bool CheckFirstArgument = - S.getLangOpts().OpenCL - ? RWGSXDimExpr->getResultAsAPSInt() > MWGSZDimExpr->getResultAsAPSInt() - : RWGSXDimExpr->getResultAsAPSInt() > MWGSXDimExpr->getResultAsAPSInt(); - bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt() > MWGSYDimExpr->getResultAsAPSInt(); - bool CheckThirdArgument = - S.getLangOpts().OpenCL - ? RWGSZDimExpr->getResultAsAPSInt() > MWGSXDimExpr->getResultAsAPSInt() - : RWGSZDimExpr->getResultAsAPSInt() > MWGSZDimExpr->getResultAsAPSInt(); + bool CheckFirstArgument = S.getLangOpts().OpenCL + ? RWGSXDimExpr->getResultAsAPSInt() > + MWGSZDimExpr->getResultAsAPSInt() + : RWGSXDimExpr->getResultAsAPSInt() > + MWGSXDimExpr->getResultAsAPSInt(); + + bool CheckSecondArgument = + RWGSYDimExpr->getResultAsAPSInt() > MWGSYDimExpr->getResultAsAPSInt(); + + bool CheckThirdArgument = S.getLangOpts().OpenCL + ? RWGSZDimExpr->getResultAsAPSInt() > + MWGSXDimExpr->getResultAsAPSInt() + : RWGSZDimExpr->getResultAsAPSInt() > + MWGSZDimExpr->getResultAsAPSInt(); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3491,8 +3494,9 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim(), XDim, YDim, ZDim)) { + if (checkWorkGroupSizeAttrValues(*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); @@ -3575,9 +3579,9 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim(), A.getXDim(), - A.getYDim(), A.getZDim())) { + if (checkWorkGroupSizeAttrValues(*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); @@ -3613,12 +3617,9 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, // a declaration along with reqd_work_group_size attribute, // check to see if values of reqd_work_group_size arguments are // equal or greater than values of reqd_work_group_size attribute arguments. -static bool checkReqdWorkGroupSizeAttrValues(Sema &S, const Expr *RWGSXDim, - const Expr *RWGSYDim, - const Expr *RWGSZDim, - const Expr *R1WGSXDim, - const Expr *R1WGSYDim, - const Expr *R1WGSZDim) { +static bool checkReqdWorkGroupSizeAttrValues( + Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, + const Expr *R1WGSXDim, const Expr *R1WGSYDim, const Expr *R1WGSZDim) { // If any of the operand is still value dependent, we can't test anything. const auto *RWGSXDimExpr = dyn_cast(RWGSXDim); const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); @@ -3626,21 +3627,27 @@ static bool checkReqdWorkGroupSizeAttrValues(Sema &S, const Expr *RWGSXDim, const auto *R1WGSXDimExpr = dyn_cast(R1WGSXDim); const auto *R1WGSYDimExpr = dyn_cast(R1WGSYDim); const auto *R1WGSZDimExpr = dyn_cast(R1WGSZDim); - if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !R1WGSXDimExpr || !R1WGSYDimExpr || - !R1WGSZDimExpr) + + if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !R1WGSXDimExpr || + !R1WGSYDimExpr || !R1WGSZDimExpr) return false; // Otherwise, check if value of reqd_work_group_size argument is // less than value of reqd_work_group_size attribute argument. - bool CheckFirstArgument = - S.getLangOpts().OpenCL - ? RWGSXDimExpr->getResultAsAPSInt() < R1WGSZDimExpr->getResultAsAPSInt() - : RWGSXDimExpr->getResultAsAPSInt() < R1WGSXDimExpr->getResultAsAPSInt(); - bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt() < R1WGSYDimExpr->getResultAsAPSInt(); - bool CheckThirdArgument = - S.getLangOpts().OpenCL - ? RWGSZDimExpr->getResultAsAPSInt() < R1WGSXDimExpr->getResultAsAPSInt() - : RWGSZDimExpr->getResultAsAPSInt() < R1WGSZDimExpr->getResultAsAPSInt(); + bool CheckFirstArgument = S.getLangOpts().OpenCL + ? RWGSXDimExpr->getResultAsAPSInt() < + R1WGSZDimExpr->getResultAsAPSInt() + : RWGSXDimExpr->getResultAsAPSInt() < + R1WGSXDimExpr->getResultAsAPSInt(); + + bool CheckSecondArgument = + RWGSYDimExpr->getResultAsAPSInt() < R1WGSYDimExpr->getResultAsAPSInt(); + + bool CheckThirdArgument = S.getLangOpts().OpenCL + ? RWGSZDimExpr->getResultAsAPSInt() < + R1WGSXDimExpr->getResultAsAPSInt() + : RWGSZDimExpr->getResultAsAPSInt() < + R1WGSZDimExpr->getResultAsAPSInt(); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3655,8 +3662,7 @@ static bool checkReqdWorkGroupSizeAttrValues(Sema &S, const Expr *RWGSXDim, // increments the fastest, and in SYCL, the last argument is the index that // increments the fastest. static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, - const Expr *RWGSXDim, - const Expr *RWGSZDim) { + 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); @@ -3667,20 +3673,20 @@ static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, // Otherwise, check which argument increments the fastest // in OpenCL vs SYCL mode. - unsigned WorkGroupSize = S.getLangOpts().OpenCL - ? (RWGSXDimExpr->getResultAsAPSInt()).getZExtValue() - : (RWGSZDimExpr->getResultAsAPSInt()).getZExtValue(); + 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 must evenly divide the index that increments fastest in the // reqd_work_group_size attribute. - return WorkGroupSize % (NSWIValueExpr->getResultAsAPSInt()).getSExtValue() != 0; + return WorkGroupSize % (NSWIValueExpr->getResultAsAPSInt()).getSExtValue() != + 0; } -void Sema::AddReqdWorkGroupSizeAttr(Decl *D, - const AttributeCommonInfo &CI, - Expr *XDim, Expr *YDim, - Expr *ZDim) { +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 * { @@ -3734,8 +3740,8 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, XDim, YDim, ZDim, - DeclAttr->getXDim(), DeclAttr->getYDim(), + if (checkWorkGroupSizeAttrValues(*this, XDim, YDim, ZDim, + DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) << CI << DeclAttr; @@ -3759,9 +3765,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkReqdWorkGroupSizeAttrValues(*this, XDim, YDim, ZDim, - DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim())) { + if (checkReqdWorkGroupSizeAttrValues( + *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); @@ -3811,14 +3817,13 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, [](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) { +ReqdWorkGroupSizeAttr * +Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // If the declaration has a ReqdWorkGroupSizeAttr, // check to see if the attribute holds equal values to // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr @@ -3846,10 +3851,9 @@ ReqdWorkGroupSizeAttr *Sema::MergeReqdWorkGroupSizeAttr( // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, A.getXDim(), - A.getYDim(), A.getZDim(), - DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim())) { + if (checkWorkGroupSizeAttrValues( + *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); @@ -3871,9 +3875,9 @@ ReqdWorkGroupSizeAttr *Sema::MergeReqdWorkGroupSizeAttr( // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkReqdWorkGroupSizeAttrValues(*this, A.getXDim(), A.getYDim(), A.getZDim(), - DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim())) { + if (checkReqdWorkGroupSizeAttrValues( + *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); @@ -3895,7 +3899,8 @@ ReqdWorkGroupSizeAttr *Sema::MergeReqdWorkGroupSizeAttr( // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (CheckWorkGroupSize(*this, DeclAttr->getValue(), A.getXDim(), A.getZDim())) { + 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); @@ -3926,8 +3931,8 @@ ReqdWorkGroupSizeAttr *Sema::MergeReqdWorkGroupSizeAttr( return nullptr; } - return ::new (Context) ReqdWorkGroupSizeAttr( - Context, A, A.getXDim(), A.getYDim(), A.getZDim()); + return ::new (Context) + ReqdWorkGroupSizeAttr(Context, A, A.getXDim(), A.getYDim(), A.getZDim()); } static void handleReqdWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { @@ -4118,9 +4123,10 @@ void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (CheckWorkGroupSize(*this, E, DeclAttr->getXDim(), DeclAttr->getZDim())) { + if (CheckWorkGroupSize(*this, E, DeclAttr->getXDim(), + DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_sycl_num_kernel_wrong_reqd_wg_size) - << CI << DeclAttr; + << CI << DeclAttr; Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); return; } @@ -4161,13 +4167,14 @@ SYCLIntelNumSimdWorkItemsAttr *Sema::MergeSYCLIntelNumSimdWorkItemsAttr( // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. 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) + 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; - } - } + Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); + return nullptr; + } + } return ::new (Context) SYCLIntelNumSimdWorkItemsAttr(Context, A, A.getValue()); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index fa61e72634d7c..85f298541ff67 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -846,7 +846,7 @@ static void instantiateReqdWorkGroupSizeAttr( return; S.AddReqdWorkGroupSizeAttr(New, *A, XResult.get(), YResult.get(), - ZResult.get()); + ZResult.get()); } // This doesn't take any template parameters, but we have a custom action that @@ -1074,8 +1074,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *ReqdWorkGroupSize = dyn_cast(TmplAttr)) { - instantiateReqdWorkGroupSizeAttr( - *this, TemplateArgs, ReqdWorkGroupSize, New); + instantiateReqdWorkGroupSizeAttr(*this, TemplateArgs, ReqdWorkGroupSize, + New); continue; } if (const auto *SYCLIntelMaxWorkGroupSize = diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 097acce521145..5253de2411dc2 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -175,14 +175,16 @@ struct TRIFuncObjBad12 { struct TRIFuncObjBad13 { [[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 {} + [[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}} + [[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 requires a positive integral compile time constant expression}} - void operator()() const {} + void + operator()() const {} }; struct TRIFuncObjBad15 { @@ -204,9 +206,10 @@ struct TRIFuncObjBad17 { }; 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-error{{'reqd_work_group_size' attribute requires a positive integral compile time constant expression}} - void operator()() const {} + [[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 TRIFuncObjBad19 { @@ -215,7 +218,8 @@ struct TRIFuncObjBad19 { }; [[sycl::reqd_work_group_size(10, 5, 9)]] // expected-note{{conflicting attribute is here}} -void TRIFuncObjBad19::operator()() const {} +void +TRIFuncObjBad19::operator()() const {} struct TRIFuncObjBad20 { [[sycl::reqd_work_group_size(10, 5, 9)]] void // expected-note{{conflicting attribute is here}} @@ -223,7 +227,8 @@ struct TRIFuncObjBad20 { }; [[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 TRIFuncObjBad20::operator()() const {} +void +TRIFuncObjBad20::operator()() const {} #endif // TRIGGER_ERROR // If the declaration has a [[sycl::reqd_work_group_size()]] // or [[cl::reqd_work_group_size()]] or @@ -260,7 +265,7 @@ struct TRIFuncObjGood5 { }; [[sycl::reqd_work_group_size(3, 10, 5)]] -void TRIFuncObjGood5::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 @@ -268,7 +273,7 @@ struct TRIFuncObjGood6 { }; [[intel::num_simd_work_items(5)]] -void TRIFuncObjGood6::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}} \ @@ -413,8 +418,8 @@ 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()); h.single_task(TRIFuncObjBad20()); diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 6fa4bedf404aa..5baf19797bf71 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -24,8 +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)]] void func3(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[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-warning {{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 d4967ae0aeaaa..edce0eb43ab0a 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 @@ -27,8 +27,7 @@ class Functor32 { public: // expected-note@+2{{conflicting attribute is here}} // expected-error@+2{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - [[sycl::reqd_work_group_size(32, 1, 1)]] - [[sycl::reqd_work_group_size(1, 1, 32)]] void + [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; #endif // TRIGGER_ERROR @@ -49,8 +48,7 @@ struct TRIFuncObjGood { operator()() const; }; -[[sycl::reqd_work_group_size(1, 2, 3)]] -void TRIFuncObjGood::operator()() const {} +[[sycl::reqd_work_group_size(1, 2, 3)]] void TRIFuncObjGood::operator()() const {} int main() { q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index a690c5a989ab9..7402248de76f2 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -48,7 +48,7 @@ class Functor16 { #ifdef TRIGGER_ERROR class Functor32 { public: - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} operator()() const {} }; diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 892c9c40a24aa..85aa5cdb98641 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -39,7 +39,8 @@ struct TRIFuncObjGood3 { }; [[sycl::reqd_work_group_size(4, 4)]] // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} -void TRIFuncObjGood3::operator()() const {} +void +TRIFuncObjGood3::operator()() const {} // Show that the attribute works on member functions. class Functor { @@ -90,8 +91,9 @@ void instantiate() { // equal or greater than values coming from reqd_work_group_size attribute. [[sycl::reqd_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f10() {} // OK -[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 8)]] void f11() {}; // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 8)]] void +f11(){}; // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} [[sycl::reqd_work_group_size(32, 32, 1)]] [[sycl::reqd_work_group_size(32, 32)]] void f12() {} // OK @@ -115,4 +117,5 @@ struct TRIFuncObjBad { }; [[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} -void TRIFuncObjBad::operator()() const {} +void +TRIFuncObjBad::operator()() const {} 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 3123a67646f25..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 @@ -188,11 +188,11 @@ int check1() { 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 + 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; } From 4299ef6e3f6726052057f87ccdcd1281c8b8be52 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 11:00:39 -0800 Subject: [PATCH 03/23] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 8 ++++---- clang/lib/Sema/SemaDeclAttr.cpp | 4 ++-- clang/test/SemaSYCL/num_simd_work_items_device.cpp | 2 -- .../SemaSYCL/reqd-work-group-size-device-direct-prop.cpp | 2 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 6 +++--- 5 files changed, 10 insertions(+), 12 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 1dd7f3b628ee0..66b601fbb17d6 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -638,13 +638,13 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, if (const ReqdWorkGroupSizeAttr *A = FD->getAttr()) { // Attributes arguments (first and third) are reversed on SYCLDevice. llvm::Metadata *AttrMDArgs[] = { - llvm::ConstantAsMetadata::get(Builder.getInt( + llvm::ConstantAsMetadata::get(Builder.getInt( getLangOpts().SYCLIsDevice ? *A->getZDimVal() : *A->getXDimVal())), - llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())), + 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)); + Fn->setMetadata("reqd_work_group_size", + llvm::MDNode::get(Context, AttrMDArgs)); } bool IsKernelOrDevice = diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index a9cff77c78e94..1350467be7185 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3865,12 +3865,12 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // a declaration along with reqd_work_group_size attribute, // check to see if values of reqd_work_group_size arguments are // equal or greater than values of reqd_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. diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 5253de2411dc2..08b44c80ed9ac 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -264,7 +264,6 @@ struct TRIFuncObjGood5 { operator()() const; }; -[[sycl::reqd_work_group_size(3, 10, 5)]] [[sycl::reqd_work_group_size(3, 10, 5)]] void TRIFuncObjGood5::operator()() const {} struct TRIFuncObjGood6 { @@ -272,7 +271,6 @@ struct TRIFuncObjGood6 { operator()() const; }; -[[intel::num_simd_work_items(5)]] [[intel::num_simd_work_items(5)]] void TRIFuncObjGood6::operator()() const {} [[intel::num_simd_work_items(2)]] 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 edce0eb43ab0a..d464c1193ea73 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 @@ -26,7 +26,7 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro class Functor32 { public: // expected-note@+2{{conflicting attribute is here}} - // expected-error@+2{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 85aa5cdb98641..c7291e3a579bf 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -91,9 +91,9 @@ void instantiate() { // equal or greater than values coming from reqd_work_group_size attribute. [[sycl::reqd_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f10() {} // OK -[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 8)]] void -f11(){}; // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +f11(){}; [[sycl::reqd_work_group_size(32, 32, 1)]] [[sycl::reqd_work_group_size(32, 32)]] void f12() {} // OK From a7833f1631e2c555260358042c598e61e22e71d7 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 16:03:20 -0800 Subject: [PATCH 04/23] Fix the clang-tidy build after work_group_size attr changes Signed-off-by: Soumi Manna --- .../clang-tidy/altera/SingleWorkItemBarrierCheck.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp index 521126f990e97..12bc436f8999e 100644 --- a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp +++ b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp @@ -57,9 +57,9 @@ 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. From 4588ed76c7a823b068d818e42973398b37a725c2 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 16:10:48 -0800 Subject: [PATCH 05/23] Fix Format errors with clang-tidy build fix Signed-off-by: Soumi Manna --- .../clang-tidy/altera/SingleWorkItemBarrierCheck.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp index 12bc436f8999e..4b26e461b9d6d 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() > 1 || - *Attribute->getYDimVal() > 1 || - *Attribute->getZDimVal() > 1) + if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 || + *Attribute->getZDimVal() > 1) IsNDRange = true; } if (IsNDRange) // No warning if kernel is treated as an NDRange. From 9d18d2d692b6872c6cccf331baeb851d3e5f036c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Mar 2022 16:16:38 -0800 Subject: [PATCH 06/23] Fix Format errors with clang-tidy build fix Signed-off-by: Soumi Manna --- .../clang-tidy/altera/SingleWorkItemBarrierCheck.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp index 4b26e461b9d6d..87bf2dd622880 100644 --- a/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp +++ b/clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp @@ -57,8 +57,8 @@ void SingleWorkItemBarrierCheck::check(const MatchFinder::MatchResult &Result) { bool IsNDRange = false; if (MatchedDecl->hasAttr()) { const auto *Attribute = MatchedDecl->getAttr(); - if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 || - *Attribute->getZDimVal() > 1) + if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 || + *Attribute->getZDimVal() > 1) IsNDRange = true; } if (IsNDRange) // No warning if kernel is treated as an NDRange. From c89248154295a8c03f8c7234abdf4bbb635cc14f Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 13 Mar 2022 08:08:49 -0700 Subject: [PATCH 07/23] address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/TargetInfo.cpp | 24 ++++++------------------ clang/lib/Sema/SemaDeclAttr.cpp | 2 +- 2 files changed, 7 insertions(+), 19 deletions(-) diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index e60ae92ab1f40..5e722a2e69f60 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -8372,15 +8372,9 @@ void TCETargetCodeGenInfo::setTargetAttributes( SmallVector Operands; Operands.push_back(llvm::ConstantAsMetadata::get(F)); - const auto *XDimExpr = cast(Attr->getXDim()); - const auto *YDimExpr = cast(Attr->getYDim()); - const auto *ZDimExpr = cast(Attr->getZDim()); - Optional XDimVal = XDimExpr->getResultAsAPSInt(); - Optional YDimVal = YDimExpr->getResultAsAPSInt(); - Optional ZDimVal = ZDimExpr->getResultAsAPSInt(); - unsigned XDim = XDimVal->getZExtValue(); - unsigned YDim = YDimVal->getZExtValue(); - unsigned ZDim = ZDimVal->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)))); @@ -9260,15 +9254,9 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( Max = FlatWGS->getMax()->EvaluateKnownConstInt(Ctx).getExtValue(); } if (ReqdWGS) { - const auto *XDimExpr = cast(ReqdWGS->getXDim()); - const auto *YDimExpr = cast(ReqdWGS->getYDim()); - const auto *ZDimExpr = cast(ReqdWGS->getZDim()); - Optional XDimVal = XDimExpr->getResultAsAPSInt(); - Optional YDimVal = YDimExpr->getResultAsAPSInt(); - Optional ZDimVal = ZDimExpr->getResultAsAPSInt(); - XDim = XDimVal->getZExtValue(); - YDim = YDimVal->getZExtValue(); - ZDim = ZDimVal->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/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 1350467be7185..ec9e8fd377f1b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3681,7 +3681,7 @@ static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, // Check if 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. - return WorkGroupSize % (NSWIValueExpr->getResultAsAPSInt()).getSExtValue() != + return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getSExtValue() != 0; } From 7f2135cdd54da532dc3ac5d3933f771186b9e7b9 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 13 Mar 2022 08:22:32 -0700 Subject: [PATCH 08/23] Update num_simd_work_items attribute so that it takes ZExtValue instead of the SExtValue Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 2 +- clang/lib/Sema/SemaDeclAttr.cpp | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 66b601fbb17d6..bfd242f045b43 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -706,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/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index ec9e8fd377f1b..8321e9abc116f 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3681,8 +3681,7 @@ static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, // Check if 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. - return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getSExtValue() != - 0; + return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getZExtValue() != 0; } void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, From 88f0d6183327114b4d03d02c378679b31dfc05e4 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 13 Mar 2022 12:45:15 -0700 Subject: [PATCH 09/23] Add comparator as a template argument and get rid of one function Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 96 ++++++++++++--------------------- 1 file changed, 34 insertions(+), 62 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 8321e9abc116f..2d7e4d4ffca36 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3407,10 +3407,16 @@ 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 +// 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 arguments are // equal or less than values of max_work_group_size attribute arguments. +// If the 'reqd_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 arguments are +// equal or greater than values of 'reqd_work_group_size' attribute arguments. + +template static bool checkWorkGroupSizeAttrValues( Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, const Expr *MWGSXDim, const Expr *MWGSYDim, const Expr *MWGSZDim) { @@ -3426,22 +3432,27 @@ static bool checkWorkGroupSizeAttrValues( !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. - bool CheckFirstArgument = S.getLangOpts().OpenCL - ? RWGSXDimExpr->getResultAsAPSInt() > - MWGSZDimExpr->getResultAsAPSInt() - : RWGSXDimExpr->getResultAsAPSInt() > - MWGSXDimExpr->getResultAsAPSInt(); + // Otherwise, check if value of reqd_work_group_size argument is greater + // than value of max_work_group_size attribute argument. + // or check if value of reqd_work_group_size argument is less than value + // of reqd_work_group_size attribute argument. - bool CheckSecondArgument = - RWGSYDimExpr->getResultAsAPSInt() > MWGSYDimExpr->getResultAsAPSInt(); + bool CheckFirstArgument = + S.getLangOpts().OpenCL + ? COMPARATOR() (RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), + MWGSZDimExpr->getResultAsAPSInt().getZExtValue()) + : COMPARATOR() (RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), + MWGSXDimExpr->getResultAsAPSInt().getZExtValue()); - bool CheckThirdArgument = S.getLangOpts().OpenCL - ? RWGSZDimExpr->getResultAsAPSInt() > - MWGSXDimExpr->getResultAsAPSInt() - : RWGSZDimExpr->getResultAsAPSInt() > - MWGSZDimExpr->getResultAsAPSInt(); + bool CheckSecondArgument = COMPARATOR() + (RWGSYDimExpr->getResultAsAPSInt().getZExtValue(), MWGSYDimExpr->getResultAsAPSInt().getZExtValue()); + + bool CheckThirdArgument = COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), MWGSZDimExpr->getResultAsAPSInt().getZExtValue()); + S.getLangOpts().OpenCL + ? COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), + MWGSXDimExpr->getResultAsAPSInt().getZExtValue()) + : COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), + MWGSZDimExpr->getResultAsAPSInt().getZExtValue()); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3494,7 +3505,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), + if (checkWorkGroupSizeAttrValues>(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3579,7 +3590,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), + if (checkWorkGroupSizeAttrValues>(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3613,45 +3624,6 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, } // Handles reqd_work_group_size. -// If the reqd_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 arguments are -// equal or greater than values of reqd_work_group_size attribute arguments. -static bool checkReqdWorkGroupSizeAttrValues( - Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, - const Expr *R1WGSXDim, const Expr *R1WGSYDim, const Expr *R1WGSZDim) { - // If any of the operand is still value dependent, we can't test anything. - const auto *RWGSXDimExpr = dyn_cast(RWGSXDim); - const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); - const auto *RWGSZDimExpr = dyn_cast(RWGSZDim); - const auto *R1WGSXDimExpr = dyn_cast(R1WGSXDim); - const auto *R1WGSYDimExpr = dyn_cast(R1WGSYDim); - const auto *R1WGSZDimExpr = dyn_cast(R1WGSZDim); - - if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !R1WGSXDimExpr || - !R1WGSYDimExpr || !R1WGSZDimExpr) - return false; - - // Otherwise, check if value of reqd_work_group_size argument is - // less than value of reqd_work_group_size attribute argument. - bool CheckFirstArgument = S.getLangOpts().OpenCL - ? RWGSXDimExpr->getResultAsAPSInt() < - R1WGSZDimExpr->getResultAsAPSInt() - : RWGSXDimExpr->getResultAsAPSInt() < - R1WGSXDimExpr->getResultAsAPSInt(); - - bool CheckSecondArgument = - RWGSYDimExpr->getResultAsAPSInt() < R1WGSYDimExpr->getResultAsAPSInt(); - - bool CheckThirdArgument = S.getLangOpts().OpenCL - ? RWGSZDimExpr->getResultAsAPSInt() < - R1WGSXDimExpr->getResultAsAPSInt() - : RWGSZDimExpr->getResultAsAPSInt() < - R1WGSZDimExpr->getResultAsAPSInt(); - - return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; -} - // 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 @@ -3739,7 +3711,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, XDim, YDim, ZDim, + if (checkWorkGroupSizeAttrValues>(*this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3764,7 +3736,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkReqdWorkGroupSizeAttrValues( + if (checkWorkGroupSizeAttrValues>( *this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3850,7 +3822,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues( + if (checkWorkGroupSizeAttrValues>( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3874,7 +3846,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkReqdWorkGroupSizeAttrValues( + if (checkWorkGroupSizeAttrValues>( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) From dfcbd02a23525d757f4a2575c7dfafcfd6bc45d8 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 13 Mar 2022 13:23:06 -0700 Subject: [PATCH 10/23] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 161 ++++++++++++++++---------------- 1 file changed, 81 insertions(+), 80 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 2d7e4d4ffca36..c89d3c375dd37 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3407,8 +3407,8 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, ZDimExpr->getResultAsAPSInt() != 1)); } -// If the 'max_work_group_size' attribute is specified on -// a declaration along with 'reqd_work_group_size' attribute, +// 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 arguments are // equal or less than values of max_work_group_size attribute arguments. // If the 'reqd_work_group_size' attribute is specified on @@ -3416,43 +3416,44 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // check to see if values of reqd_work_group_size arguments are // equal or greater than values of 'reqd_work_group_size' attribute arguments. -template +template static bool checkWorkGroupSizeAttrValues( Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, - const Expr *MWGSXDim, const Expr *MWGSYDim, const Expr *MWGSZDim) { + const Expr *WGSXDim, const Expr *WGSYDim, const Expr *WGSZDim) { // If any of the operand is still value dependent, we can't test anything. 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); + const auto *WGSXDimExpr = dyn_cast(WGSXDim); + const auto *WGSYDimExpr = dyn_cast(WGSYDim); + const auto *WGSZDimExpr = dyn_cast(WGSZDim); 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. - // or check if value of reqd_work_group_size argument is less than value - // of reqd_work_group_size attribute argument. + // Otherwise, check if value of 'reqd_work_group_size' attribute argument + // is greater than value of 'max_work_group_size' attribute argument. + // or check if value of 'reqd_work_group_size' attribute argument is less + // than value of 'reqd_work_group_size' attribute argument. - bool CheckFirstArgument = - S.getLangOpts().OpenCL - ? COMPARATOR() (RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), - MWGSZDimExpr->getResultAsAPSInt().getZExtValue()) - : COMPARATOR() (RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), - MWGSXDimExpr->getResultAsAPSInt().getZExtValue()); + bool CheckFirstArgument = + S.getLangOpts().OpenCL + ? Comparator()(RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), + WGSZDimExpr->getResultAsAPSInt().getZExtValue()) + : Comparator()(RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), + WGSXDimExpr->getResultAsAPSInt().getZExtValue()); - bool CheckSecondArgument = COMPARATOR() - (RWGSYDimExpr->getResultAsAPSInt().getZExtValue(), MWGSYDimExpr->getResultAsAPSInt().getZExtValue()); + bool CheckSecondArgument = + Comparator()(RWGSYDimExpr->getResultAsAPSInt().getZExtValue(), + WGSYDimExpr->getResultAsAPSInt().getZExtValue()); - bool CheckThirdArgument = COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), MWGSZDimExpr->getResultAsAPSInt().getZExtValue()); - S.getLangOpts().OpenCL - ? COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), - MWGSXDimExpr->getResultAsAPSInt().getZExtValue()) - : COMPARATOR() (RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), - MWGSZDimExpr->getResultAsAPSInt().getZExtValue()); + bool CheckThirdArgument = + S.getLangOpts().OpenCL + ? Comparator()(RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), + WGSXDimExpr->getResultAsAPSInt().getZExtValue()) + : Comparator()(RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), + WGSZDimExpr->getResultAsAPSInt().getZExtValue()); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3490,10 +3491,10 @@ 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. + // 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. // // 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 @@ -3505,9 +3506,9 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>(*this, DeclAttr->getXDim(), - DeclAttr->getYDim(), DeclAttr->getZDim(), - XDim, YDim, ZDim)) { + if (checkWorkGroupSizeAttrValues>( + *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); @@ -3575,10 +3576,10 @@ 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 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 @@ -3590,9 +3591,9 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>(*this, DeclAttr->getXDim(), - DeclAttr->getYDim(), DeclAttr->getZDim(), - A.getXDim(), A.getYDim(), A.getZDim())) { + if (checkWorkGroupSizeAttrValues>( + *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); @@ -3624,10 +3625,10 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, } // Handles reqd_work_group_size. -// 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. +// 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 @@ -3650,9 +3651,9 @@ static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, ? (RWGSXDimExpr->getResultAsAPSInt()).getZExtValue() : (RWGSZDimExpr->getResultAsAPSInt()).getZExtValue(); - // Check if the required work group size specified by num_simd_work_items + // Check if 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. + // 'reqd_work_group_size' attribute. return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getZExtValue() != 0; } @@ -3697,10 +3698,10 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } } - // If the max_work_group_size attribute is specified on - // a declaration along with reqd_work_group_size attribute - // 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 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 @@ -3711,9 +3712,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>(*this, XDim, YDim, ZDim, - DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim())) { + if (checkWorkGroupSizeAttrValues>( + *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); @@ -3721,10 +3722,10 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } } - // If the reqd_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 arguments are - // equal or greater than values of reqd_work_group_size attribute + // If the 'reqd_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 greater than values of 'reqd_work_group_size' attribute // arguments. // // The arguments to reqd_work_group_size are ordered based on which index @@ -3746,10 +3747,10 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } } - // 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. + // 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 @@ -3808,10 +3809,10 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { } } - // 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 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 @@ -3832,10 +3833,10 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { } } - // If the reqd_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 arguments are - // equal or greater than values of reqd_work_group_size attribute arguments. + // If the 'reqd_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 greater than values of 'reqd_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 @@ -3856,10 +3857,10 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { } } - // 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. + // 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 @@ -4080,10 +4081,10 @@ 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. + // 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 @@ -4124,10 +4125,10 @@ SYCLIntelNumSimdWorkItemsAttr *Sema::MergeSYCLIntelNumSimdWorkItemsAttr( } } - // 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. + // 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 From 4032496fc9736ed8204565b5e4e8bf598b301baf Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 13 Mar 2022 13:33:02 -0700 Subject: [PATCH 11/23] Fix build errors Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index c89d3c375dd37..54c79bca9f1c1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3428,8 +3428,8 @@ static bool checkWorkGroupSizeAttrValues( const auto *WGSYDimExpr = dyn_cast(WGSYDim); const auto *WGSZDimExpr = dyn_cast(WGSZDim); - if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !MWGSXDimExpr || - !MWGSYDimExpr || !MWGSZDimExpr) + if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !WGSXDimExpr || + !WGSYDimExpr || !WGSZDimExpr) return false; // Otherwise, check if value of 'reqd_work_group_size' attribute argument From 87314f906c550a0e3fe3a1d3cc2676f8692e3c5d Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 14 Mar 2022 17:58:11 -0700 Subject: [PATCH 12/23] Fix regressions Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 93 +++++++++---------- .../redeclaration-attribute-propagation.cpp | 7 +- ...eqd-work-group-size-device-direct-prop.cpp | 10 +- .../SemaSYCL/reqd-work-group-size-device.cpp | 7 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 29 ++++-- 5 files changed, 78 insertions(+), 68 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 54c79bca9f1c1..a083d7340342c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3722,31 +3722,6 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } } - // If the 'reqd_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 greater than values of 'reqd_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. - if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( - *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 @@ -3780,8 +3755,31 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, if (llvm::is_contained(Results, DupArgResult::Different)) { Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI; Diag(Existing->getLoc(), diag::note_previous_attribute); + } + + // If the 'reqd_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 greater than values of 'reqd_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. + if (checkWorkGroupSizeAttrValues>( + *this, XDim, YDim, ZDim, Existing->getXDim(), Existing->getYDim(), + Existing->getZDim())) { + Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) + << CI << Existing; + Diag(Existing->getLoc(), diag::note_conflicting_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. @@ -3833,30 +3831,6 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { } } - // If the 'reqd_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 greater than values of 'reqd_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. - if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( - *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 @@ -3892,6 +3866,27 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { if (llvm::is_contained(Results, DupArgResult::Different)) { Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A; Diag(A.getLoc(), diag::note_previous_attribute); + } + + // If the 'reqd_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 greater than values of 'reqd_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. + if (checkWorkGroupSizeAttrValues>( + *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); return nullptr; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 5baf19797bf71..6165134edeb8a 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -24,9 +24,10 @@ 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)]] void func3(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} - +[[sycl::reqd_work_group_size(4, 4, 4)]] void func3(); // expected-note {{previous attribute is here}} \ + // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ + // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // fourth case - expect warning. [[intel::max_work_group_size(4, 4, 4)]] void func4(); // expected-note {{previous attribute is here}} [[intel::max_work_group_size(8, 8, 8)]] void func4() {} // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} 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 d464c1193ea73..6653b8fefe781 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,12 +22,14 @@ __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@+2{{conflicting attribute is here}} - // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void + // 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 {{previous attribute is here}} + [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; #endif // TRIGGER_ERROR diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index 7402248de76f2..4ddfcc259e1ba 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -48,8 +48,11 @@ class Functor16 { #ifdef TRIGGER_ERROR class Functor32 { public: - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}} - [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // 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 {{previous attribute is here}} + [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; #endif diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index c7291e3a579bf..cf09a7bc91362 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -34,11 +34,13 @@ struct TRIFuncObjGood2 { }; struct TRIFuncObjGood3 { - [[sycl::reqd_work_group_size(8, 8)]] void // expected-note {{previous attribute is here}} + [[sycl::reqd_work_group_size(8, 8)]] void // expected-note {{previous attribute is here}} \ + // expected-note {{conflicting attribute is here}} operator()() const; }; -[[sycl::reqd_work_group_size(4, 4)]] // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[sycl::reqd_work_group_size(4, 4)]] // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ + // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} void TRIFuncObjGood3::operator()() const {} @@ -91,31 +93,38 @@ void instantiate() { // equal or greater than values coming from reqd_work_group_size attribute. [[sycl::reqd_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f10() {} // OK -[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} \ + // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} \ + // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} f11(){}; [[sycl::reqd_work_group_size(32, 32, 1)]] [[sycl::reqd_work_group_size(32, 32)]] void f12() {} // OK // Test that template redeclarations also get diagnosed properly. template -[[sycl::reqd_work_group_size(64, 1, 1)]] void f13(); // #f13conflict - +[[sycl::reqd_work_group_size(64, 1, 1)]] void f13(); // #f13conflict \ + // #f13prev template -[[sycl::reqd_work_group_size(X, Y, Z)]] void f13() {} // #f13 +[[sycl::reqd_work_group_size(X, Y, Z)]] void f13() {} // #f13err \ + // #f13warn void test() { f13<64, 1, 1>(); // OK, args are the same on the redecl. - // expected-error@#f13 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // expected-error@#f13err {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // expected-warning@#f13warn {{attribute 'reqd_work_group_size' is already applied with different arguments}} // expected-note@#f13conflict {{conflicting attribute is here}} + // expected-note@#f13prev {{previous attribute is here}} f13<1, 1, 64>(); // expected-note {{in instantiation}} } struct TRIFuncObjBad { - [[sycl::reqd_work_group_size(32, 1, 1)]] void // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(32, 1, 1)]] void // expected-note {{conflicting attribute is here}} \ + // expected-note {{previous attribute is here}} operator()() const; }; -[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} \ + // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} void TRIFuncObjBad::operator()() const {} From ea792e3e891524e7fa4b6c6131e8a7065a1815da Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 14 Mar 2022 18:04:43 -0700 Subject: [PATCH 13/23] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index a083d7340342c..ca7d86af4d261 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3870,8 +3870,9 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // If the 'reqd_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 greater than values of 'reqd_work_group_size' attribute arguments. + // check to see if values of 'reqd_work_group_size' attribute arguments + // are equal or greater than values of 'reqd_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 @@ -3883,7 +3884,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (checkWorkGroupSizeAttrValues>( *this, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { + 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); From 88ce3de486cb40d2928ad91c931fbe58ebfa4623 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 15 Mar 2022 11:30:58 -0700 Subject: [PATCH 14/23] Update comments Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 118 ++++++++++++++++++-------------- 1 file changed, 68 insertions(+), 50 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index ca7d86af4d261..394ce829da650 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3407,15 +3407,18 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, ZDimExpr->getResultAsAPSInt() != 1)); } -// 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 arguments are -// equal or less than values of max_work_group_size attribute arguments. -// If the 'reqd_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 arguments are -// equal or greater than values of 'reqd_work_group_size' attribute arguments. - +// 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 greater than values of +// 'max_work_group_size' attribute arguments. +// +// If the 'reqd_work_group_size' attribute is specified multiple times on a +// declaration, check if the values of 'reqd_work_group_size' attribute +// arguments specified earlier are less than the values of +// 'reqd_work_group_size' attribute arguments specified after. template static bool checkWorkGroupSizeAttrValues( Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, @@ -3432,10 +3435,9 @@ static bool checkWorkGroupSizeAttrValues( !WGSYDimExpr || !WGSZDimExpr) return false; - // Otherwise, check if value of 'reqd_work_group_size' attribute argument - // is greater than value of 'max_work_group_size' attribute argument. - // or check if value of 'reqd_work_group_size' attribute argument is less - // than value of 'reqd_work_group_size' attribute argument. + // Otherwise, compare the first set of X, Y, and Z dimension values with the + // second set of values passed in. The comparison is made using the operator + // passed as the template parameter. bool CheckFirstArgument = S.getLangOpts().OpenCL @@ -3491,10 +3493,14 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, if (!XDim || !YDim || !ZDim) return; - // 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. + // 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. The + // comparison is made using the operator passed as the template parameter. // // 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 @@ -3576,10 +3582,14 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( 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. + // 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. The + // comparison is made using the operator passed as the template parameter. // // 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 @@ -3625,10 +3635,10 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, } // 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. +// 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 @@ -3698,10 +3708,14 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } } - // 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. + // 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. The + // comparison is made using the operator passed as the template parameter. // // 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 @@ -3709,8 +3723,8 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // 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. + // mode. All spellings of reqd_work_group_size attribute (regardless of + // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { if (checkWorkGroupSizeAttrValues>( *this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), @@ -3723,9 +3737,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, } // 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. + // 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 @@ -3757,11 +3771,11 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Diag(Existing->getLoc(), diag::note_previous_attribute); } - // If the 'reqd_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 greater than values of 'reqd_work_group_size' attribute - // arguments. + // If the 'reqd_work_group_size' attribute is specified multiple times on a + // declaration, check if the values of 'reqd_work_group_size' attribute + // arguments specified earlier are less than the values of + // 'reqd_work_group_size' attribute arguments specified after. The + // comparison is made using the operator passed as the template parameter. // // 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 @@ -3807,10 +3821,14 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { } } - // 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. + // 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. The + // comparison is made using the operator passed as the template parameter. // // 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 @@ -3818,8 +3836,8 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // 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. + // mode. All spellings of reqd_work_group_size attribute (regardless of + // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { if (checkWorkGroupSizeAttrValues>( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), @@ -3868,11 +3886,11 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { Diag(A.getLoc(), diag::note_previous_attribute); } - // If the 'reqd_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 greater than values of 'reqd_work_group_size' attribute - // arguments. + // If the 'reqd_work_group_size' attribute is specified multiple times on a + // declaration, check if the values of 'reqd_work_group_size' attribute + // arguments specified earlier are less than the values of + // 'reqd_work_group_size' attribute arguments specified after. The + // comparison is made using the operator passed as the template parameter. // // 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 From 6acca53d632d834dd6210f110d4297da31082b50 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 15 Mar 2022 11:44:38 -0700 Subject: [PATCH 15/23] remov extra space Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 394ce829da650..1336cfd12ddf8 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3417,7 +3417,7 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // // If the 'reqd_work_group_size' attribute is specified multiple times on a // declaration, check if the values of 'reqd_work_group_size' attribute -// arguments specified earlier are less than the values of +// arguments specified earlier are less than the values of // 'reqd_work_group_size' attribute arguments specified after. template static bool checkWorkGroupSizeAttrValues( From 77f4283fbbfb0148ae219969500b062bf9c996b0 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 06:41:59 -0700 Subject: [PATCH 16/23] update comments and fix regression when reqd_work_group_size attribute is applied on multiple times on a given declaration Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 32 ++++++++++++-------- clang/test/SemaSYCL/reqd_work_group_size.cpp | 20 +++++++++--- 2 files changed, 34 insertions(+), 18 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index fd79be514a808..d3bc1d7436d42 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3383,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. @@ -3417,7 +3422,7 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // // If the 'reqd_work_group_size' attribute is specified multiple times on a // declaration, check if the values of 'reqd_work_group_size' attribute -// arguments specified earlier are less than the values of +// arguments specified earlier are less than or greater than the values of // 'reqd_work_group_size' attribute arguments specified after. template static bool checkWorkGroupSizeAttrValues( @@ -3628,6 +3633,7 @@ 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), @@ -3773,7 +3779,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // If the 'reqd_work_group_size' attribute is specified multiple times on a // declaration, check if the values of 'reqd_work_group_size' attribute - // arguments specified earlier are less than the values of + // arguments specified earlier are less than or greater than the values of // 'reqd_work_group_size' attribute arguments specified after. The // comparison is made using the operator passed as the template parameter. // @@ -3785,9 +3791,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // __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. - if (checkWorkGroupSizeAttrValues>( - *this, XDim, YDim, ZDim, Existing->getXDim(), Existing->getYDim(), - Existing->getZDim())) { + if (checkWorkGroupSizeAttrValues>( + *this, Existing->getXDim(), Existing->getYDim(), + Existing->getZDim(), XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) << CI << Existing; Diag(Existing->getLoc(), diag::note_conflicting_attribute); @@ -3888,7 +3894,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // If the 'reqd_work_group_size' attribute is specified multiple times on a // declaration, check if the values of 'reqd_work_group_size' attribute - // arguments specified earlier are less than the values of + // arguments specified earlier are less than or greater than the values of // 'reqd_work_group_size' attribute arguments specified after. The // comparison is made using the operator passed as the template parameter. // @@ -3900,7 +3906,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // __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. - if (checkWorkGroupSizeAttrValues>( + if (checkWorkGroupSizeAttrValues>( *this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index cf09a7bc91362..00c8facfbf818 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -15,19 +15,25 @@ [[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-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[sycl::reqd_work_group_size(6, 6, 6)]] // expected-note {{previous attribute is here}} \ + // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(16, 16, 16)]] void // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ + // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} f4() {} // Catch the easy case where the attributes are all specified at once with // different arguments. struct TRIFuncObjGood1 { + // expected-note@+4 {{conflicting attribute is here}} + // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note@+2 {{previous attribute is here}} // expected-warning@+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@+4 {{conflicting attribute is here}} + // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note@+2 {{previous attribute is here}} // expected-warning@+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 {} @@ -48,13 +54,14 @@ TRIFuncObjGood3::operator()() const {} 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-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} + [[sycl::reqd_work_group_size(16, 16, 16)]] [[sycl::reqd_work_group_size(32, 32, 32)]] void operator()(int) const; // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note {{conflicting 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)]] [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()(int) const; // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} expected-note {{conflicting attribute is here}} + [[intel::max_work_group_size(16, 16, 16)]] // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()(int) const; // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} }; // Ensure that template arguments behave appropriately based on instantiations. @@ -63,7 +70,8 @@ template // Test that template redeclarations also get diagnosed properly. template -[[sycl::reqd_work_group_size(1, 1, 1)]] void f7(); // #f7prev +[[sycl::reqd_work_group_size(1, 1, 1)]] void f7(); // #f7prev \ + // #f7conflict template [[sycl::reqd_work_group_size(X, Y, Z)]] void f7() {} // #f7 @@ -82,6 +90,8 @@ void instantiate() { // 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-note@#f7conflict {{conflicting attribute is here}} + // expected-error@#f7 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-warning@#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}} From bc1da4e5db4e3d1f2f093896e07de27020eeef2b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 06:56:18 -0700 Subject: [PATCH 17/23] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 9 +++++---- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index c991ba29cdeb2..7f0715c2e3eea 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3793,7 +3793,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // (regardless of syntax used) follow the SYCL rules when in SYCL mode. if (checkWorkGroupSizeAttrValues>( *this, Existing->getXDim(), Existing->getYDim(), - Existing->getZDim(), XDim, YDim, ZDim)) { + Existing->getZDim(), XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) << CI << Existing; Diag(Existing->getLoc(), diag::note_conflicting_attribute); diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 00c8facfbf818..77e02a00c0525 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -24,8 +24,8 @@ f4() {} // Catch the easy case where the attributes are all specified at once with // different arguments. struct TRIFuncObjGood1 { - // expected-note@+4 {{conflicting attribute is here}} - // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} + // expected-note@+4 {{conflicting attribute is here}} + // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note@+2 {{previous attribute is here}} // expected-warning@+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 {} @@ -60,8 +60,9 @@ class Functor { 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 operator()(int) const; // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} + [[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. From a2d1c238bfaf955d43c1c7869c19b31a6d0caf70 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 09:34:19 -0700 Subject: [PATCH 18/23] address @elizabeth review comments Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 114 +++++------------- .../redeclaration-attribute-propagation.cpp | 7 +- ...eqd-work-group-size-device-direct-prop.cpp | 7 +- .../SemaSYCL/reqd-work-group-size-device.cpp | 7 +- clang/test/SemaSYCL/reqd_work_group_size.cpp | 69 ++++------- 5 files changed, 59 insertions(+), 145 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 7f0715c2e3eea..920f71941b4a4 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3417,50 +3417,42 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // // 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 greater than values of +// 'reqd_work_group_size' attribute arguments are equal and less than values of // 'max_work_group_size' attribute arguments. -// -// If the 'reqd_work_group_size' attribute is specified multiple times on a -// declaration, check if the values of 'reqd_work_group_size' attribute -// arguments specified earlier are less than or greater than the values of -// 'reqd_work_group_size' attribute arguments specified after. -template static bool checkWorkGroupSizeAttrValues( Sema &S, const Expr *RWGSXDim, const Expr *RWGSYDim, const Expr *RWGSZDim, - const Expr *WGSXDim, const Expr *WGSYDim, const Expr *WGSZDim) { + 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 *RWGSXDimExpr = dyn_cast(RWGSXDim); const auto *RWGSYDimExpr = dyn_cast(RWGSYDim); const auto *RWGSZDimExpr = dyn_cast(RWGSZDim); - const auto *WGSXDimExpr = dyn_cast(WGSXDim); - const auto *WGSYDimExpr = dyn_cast(WGSYDim); - const auto *WGSZDimExpr = dyn_cast(WGSZDim); + const auto *MWGSXDimExpr = dyn_cast(MWGSXDim); + const auto *MWGSYDimExpr = dyn_cast(MWGSYDim); + const auto *MWGSZDimExpr = dyn_cast(MWGSZDim); - if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !WGSXDimExpr || - !WGSYDimExpr || !WGSZDimExpr) + if (!RWGSXDimExpr || !RWGSYDimExpr || !RWGSZDimExpr || !MWGSXDimExpr || + !MWGSYDimExpr || !MWGSZDimExpr) return false; - // Otherwise, compare the first set of X, Y, and Z dimension values with the - // second set of values passed in. The comparison is made using the operator - // passed as the template parameter. - + // 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 - ? Comparator()(RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), - WGSZDimExpr->getResultAsAPSInt().getZExtValue()) - : Comparator()(RWGSXDimExpr->getResultAsAPSInt().getZExtValue(), - WGSXDimExpr->getResultAsAPSInt().getZExtValue()); + ? RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSZDimExpr->getResultAsAPSInt().getZExtValue() + : RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSXDimExpr->getResultAsAPSInt().getZExtValue(); bool CheckSecondArgument = - Comparator()(RWGSYDimExpr->getResultAsAPSInt().getZExtValue(), - WGSYDimExpr->getResultAsAPSInt().getZExtValue()); + RWGSYDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSYDimExpr->getResultAsAPSInt().getZExtValue(); bool CheckThirdArgument = S.getLangOpts().OpenCL - ? Comparator()(RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), - WGSXDimExpr->getResultAsAPSInt().getZExtValue()) - : Comparator()(RWGSZDimExpr->getResultAsAPSInt().getZExtValue(), - WGSZDimExpr->getResultAsAPSInt().getZExtValue()); + ? RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSXDimExpr->getResultAsAPSInt().getZExtValue() + : RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSZDimExpr->getResultAsAPSInt().getZExtValue(); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3504,8 +3496,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // 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. The - // comparison is made using the operator passed as the template parameter. + // are greater 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 @@ -3517,7 +3508,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( + if (checkWorkGroupSizeAttrValues( *this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3593,8 +3584,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // 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. The - // comparison is made using the operator passed as the template parameter. + // are greater 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 @@ -3606,7 +3596,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( + if (checkWorkGroupSizeAttrValues( *this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3720,8 +3710,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // 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. The - // comparison is made using the operator passed as the template parameter. + // are greater 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 @@ -3732,7 +3721,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute (regardless of // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( + if (checkWorkGroupSizeAttrValues( *this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3773,30 +3762,8 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // 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::warn_duplicate_attribute) << CI; + Diag(CI.getLoc(), diag::err_duplicate_attribute) << CI; Diag(Existing->getLoc(), diag::note_previous_attribute); - } - - // If the 'reqd_work_group_size' attribute is specified multiple times on a - // declaration, check if the values of 'reqd_work_group_size' attribute - // arguments specified earlier are less than or greater than the values of - // 'reqd_work_group_size' attribute arguments specified after. The - // comparison is made using the operator passed as the template parameter. - // - // 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. - if (checkWorkGroupSizeAttrValues>( - *this, Existing->getXDim(), Existing->getYDim(), - Existing->getZDim(), XDim, YDim, ZDim)) { - Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) - << CI << Existing; - Diag(Existing->getLoc(), diag::note_conflicting_attribute); return; } @@ -3833,8 +3800,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // 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. The - // comparison is made using the operator passed as the template parameter. + // are greater 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 @@ -3845,7 +3811,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // mode. All spellings of reqd_work_group_size attribute (regardless of // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues>( + if (checkWorkGroupSizeAttrValues( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3888,30 +3854,8 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // 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::warn_duplicate_attribute) << &A; + Diag(DeclAttr->getLoc(), diag::err_duplicate_attribute) << &A; Diag(A.getLoc(), diag::note_previous_attribute); - } - - // If the 'reqd_work_group_size' attribute is specified multiple times on a - // declaration, check if the values of 'reqd_work_group_size' attribute - // arguments specified earlier are less than or greater than the values of - // 'reqd_work_group_size' attribute arguments specified after. The - // comparison is made using the operator passed as the template parameter. - // - // 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. - if (checkWorkGroupSizeAttrValues>( - *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); return nullptr; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 6165134edeb8a..1a5357a9f0ab2 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -24,10 +24,9 @@ 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)]] void func3(); // expected-note {{previous attribute is here}} \ - // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ - // 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}} [[intel::max_work_group_size(8, 8, 8)]] void func4() {} // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} 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 6653b8fefe781..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 @@ -25,11 +25,8 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro // 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 {{previous 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 diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index 4ddfcc259e1ba..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 {{previous 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 diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 77e02a00c0525..4ccb53b819bfc 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -15,38 +15,30 @@ [[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}} \ - // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16, 16, 16)]] void // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ - // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[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@+4 {{conflicting attribute is here}} - // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note@+2 {{previous attribute is here}} - // expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // 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@+4 {{conflicting attribute is here}} - // expected-error@+3 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note@+2 {{previous attribute is here}} - // expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // 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}} \ - // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(8, 8)]] void // expected-note {{previous attribute is here}} operator()() const; }; -[[sycl::reqd_work_group_size(4, 4)]] // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ - // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} void TRIFuncObjGood3::operator()() const {} @@ -54,7 +46,7 @@ TRIFuncObjGood3::operator()() const {} 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-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} // expected-note {{conflicting attribute is here}} + [[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 { @@ -71,8 +63,7 @@ template // Test that template redeclarations also get diagnosed properly. template -[[sycl::reqd_work_group_size(1, 1, 1)]] void f7(); // #f7prev \ - // #f7conflict +[[sycl::reqd_work_group_size(1, 1, 1)]] void f7(); // #f7prev template [[sycl::reqd_work_group_size(X, Y, Z)]] void f7() {} // #f7 @@ -82,7 +73,7 @@ template 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-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[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 @@ -91,51 +82,37 @@ void instantiate() { // 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-note@#f7conflict {{conflicting attribute is here}} - // expected-error@#f7 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - // expected-warning@#f7 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // 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}} } -// If the reqd_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 arguments are -// equal or greater than values coming from reqd_work_group_size attribute. -[[sycl::reqd_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f10() {} // OK +// Tests for 'reqd_work_group_size' attribute duplication. -[[sycl::reqd_work_group_size(8)]] // expected-note {{conflicting attribute is here}} \ - // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} \ - // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} -f11(){}; +[[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 f12() {} // OK +[[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 f13(); // #f13conflict \ - // #f13prev +[[sycl::reqd_work_group_size(64, 1, 1)]] void f10(); // #f10prev template -[[sycl::reqd_work_group_size(X, Y, Z)]] void f13() {} // #f13err \ - // #f13warn +[[sycl::reqd_work_group_size(X, Y, Z)]] void f10() {} // #f10err void test() { - f13<64, 1, 1>(); // OK, args are the same on the redecl. - // expected-error@#f13err {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - // expected-warning@#f13warn {{attribute 'reqd_work_group_size' is already applied with different arguments}} - // expected-note@#f13conflict {{conflicting attribute is here}} - // expected-note@#f13prev {{previous attribute is here}} - f13<1, 1, 64>(); // expected-note {{in instantiation}} + 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 {{conflicting attribute is here}} \ - // expected-note {{previous attribute is here}} + [[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 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} \ - // expected-warning {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[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 {} From 16d2aec75f95a564fcc10c7e3e559bd10ec741cc Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 09:45:53 -0700 Subject: [PATCH 19/23] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 31 +++++++++++++++---------------- 1 file changed, 15 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 920f71941b4a4..a2864b04eba8a 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3439,20 +3439,19 @@ static bool checkWorkGroupSizeAttrValues( bool CheckFirstArgument = S.getLangOpts().OpenCL ? RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > - MWGSZDimExpr->getResultAsAPSInt().getZExtValue() + MWGSZDimExpr->getResultAsAPSInt().getZExtValue() : RWGSXDimExpr->getResultAsAPSInt().getZExtValue() > - MWGSXDimExpr->getResultAsAPSInt().getZExtValue(); + MWGSXDimExpr->getResultAsAPSInt().getZExtValue(); - bool CheckSecondArgument = - RWGSYDimExpr->getResultAsAPSInt().getZExtValue() > - MWGSYDimExpr->getResultAsAPSInt().getZExtValue(); + bool CheckSecondArgument = RWGSYDimExpr->getResultAsAPSInt().getZExtValue() > + MWGSYDimExpr->getResultAsAPSInt().getZExtValue(); bool CheckThirdArgument = S.getLangOpts().OpenCL ? RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > - MWGSXDimExpr->getResultAsAPSInt().getZExtValue() + MWGSXDimExpr->getResultAsAPSInt().getZExtValue() : RWGSZDimExpr->getResultAsAPSInt().getZExtValue() > - MWGSZDimExpr->getResultAsAPSInt().getZExtValue(); + MWGSZDimExpr->getResultAsAPSInt().getZExtValue(); return CheckFirstArgument || CheckSecondArgument || CheckThirdArgument; } @@ -3508,9 +3507,9 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues( - *this, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim(), XDim, YDim, ZDim)) { + if (checkWorkGroupSizeAttrValues(*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); @@ -3596,9 +3595,9 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues( - *this, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { + if (checkWorkGroupSizeAttrValues(*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); @@ -3721,9 +3720,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute (regardless of // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues( - *this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), - DeclAttr->getZDim())) { + if (checkWorkGroupSizeAttrValues(*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); From 34ba9309fef5a2f0a5bd7782d5822906831d88b9 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 12:18:27 -0700 Subject: [PATCH 20/23] Update tests as per review comments Signed-off-by: Soumi Manna --- .../intel-max-global-work-dim-device.cpp | 45 ++++++++----------- ...eqd-work-group-size-device-direct-prop.cpp | 32 ++++++------- .../intel-reqd-work-group-size-device.cpp | 39 +++------------- .../SemaSYCL/num_simd_work_items_device.cpp | 31 +++++-------- clang/test/SemaSYCL/reqd_work_group_size.cpp | 12 +++++ 5 files changed, 61 insertions(+), 98 deletions(-) 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 21b868b84f742..4387b764ecb5f 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -119,21 +119,13 @@ struct TRIFuncObjBad1 { [[intel::max_global_work_dim(0)]] void TRIFuncObjBad1::operator()() const {} -struct TRIFuncObjBad2 { - [[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 TRIFuncObjBad2::operator()() const {} - // 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'}} @@ -141,88 +133,88 @@ 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 { +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 { +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() { @@ -460,9 +452,8 @@ int main() { h.single_task(TRIFuncObjBad12()); h.single_task(TRIFuncObjBad13()); h.single_task(TRIFuncObjBad14()); - h.single_task(TRIFuncObjBad15()); - 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/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp index d9ab0c294d812..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 @@ -30,18 +30,6 @@ class Functor32 { // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} }; - -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 {} -}; #endif // TRIGGER_ERROR class Functor16 { @@ -109,13 +97,19 @@ int main() { FunctorAttr fattr; h.single_task(fattr); -#ifdef TRIGGER_ERROR - Functor33 f33; - h.single_task(f33); - - Functor30 f30; - h.single_task(f30); -#endif // TRIGGER_ERROR + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + 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 f3b0a434262de..1dea68a817583 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -33,19 +33,7 @@ void bar() { [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR -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 {} -}; - [[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} - -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 {} -}; #endif // TRIGGER_ERROR class Functor16 { @@ -53,11 +41,6 @@ class Functor16 { [[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 {} @@ -96,38 +79,30 @@ int main() { FunctorAttr fattr; h.single_task(fattr); -#ifdef TRIGGER_ERROR - Functor33 f33; - h.single_task(f33); - - Functor30 f30; - h.single_task(f30); -#endif // TRIGGER_ERROR - - 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(); }); @@ -180,7 +155,7 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 128 // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name7 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 // CHECK: ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 32 diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 08b44c80ed9ac..19b7508bd297b 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -181,54 +181,47 @@ struct TRIFuncObjBad13 { }; 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 requires a positive integral compile time constant expression}} - 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 TRIFuncObjBad18 { +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 TRIFuncObjBad19 { +struct TRIFuncObjBad18 { [[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 -TRIFuncObjBad19::operator()() const {} +TRIFuncObjBad18::operator()() const {} -struct TRIFuncObjBad20 { +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 -TRIFuncObjBad20::operator()() const {} +TRIFuncObjBad19::operator()() const {} #endif // TRIGGER_ERROR // If the declaration has a [[sycl::reqd_work_group_size()]] // or [[cl::reqd_work_group_size()]] or @@ -420,11 +413,9 @@ int main() { h.single_task(TRIFuncObjBad19()); - h.single_task(TRIFuncObjBad20()); - #endif // TRIGGER_ERROR - h.single_task(TRIFuncObjGood5()); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel31 + h.single_task(TRIFuncObjGood5()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel30 // CHECK: SYCLIntelNumSimdWorkItemsAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 5 @@ -440,8 +431,8 @@ int main() { // CHECK-NEXT: value: Int 5 // CHECK-NEXT: IntegerLiteral{{.*}}5{{$}} - h.single_task(TRIFuncObjGood6()); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel32 + h.single_task(TRIFuncObjGood6()); + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel31 // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 3 diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 4ccb53b819bfc..318cbf4efeeb0 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -6,6 +6,18 @@ [[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() {} From 5b924a0c5b2eb719d9b1cdb5994280bd024ea0e8 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 12:36:09 -0700 Subject: [PATCH 21/23] Fix format errors Signed-off-by: Soumi Manna --- .../SemaSYCL/intel-max-global-work-dim-device.cpp | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) 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 4387b764ecb5f..1dc5ccca9a53c 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -157,7 +157,8 @@ struct TRIFuncObjBad6 { }; [[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 TRIFuncObjBad6::operator()() const {} +void +TRIFuncObjBad6::operator()() const {} struct TRIFuncObjBad7 { [[intel::max_global_work_dim(0)]] void @@ -165,7 +166,8 @@ struct TRIFuncObjBad7 { }; [[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 +TRIFuncObjBad7::operator()() const {} struct TRIFuncObjBad8 { [[intel::max_global_work_dim(0)]] void @@ -173,7 +175,8 @@ struct TRIFuncObjBad8 { }; [[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 TRIFuncObjBad8::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. @@ -454,7 +457,7 @@ int main() { h.single_task(TRIFuncObjBad14()); h.single_task( - []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} + []() [[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; From 1137c2d48c87431a8da942312f78aa0bb08753fe Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 18 Mar 2022 12:50:09 -0700 Subject: [PATCH 22/23] update function name to make it more readable Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index a2864b04eba8a..b617aab76ffe9 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3419,7 +3419,7 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // with 'max_work_group_size' attribute, check to see if values of // 'reqd_work_group_size' attribute arguments are equal and less than values of // 'max_work_group_size' attribute arguments. -static bool checkWorkGroupSizeAttrValues( +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. @@ -3507,7 +3507,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), + if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), XDim, YDim, ZDim)) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3595,7 +3595,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // __attribute__((reqd_work_group_size)) is only available in OpenCL mode // and follows the OpenCL rules. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, DeclAttr->getXDim(), + if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), A.getXDim(), A.getYDim(), A.getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3720,7 +3720,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // mode. All spellings of reqd_work_group_size attribute (regardless of // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues(*this, XDim, YDim, ZDim, + if (checkMaxAllowedWorkGroupSize(*this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) @@ -3810,7 +3810,7 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // mode. All spellings of reqd_work_group_size attribute (regardless of // syntax used) follow the SYCL rules when in SYCL mode. if (const auto *DeclAttr = D->getAttr()) { - if (checkWorkGroupSizeAttrValues( + if (checkMaxAllowedWorkGroupSize( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())) { Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) From f376e43452ffd1a505bb30db76e576eb766f8a4a Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 21 Mar 2022 09:20:32 -0700 Subject: [PATCH 23/23] Update comments --- clang/lib/Sema/SemaDeclAttr.cpp | 128 ++++++++------------------------ 1 file changed, 32 insertions(+), 96 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index b617aab76ffe9..93706678dc8ca 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3417,8 +3417,17 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, // // 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 and less than values of -// 'max_work_group_size' attribute arguments. +// '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) { @@ -3491,21 +3500,11 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, // 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 + // '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. - // - // 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 *DeclAttr = D->getAttr()) { if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), @@ -3518,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)) { @@ -3579,21 +3578,11 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( // 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 + // '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. - // - // 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 *DeclAttr = D->getAttr()) { if (checkMaxAllowedWorkGroupSize(*this, DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim(), @@ -3605,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())) { @@ -3639,6 +3627,10 @@ static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, // 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. @@ -3657,7 +3649,7 @@ static bool CheckWorkGroupSize(Sema &S, const Expr *NSWIValue, : (RWGSZDimExpr->getResultAsAPSInt()).getZExtValue(); // Check if the required work group size specified by 'num_simd_work_items' - // attribute must evenly divide the index that increments fastest in the + // attribute evenly divides the index that increments fastest in the // 'reqd_work_group_size' attribute. return WorkGroupSize % NSWIValueExpr->getResultAsAPSInt().getZExtValue() != 0; } @@ -3694,7 +3686,7 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, return; // If the declaration has a ReqdWorkGroupSizeAttr, 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)) { @@ -3705,20 +3697,11 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // 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 + // '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. - // - // 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. if (const auto *DeclAttr = D->getAttr()) { if (checkMaxAllowedWorkGroupSize(*this, XDim, YDim, ZDim, DeclAttr->getXDim(), DeclAttr->getYDim(), @@ -3734,15 +3717,6 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, // 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. 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) @@ -3780,10 +3754,9 @@ void Sema::AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, ReqdWorkGroupSizeAttr * Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { - // If the declaration has a ReqdWorkGroupSizeAttr, - // 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 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())) { @@ -3800,15 +3773,6 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // // We emit diagnostic if values of 'reqd_work_group_size' attribute arguments // are greater 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. if (const auto *DeclAttr = D->getAttr()) { if (checkMaxAllowedWorkGroupSize( *this, A.getXDim(), A.getYDim(), A.getZDim(), DeclAttr->getXDim(), @@ -3824,15 +3788,6 @@ Sema::MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A) { // 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. if (const auto *DeclAttr = D->getAttr()) { if (CheckWorkGroupSize(*this, DeclAttr->getValue(), A.getXDim(), A.getZDim())) { @@ -4048,15 +4003,6 @@ void Sema::AddSYCLIntelNumSimdWorkItemsAttr(Decl *D, // 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. if (const auto *DeclAttr = D->getAttr()) { if (CheckWorkGroupSize(*this, E, DeclAttr->getXDim(), DeclAttr->getZDim())) { @@ -4092,15 +4038,6 @@ SYCLIntelNumSimdWorkItemsAttr *Sema::MergeSYCLIntelNumSimdWorkItemsAttr( // 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. if (const auto *DeclAttr = D->getAttr()) { if (CheckWorkGroupSize(*this, A.getValue(), DeclAttr->getXDim(), DeclAttr->getZDim())) { @@ -4344,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) || @@ -4376,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,