diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 155ce2a21682e..18ef9b2fb7fa8 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11597,7 +11597,7 @@ def err_sycl_non_constant_array_type : Error< def err_conflicting_sycl_kernel_attributes : Error< "conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function">; def err_conflicting_sycl_function_attributes : Error< - "%0 attribute conflicts with '%1' attribute">; + "%0 attribute conflicts with %1 attribute">; def err_sycl_function_attribute_mismatch : Error< "SYCL kernel without %0 attribute can't call a function with this attribute">; def err_sycl_x_y_z_arguments_must_be_one : Error< diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index e2946db2dad3d..6eb1f0f6bdc51 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3268,12 +3268,31 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { 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()) { - if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()) && - (getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal()) && - (getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) { + 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->getSpelling(); + << AL << A; + S.Diag(A->getLocation(), diag::note_conflicting_attribute); Result &= false; } } @@ -3286,7 +3305,8 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) { (getExprValue(AL.getArgAsExpr(2), Ctx) >= getExprValue(A->getZDim(), Ctx)))) { S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes) - << AL << A->getSpelling(); + << AL << A; + S.Diag(A->getLocation(), diag::note_conflicting_attribute); Result &= false; } } @@ -3562,6 +3582,23 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim, ZDimExpr->getResultAsAPSInt() != 1)); } +// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on +// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] +// attribute, check to see if values of reqd_work_group_size arguments are +// equal or less than values of max_work_group_size attribute arguments. +static bool checkWorkGroupSizeAttrValues(const Expr *RWGS, const Expr *MWGS) { + // 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) + 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(); +} + void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XDim, Expr *YDim, @@ -3595,6 +3632,40 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, if (!XDim || !YDim || !ZDim) return; + // If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on + // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] + // attribute, check to see if values of reqd_work_group_size arguments are + // equal or less than values of max_work_group_size attribute arguments. + // + // The arguments to reqd_work_group_size are ordered based on which index + // increments the fastest. In OpenCL, the first argument is the index that + // increments the fastest, and in SYCL, the last argument is the index that + // increments the fastest. + // + // [[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()) { + 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) { + Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes) + << CI << DeclAttr; + Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute); + 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. @@ -3655,6 +3726,40 @@ 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. + // + // 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()) { + 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) { + Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes) + << DeclAttr << &A; + Diag(A.getLoc(), diag::note_conflicting_attribute); + 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 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 70b8553f915ff..a07324db6ced4 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp @@ -35,12 +35,14 @@ struct Func { #ifdef TRIGGER_ERROR struct DAFuncObj { - [[intel::max_work_group_size(4, 4, 4)]] - [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \ - // expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \ - // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}} - void operator()() const {} + [[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}} + [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \ + // expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \ + // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}} + void + operator()() const {} }; + #endif // TRIGGER_ERROR int main() { diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp index 8c8088a105813..524884b00675c 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp @@ -30,6 +30,11 @@ class Functor { [[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}} }; +class FunctorC { +public: + [[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(64, 64, 64)]] void operator()() const; + [[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(16, 16, 16)]] void operator()(int) const; // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} expected-note {{conflicting attribute is here}} +}; // Ensure that template arguments behave appropriately based on instantiations. template [[intel::max_work_group_size(N, 1, 1)]] void f6(); // #f6 @@ -59,3 +64,48 @@ void instantiate() { // expected-note@#f7prev {{previous attribute is here}} f7<2, 2, 2>(); // expected-note {{in instantiation}} } + +// 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 coming from max_work_group_size attribute. +[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}} +[[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +void +f9() {} + +[[intel::max_work_group_size(4, 4, 4)]] void f10(); +[[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK + +[[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. + +[[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}} +f13() {} + +[[intel::max_work_group_size(16, 16, 16)]] void f14(); // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} + +[[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ + // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} +[[intel::max_work_group_size(1, 2, 3)]] void +f15() {} // OK + +[[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} + +[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK + +[[sycl::reqd_work_group_size(16)]] // expected-note {{conflicting attribute is here}} +[[intel::max_work_group_size(1, 1, 16)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} +f18(); + +[[intel::max_work_group_size(16, 16, 1)]] void f19(); +[[sycl::reqd_work_group_size(16, 16)]] void f19(); // OK diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index ccdbfbfb9f20e..af8730ea98700 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -20,13 +20,8 @@ func1(); #else //second case - expect error -[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}} -void -func2(); - -[[sycl::reqd_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}} -void -func2() {} +[[intel::max_work_group_size(4, 4, 4)]] void func2(); // expected-note {{conflicting attribute is here}} +[[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}} @@ -36,7 +31,7 @@ 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}} +func3() {} // 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}} @@ -77,7 +72,7 @@ int main() { #else h.single_task( - []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} + []() { func2(); }); h.single_task( []() { func3(); }); 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 88b8ab6a3a3a1..8bc681b5b690d 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,9 +26,11 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro class Functor32 { public: // expected-note@+3{{conflicting attribute is here}} - // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}} - // 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 {} + // 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 + 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 700f861c7fce9..634257a4aebe1 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -49,9 +49,11 @@ class Functor16 { class Functor32 { public: // expected-note@+3{{conflicting attribute is here}} - // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}} - // 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 {} + // 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 + operator()() const {} }; #endif class Functor16x16x16 {