diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 2e49675b5f68d..1e94b129fc387 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12358,6 +12358,8 @@ def err_bit_cast_type_size_mismatch : Error< "__builtin_bit_cast source size does not equal destination size (%0 vs %1)">; // SYCL-specific diagnostics +def warn_sycl_incorrect_use_attribute_non_kernel_function : Warning< + "%0 attribute can only be applied to a SYCL kernel function">, InGroup; def warn_sycl_kernel_num_of_template_params : Warning< "'sycl_kernel' attribute only applies to a function template with at least" " two template parameters">, InGroup; diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index a7b83937626c2..451d24f087c44 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -259,6 +259,8 @@ class SemaSYCL : public SemaBase { // useful notes that shows where the kernel was called. bool DiagnosingSYCLKernel = false; + llvm::DenseSet SYCLKernelFunctions; + public: SemaSYCL(Sema &S); @@ -300,6 +302,10 @@ class SemaSYCL : public SemaBase { void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); } llvm::SetVector &syclDeviceDecls() { return SyclDeviceDecls; } + void addSYCLKernelFunction(const FunctionDecl *FD) { + SYCLKernelFunctions.insert(FD); + } + /// Lazily creates and returns SYCL integration header instance. SYCLIntegrationHeader &getSyclIntegrationHeader() { if (SyclIntHeader == nullptr) @@ -375,6 +381,8 @@ class SemaSYCL : public SemaBase { SourceLocation Loc, DeviceDiagnosticReason Reason); + void performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD); + /// Tells whether given variable is a SYCL explicit SIMD extension's "private /// global" variable - global variable in the private address space. bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) { diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index cd655ef33f870..ceb89c52c985a 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1909,6 +1909,10 @@ class DeferredDiagnosticsEmitter void checkFunc(SourceLocation Loc, FunctionDecl *FD) { auto &Done = DoneMap[InOMPDeviceContext > 0 ? 1 : 0]; FunctionDecl *Caller = UsePath.empty() ? nullptr : UsePath.back(); + + if (!Caller && S.LangOpts.SYCLIsDevice) + S.SYCL().performSYCLDelayedAttributesAnalaysis(FD); + if ((!ShouldEmitRootNode && !S.getLangOpts().OpenMP && !Caller) || S.shouldIgnoreInHostDeviceCheck(FD) || InUsePath.count(FD)) return; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 85fbe63cf9be3..ebc58fa1fb21f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -882,6 +882,8 @@ class SingleDeviceFunctionTracker { // having a kernel lambda with a lambda call inside of it. KernelBody = CurrentDecl; } + if (KernelBody) + Parent.SemaSYCLRef.addSYCLKernelFunction(KernelBody); } // Recurse. @@ -6852,3 +6854,19 @@ ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc, return BuildUniqueStableNameExpr(OpLoc, LParen, RParen, TSI); } + +void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) { + if (SYCLKernelFunctions.contains(FD)) + return; + + for (const auto *KernelAttr : std::vector{ + FD->getAttr(), + FD->getAttr(), + FD->getAttr(), + FD->getAttr()}) { + if (KernelAttr) + Diag(KernelAttr->getLoc(), + diag::warn_sycl_incorrect_use_attribute_non_kernel_function) + << KernelAttr; + } +} diff --git a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp index a0cbcef3e83be..6737fe7f2aa70 100644 --- a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp +++ b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp @@ -17,7 +17,8 @@ // Produce a conflicting attribute warning when the args are different. [[sycl::work_group_size_hint(4, 1, 1)]] void f3(); // expected-note {{previous attribute is here}} -[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} +[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} \ +// expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} // 1 and 2 dim versions [[sycl::work_group_size_hint(2)]] void f4(); // ok @@ -70,10 +71,13 @@ void instantiate() { f8<0>(); // expected-note {{in instantiation}} #endif + // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} f9<1, 1, 1>(); // OK, args are the same on the redecl. // expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}} // expected-note@#f9prev {{previous attribute is here}} + // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} + f9<1, 2, 3>(); // expected-note {{in instantiation}} } @@ -97,14 +101,14 @@ class Functor16x2x1 { class Functor4x4x4 { public: - [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {}; + [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} }; // Checking whether propagation of the attribute happens or not, according to the SYCL version. #if defined(EXPECT_PROP) // if attribute is propagated, then we expect errors here void f8x8x8(){}; #else // otherwise no error -[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; +[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} #endif class FunctorNoProp { public: diff --git a/clang/test/SemaSYCL/device_has.cpp b/clang/test/SemaSYCL/device_has.cpp index 65dfa82c3ae09..f9557000a23d6 100644 --- a/clang/test/SemaSYCL/device_has.cpp +++ b/clang/test/SemaSYCL/device_has.cpp @@ -17,7 +17,6 @@ enum class aspect { [[sycl::device_has("123")]] void func1() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}} [[sycl::device_has(fake_cl::sycl::aspect::aspect1)]] void func2() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}} - [[sycl::device_has(sycl::aspect::cpu)]] void func3(); // expected-note{{previous attribute is here}} [[sycl::device_has(sycl::aspect::gpu)]] void func3() {} // expected-warning{{attribute 'device_has' is already applied}} diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp index 3eebe18788eae..36d25d34a99eb 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp @@ -70,7 +70,7 @@ void instantiate() { // 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}} +[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} void f9() {} @@ -78,7 +78,7 @@ 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 +[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[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}} @@ -91,14 +91,14 @@ f13() {} [[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?}} + // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[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 +[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(16)]] // expected-note {{conflicting attribute is here}} [[intel::max_work_group_size(16, 1, 1)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp index 218527972128b..e86b68db49cfa 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size.cpp @@ -33,17 +33,19 @@ int main() { return 0; } [[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); -[[intel::reqd_sub_group_size(16)]] void A() { +[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} +{ } -[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { +[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} A(); } // expected-note@+1 {{conflicting attribute is here}} -[[intel::reqd_sub_group_size(2)]] void sg_size2() {} +[[intel::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} -// expected-note@+2 {{conflicting attribute is here}} -// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}} +// expected-note@+3 {{conflicting attribute is here}} +// expected-error@+2 {{conflicting attributes applied to a SYCL kernel}} +// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} [[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() { sg_size2(); } @@ -67,7 +69,7 @@ int main() { // No diagnostic is emitted because the arguments match. [[intel::reqd_sub_group_size(12)]] void same(); -[[intel::reqd_sub_group_size(12)]] void same() {} // OK +[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} // No diagnostic because the attributes are synonyms with identical behavior. [[sycl::reqd_sub_group_size(12)]] void same(); // OK @@ -117,10 +119,12 @@ int check() { // Test that checks template parameter support on function. template -// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +// expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} [[intel::reqd_sub_group_size(N)]] void func3() {} template +// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} [[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}} template diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index c7e951c83f7a5..a2d2760480117 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -22,6 +22,7 @@ class Functor30 { // Tests for 'reqd_work_group_size' attribute duplication. // No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[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. @@ -29,21 +30,23 @@ class Functor30 { [[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(6, 6, 6)]] // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(16, 16, 16)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} f4() {} // Catch the easy case where the attributes are all specified at once with // different arguments. struct TRIFuncObjGood1 { - // expected-note@+2 {{previous attribute is here}} - // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-note@+3 {{previous attribute is here}} + // expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(64)]] [[sycl::reqd_work_group_size(128)]] void operator()() const {} }; struct TRIFuncObjGood2 { - // expected-note@+2 {{previous attribute is here}} - // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-note@+3 {{previous attribute is here}} + // expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}} + // expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(64, 64)]] [[sycl::reqd_work_group_size(128, 128)]] void operator()() const {} }; @@ -52,7 +55,8 @@ struct TRIFuncObjGood3 { operator()() const; }; -[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} +[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} \ +// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} void TRIFuncObjGood3::operator()() const {} @@ -73,7 +77,7 @@ class FunctorC { class Functor32 { public: - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} + [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error{{attribute 'reqd_work_group_size' is already applied with different arguments}} operator()() const {} }; @@ -105,16 +109,18 @@ void instantiate() { f7<1, 1, 1>(); // OK, args are the same on the redecl. // expected-error@#f7 {{attribute 'reqd_work_group_size' is already applied with different arguments}} // expected-note@#f7prev {{previous attribute is here}} + // expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} + // expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} f7<2, 2, 2>(); // expected-note {{in instantiation}} } // Tests for 'reqd_work_group_size' attribute duplication. -[[sycl::reqd_work_group_size(8)]] // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(8)]] // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[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)]] // expected-note {{previous attribute is here}} +[[sycl::reqd_work_group_size(32, 32, 1)]] // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(32, 32)]] void f9() {} // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} // Test that template redeclarations also get diagnosed properly. @@ -127,6 +133,8 @@ void test() { f10<64, 1, 1>(); // OK, args are the same on the redecl. // expected-error@#f10err {{attribute 'reqd_work_group_size' is already applied with different arguments}} // expected-note@#f10prev {{previous attribute is here}} + // expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} + // expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} f10<1, 1, 64>(); // expected-note {{in instantiation}} } @@ -135,7 +143,8 @@ struct TRIFuncObjBad { operator()() const; }; -[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{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}} \ +// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} void TRIFuncObjBad::operator()() const {} @@ -174,6 +183,7 @@ int main() { KernelFunctor<16, 1, 1>(); } // Test that checks template parameter support on function. +// expected-warning@+2 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} template [[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp index 6578e9e41ee2a..aa85dec244fb0 100644 --- a/clang/test/SemaSYCL/sub-group-size.cpp +++ b/clang/test/SemaSYCL/sub-group-size.cpp @@ -123,9 +123,11 @@ void calls_kernel_4() { sycl::kernel_single_task([]() { // #Kernel4 // integer-error@#AttrFunc2{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // integer-note@#Kernel4{{kernel declared here}} + // expected-warning@#AttrFunc2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}} AttrFunc2(); // integer-error@#AttrExternalDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // integer-note@#Kernel4{{kernel declared here}} + // expected-warning@#AttrExternalDefined2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}} AttrExternalDefined2(); // integer-error@#AttrExternalNotDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // integer-note@#Kernel4{{kernel declared here}} diff --git a/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp b/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp new file mode 100644 index 0000000000000..c4ff35d6e02db --- /dev/null +++ b/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp @@ -0,0 +1,78 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s + +// The test check issuing diagnostics for attributes that can not be applied to a non SYCL kernel function + +#include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp + +[[sycl::reqd_work_group_size(16)]] void f1(){ // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} +} + +[[intel::reqd_sub_group_size(12)]] void f3(){ // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} +} + +[[sycl::reqd_work_group_size(16)]] void f4(){ // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} +} + +[[sycl::work_group_size_hint(8, 8, 8)]] void f5(){}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}} + +SYCL_EXTERNAL [[sycl::reqd_work_group_size(16)]] void f6() {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} + +class Functor16 { +public: + [[sycl::reqd_work_group_size(16)]] void operator()() const; +}; +void Functor16::operator()() const { +} + +class Functor16x16 { +public: + [[sycl::reqd_work_group_size(16,16)]] void operator()() const; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} +}; +void Functor16x16::operator()() const { +} + + +class Functor16x16x16 { +public: + [[sycl::reqd_work_group_size(16,16,16)]] void operator()() const{ + + } +}; + +class FunctorSubGroupSize4 { +public: + [[intel::reqd_sub_group_size(4)]] void operator()() const{} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} +}; + +class Functor8 { +public: + void operator()() const; +}; + +[[sycl::reqd_work_group_size(8)]] void Functor8::operator()() const {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} + +int main() { + sycl::queue q; + Functor16x16 f16x16; + FunctorSubGroupSize4 fs4; + Functor8 f8; + + q.submit([&](sycl::handler& h) { + h.single_task( + []()[[sycl::reqd_work_group_size(16)]]{ // OK attribute reqd_work_group_size applied to kernel + f1(); + } + ); + }); + + + q.submit([&](sycl::handler &h) { + Functor16 f16; + Functor16x16x16 f16x16x16; + h.single_task(f16); // OK attribute reqd_work_group_size applied to kernel + h.single_task(f16x16x16); // OK attribute reqd_work_group_size applied to kernel + }); + + + return 0; +} 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 f490b6eeadbca..380a3e2425cbc 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 @@ -58,6 +58,7 @@ int main() { // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} // Test that checks template parameter support on function. +// expected-warning@+2 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} template [[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} @@ -86,6 +87,7 @@ int check() { // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} // No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} [[sycl::reqd_work_group_size(4, 4, 4)]] [[sycl::reqd_work_group_size(4, 4, 4)]] void func4() {} // CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()' // CHECK: SYCLReqdWorkGroupSizeAttr