diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a3c6547b0ae55..ddf687a3c8a22 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1406,17 +1406,20 @@ def SYCLIntelMaxWorkGroupSize : InheritableAttr { let LangOpts = [SYCLIsDevice, SYCLIsHost]; 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 = [SYCLIntelMaxWorkGroupSizeAttrDocs]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 500b676ca9869..65f5b5f560092 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10559,6 +10559,11 @@ class Sema final { const SYCLUsesAspectsAttr &A); void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI, Expr **Exprs, unsigned Size); + void AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *XDim, Expr *YDim, Expr *ZDim); + SYCLIntelMaxWorkGroupSizeAttr * + MergeSYCLIntelMaxWorkGroupSizeAttr(Decl *D, + const SYCLIntelMaxWorkGroupSizeAttr &A); /// AddAlignedAttr - Adds an aligned attribute to a particular declaration. void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E, bool IsPackExpansion); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 0f1f9578e086e..ad96b01545366 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -743,24 +743,16 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, if (const SYCLIntelMaxWorkGroupSizeAttr *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 SYCLIntelMaxWorkGroupSizeAttr arguments are reversed. - if (getLangOpts().SYCLIsDevice) - std::swap(XDimVal, ZDimVal); - - llvm::Metadata *AttrMDArgs[] = { - llvm::ConstantAsMetadata::get( - Builder.getInt32(XDimVal->getZExtValue())), - llvm::ConstantAsMetadata::get( - Builder.getInt32(YDimVal->getZExtValue())), - llvm::ConstantAsMetadata::get( - Builder.getInt32(ZDimVal->getZExtValue()))}; - Fn->setMetadata("max_work_group_size", - llvm::MDNode::get(Context, AttrMDArgs)); + // 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("max_work_group_size", + llvm::MDNode::get(Context, AttrMDArgs)); + } } if (const auto *A = FD->getAttr()) { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 72eac772af5b1..bd1209a3f0df0 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2732,6 +2732,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, NewAttr = S.MergeSYCLUsesAspectsAttr(D, *A); else if (const auto *A = dyn_cast(Attr)) NewAttr = S.MergeSYCLIntelPipeIOAttr(D, *A); + else if (const auto *A = dyn_cast(Attr)) + NewAttr = S.MergeSYCLIntelMaxWorkGroupSizeAttr(D, *A); else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr)) NewAttr = cast(Attr->clone(S.Context)); @@ -3433,8 +3435,6 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, checkDimensionsAndSetDiagnostics(*this, New, Old); - 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 3d3c868930b02..213e0daf1d49c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3205,12 +3205,9 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { ASTContext &Ctx = S.getASTContext(); 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)))) { + if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()) && + (getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal()) && + (getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) { S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes) << AL << A->getSpelling(); Result &= false; @@ -3232,19 +3229,18 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { return Result; } -// Handles reqd_work_group_size and max_work_group_size. +// 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)), [[cl::reqd_work_group_size]], and - // [[intel::max_work_group_size]] all require exactly three arguments. + // __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.getKind() == ParsedAttr::AT_SYCLIntelMaxWorkGroupSize || AL.getSyntax() == ParsedAttr::AS_GNU) { if (!AL.checkExactlyNumArgs(S, 3)) return; @@ -3308,8 +3304,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { } } - // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or - // ReqdWorkGroupSizeAttr, check to see if they hold equal values + // 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()) { @@ -3454,6 +3450,146 @@ static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) { AL.getArgAsExpr(2)); } +// 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. +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. + const auto *MGValueExpr = dyn_cast(MGValue); + const auto *XDimExpr = dyn_cast(XDim); + const auto *YDimExpr = dyn_cast(YDim); + const auto *ZDimExpr = dyn_cast(ZDim); + + if (!MGValueExpr || !XDimExpr || !YDimExpr || !ZDimExpr) + return false; + + // Otherwise, check if the attribute values are equal to one. + return (MGValueExpr->getResultAsAPSInt() == 0 && + (XDimExpr->getResultAsAPSInt() != 1 || + YDimExpr->getResultAsAPSInt() != 1 || + ZDimExpr->getResultAsAPSInt() != 1)); +} + +void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(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 SYCLIntelMaxWorkGroupSizeAttr, 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; + 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) + SYCLIntelMaxWorkGroupSizeAttr(Context, CI, XDim, YDim, ZDim)); +} + +SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr( + Decl *D, const SYCLIntelMaxWorkGroupSizeAttr &A) { + // 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; + } + + // 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 (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; + } + } + + return ::new (Context) SYCLIntelMaxWorkGroupSizeAttr( + Context, A, A.getXDim(), A.getYDim(), A.getZDim()); +} + +static void handleSYCLIntelMaxWorkGroupSize(Sema &S, Decl *D, + const ParsedAttr &AL) { + S.AddSYCLIntelMaxWorkGroupSizeAttr(D, AL, AL.getArgAsExpr(0), + AL.getArgAsExpr(1), AL.getArgAsExpr(2)); +} + void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, Expr *E) { if (!E->isValueDependent()) { @@ -10331,7 +10467,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, handleWorkGroupSize(S, D, AL); break; case ParsedAttr::AT_SYCLIntelMaxWorkGroupSize: - handleWorkGroupSize(S, D, AL); + handleSYCLIntelMaxWorkGroupSize(S, D, AL); break; case ParsedAttr::AT_IntelReqdSubGroupSize: handleIntelReqdSubGroupSize(S, D, AL); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8f45e45dc3f8f..bebd51cbfdad3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3889,9 +3889,9 @@ static void PropagateAndDiagnoseDeviceAttr( } else 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(Ctx) || + *Existing->getYDimVal() < RWGSA->getYDimVal(Ctx) || + *Existing->getZDimVal() < RWGSA->getZDimVal(Ctx)) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); @@ -3909,9 +3909,9 @@ static void PropagateAndDiagnoseDeviceAttr( auto *SIMWGSA = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { ASTContext &Ctx = S.getASTContext(); - if (Existing->getXDimVal(Ctx) > SIMWGSA->getXDimVal(Ctx) || - Existing->getYDimVal(Ctx) > SIMWGSA->getYDimVal(Ctx) || - Existing->getZDimVal(Ctx) > SIMWGSA->getZDimVal(Ctx)) { + if (Existing->getXDimVal(Ctx) > *SIMWGSA->getXDimVal() || + Existing->getYDimVal(Ctx) > *SIMWGSA->getYDimVal() || + Existing->getZDimVal(Ctx) > *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 1aab5134a3371..9faaa06ff43f5 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -814,6 +814,25 @@ static void instantiateWorkGroupSizeHintAttr( ZResult.get()); } +static void instantiateSYCLIntelMaxWorkGroupSizeAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const SYCLIntelMaxWorkGroupSizeAttr *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.AddSYCLIntelMaxWorkGroupSizeAttr(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. @@ -1045,8 +1064,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, } if (const auto *SYCLIntelMaxWorkGroupSize = dyn_cast(TmplAttr)) { - instantiateIntelSYCTripleLFunctionAttr( - *this, TemplateArgs, SYCLIntelMaxWorkGroupSize, New); + instantiateSYCLIntelMaxWorkGroupSizeAttr(*this, TemplateArgs, + SYCLIntelMaxWorkGroupSize, New); continue; } if (const auto *SYCLIntelMaxConcurrency = 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 114538c43b034..2a9d375a010b0 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -1,6 +1,8 @@ // RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// The test checks support and functionality of [[intel::max_global_work_dim()]] attribute. + #include "sycl.hpp" using namespace cl::sycl; @@ -101,20 +103,15 @@ struct TRIFuncObjGood8 { operator()() const {} }; -// FIXME: We do not have support yet for checking -// max_work_group_size and max_global_work_dim -// attributes when merging, so the test compiles without -// any diagnostic when it shouldn't. -struct TRIFuncObjBad { - [[intel::max_work_group_size(4, 4, 4)]] void +struct TRIFuncObjGood9 { + [[intel::max_work_group_size(4, 4, 4)]] void // OK operator()() const; }; -[[intel::max_global_work_dim(0)]] -void TRIFuncObjBad::operator()() const {} +[[intel::max_global_work_dim(1)]] void TRIFuncObjGood9::operator()() const {} // FIXME: We do not have support yet for checking -// max_work_group_size and max_global_work_dim +// reqd_work_group_size and max_global_work_dim // attributes when merging, so the test compiles without // any diagnostic when it shouldn't. struct TRIFuncObjBad1 { @@ -126,7 +123,7 @@ struct TRIFuncObjBad1 { void TRIFuncObjBad1::operator()() const {} // FIXME: We do not have support yet for checking -// max_work_group_size and max_global_work_dim +// 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 { @@ -228,6 +225,13 @@ struct TRIFuncObjBad14 { [[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 { + [[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 {} #endif // TRIGGER_ERROR int main() { @@ -411,7 +415,7 @@ int main() { // CHECK-NEXT: value: Int 3 // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} - h.single_task(TRIFuncObjBad()); + h.single_task(TRIFuncObjGood9()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -425,8 +429,8 @@ int main() { // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 0 - // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} h.single_task(TRIFuncObjBad1()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 @@ -494,6 +498,7 @@ int main() { h.single_task(TRIFuncObjBad12()); h.single_task(TRIFuncObjBad13()); h.single_task(TRIFuncObjBad14()); + h.single_task(TRIFuncObjBad15()); h.single_task( []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp index 4943c9177070b..70b8553f915ff 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp @@ -1,6 +1,7 @@ // RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s +// The test checks support and functionality of [[intel:::max_work_group_size()]] attribute. #include "sycl.hpp" using namespace cl::sycl; @@ -85,55 +86,45 @@ int main() { h.single_task( []() { func_do_not_ignore(); }); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -8 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // expected-warning@+2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} +#ifdef TRIGGER_ERROR h.single_task( - []() [[intel::max_work_group_size(8, 8, -8)]]{}); + []() [[intel::max_work_group_size(8, 8, -8)]] {}); // expected-error{{'max_work_group_size' attribute requires a positive integral compile time constant expression}} - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -8 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int -8 - // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // expected-warning@+2 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} h.single_task( - []() [[intel::max_work_group_size(-8, 8, -8)]]{}); -#ifdef TRIGGER_ERROR + []() [[intel::max_work_group_size(-8, 8, -8)]] {}); // expected-error 2{{'max_work_group_size' attribute requires a positive integral compile time constant expression}} + [[intel::max_work_group_size(1, 1, 1)]] int Var = 0; // expected-error{{'max_work_group_size' attribute only applies to functions}} h.single_task( - []() [[intel::max_work_group_size(0, 1, 3)]]{}); // expected-error{{'max_work_group_size' attribute must be greater than 0}} + []() [[intel::max_work_group_size(0, 1, 3)]] {}); // expected-error{{'max_work_group_size' attribute requires a positive integral compile time constant expression}} h.single_task( []() [[intel::max_work_group_size(1.2f, 1, 3)]]{}); // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} h.single_task( - []() [[intel::max_work_group_size(16, 16, 16), // expected-note{{conflicting attribute is here}} - intel::max_work_group_size(2, 2, 2)]]{}); // expected-warning{{attribute 'max_work_group_size' is already applied with different arguments}} + []() [[intel::max_work_group_size(16, 16, 16), // expected-note{{previous attribute is here}} + intel::max_work_group_size(2, 2, 2)]] {}); // expected-warning{{attribute 'max_work_group_size' is already applied with different arguments}} h.single_task( DAFuncObj()); #endif // TRIGGER_ERROR + // Ignore duplicate attribute. + h.single_task( + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10 + // CHECK: SYCLIntelMaxWorkGroupSizeAttr + // 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: SYCLIntelMaxWorkGroupSizeAttr + []() [[intel::max_work_group_size(2, 2, 2), + intel::max_work_group_size(2, 2, 2)]] {}); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp new file mode 100644 index 0000000000000..8c8088a105813 --- /dev/null +++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp @@ -0,0 +1,61 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// Check the basics. +[[intel::max_work_group_size]] void f(); // expected-error {{'max_work_group_size' attribute requires exactly 3 arguments}} +[[intel::max_work_group_size(12, 12, 12, 12)]] void f0(); // expected-error {{'max_work_group_size' attribute requires exactly 3 arguments}} +[[intel::max_work_group_size("derp", 1, 2)]] void f1(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char[5]'}} +[[intel::max_work_group_size(1, 1, 1)]] int i; // expected-error {{'max_work_group_size' attribute only applies to functions}} + +// Tests for Intel FPGA 'max_work_group_size' attribute duplication. +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +[[intel::max_work_group_size(6, 6, 6)]] [[intel::max_work_group_size(6, 6, 6)]] void f2() {} + +// No diagnostic is emitted because the arguments match. +[[intel::max_work_group_size(32, 32, 32)]] void f3(); +[[intel::max_work_group_size(32, 32, 32)]] void f3(); // OK + +// Produce a conflicting attribute warning when the args are different. +[[intel::max_work_group_size(6, 6, 6)]] // expected-note {{previous attribute is here}} +[[intel::max_work_group_size(16, 16, 16)]] void // expected-warning {{attribute 'max_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. +[[intel::max_work_group_size(16, 16, 16), intel::max_work_group_size(2, 2, 2)]] void f5(); // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} + +// Show that the attribute works on member functions. +class Functor { +public: + [[intel::max_work_group_size(16, 16, 16)]] [[intel::max_work_group_size(16, 16, 16)]] void operator()() const; + [[intel::max_work_group_size(16, 16, 16)]] [[intel::max_work_group_size(32, 32, 32)]] void operator()(int) const; // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} +}; + +// Ensure that template arguments behave appropriately based on instantiations. +template +[[intel::max_work_group_size(N, 1, 1)]] void f6(); // #f6 + +// Test that template redeclarations also get diagnosed properly. +template +[[intel::max_work_group_size(1, 1, 1)]] void f7(); // #f7prev + +template +[[intel::max_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 +[[intel::max_work_group_size(X, 1, Z)]] void f8(); // expected-note {{previous attribute is here}} +template +[[intel::max_work_group_size(X, 2, Z)]] void f8(); // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} + +void instantiate() { + f6<1>(); // OK + // expected-error@#f6 {{'max_work_group_size' attribute requires a positive integral compile time constant expression}} + f6<-1>(); // expected-note {{in instantiation}} + // expected-error@#f6 {{'max_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 'max_work_group_size' is already applied with different arguments}} + // expected-note@#f7prev {{previous attribute is here}} + f7<2, 2, 2>(); // expected-note {{in instantiation}} +} diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 56b70a64b3b7d..ccdbfbfb9f20e 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -2,6 +2,7 @@ // RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2017 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 -Wno-sycl-2017-compat | FileCheck %s +// The test checks redeclaration of [[intel:::max_work_group_size()]] and [[sycl::reqd_work_group_size()]] attributes. #include "sycl.hpp" using namespace cl::sycl; @@ -37,15 +38,9 @@ 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}} -//fourth case - expect error -[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}} -void -func4(); - -[[intel::max_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}} -void -// expected-warning@+1 {{attribute 'max_work_group_size' is already applied with different arguments}} -func4() {} // expected-error {{'max_work_group_size' attribute conflicts with ''max_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}} #endif int main() { 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 52f441a496983..42f907356d3cf 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 @@ -113,6 +113,16 @@ template template [[intel::max_global_work_dim(0)]] void func12(); +template +[[intel::max_global_work_dim(0)]] void func13(); +template +[[intel::max_work_group_size(N, N, N)]] void func13(); // expected-error {{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} + +template +[[intel::max_global_work_dim(0)]] void func14(); +template +[[intel::max_work_group_size(N, N, N)]] void func14(); + int check() { func3<3>(); // OK func3<-1>(); // expected-note {{in instantiation of function template specialization 'func3<-1>' requested here}} @@ -125,13 +135,13 @@ int check() { func10<1>(); // OK func11<1>(); // OK func12<1>(); // OK + func13<6>(); // expected-note {{in instantiation of function template specialization 'func13<6>' requested here}} + func14<1>(); // OK return 0; } // No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. -[[intel::max_global_work_dim(2)]] -[[intel::max_global_work_dim(2)]] void func13() {} - +[[intel::max_global_work_dim(2)]] [[intel::max_global_work_dim(2)]] void func15() {} // CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' // CHECK: TemplateArgument integral 3 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} @@ -141,7 +151,7 @@ int check() { // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}} func13 'void ()' +// CHECK: FunctionDecl {{.*}} {{.*}} func15 'void ()' // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 2 diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-work-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-work-group-size-template.cpp index 09a1bf4f05a33..f5e46467daceb 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-max-work-group-size-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-max-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. +[[intel::max_work_group_size(4, 4, 4)]] [[intel::max_work_group_size(4, 4, 4)]] void func4() {} +// CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()' +// CHECK: SYCLIntelMaxWorkGroupSizeAttr +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 4 +// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 4 +// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 4 +// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} +// CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr