diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9150ece385879..990561be2e6ef 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3167,7 +3167,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // // code // } // -// [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const +// [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const // { // // code // } diff --git a/clang/lib/Sema/SemaSYCLDeclAttr.cpp b/clang/lib/Sema/SemaSYCLDeclAttr.cpp index 0e23934b4597f..db0a2fdce4aad 100644 --- a/clang/lib/Sema/SemaSYCLDeclAttr.cpp +++ b/clang/lib/Sema/SemaSYCLDeclAttr.cpp @@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, return; } + // Additionally, diagnose deprecated [[intel::reqd_sub_group_size]] spelling + if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize && A.getScopeName() && + A.getScopeName()->isStr("intel")) { + diagnoseDeprecatedAttribute(A, "sycl", "reqd_sub_group_size"); + return; + } + // Diagnose SYCL 2020 spellings in later SYCL modes. if (getLangOpts().getSYCLVersion() >= LangOptions::SYCL_2020) { // All attributes in the cl vendor namespace are deprecated in favor of a diff --git a/clang/test/CodeGenSYCL/kernel-op-calls.cpp b/clang/test/CodeGenSYCL/kernel-op-calls.cpp index a8f17d8235a50..0c3a53586b46c 100644 --- a/clang/test/CodeGenSYCL/kernel-op-calls.cpp +++ b/clang/test/CodeGenSYCL/kernel-op-calls.cpp @@ -11,7 +11,7 @@ class Functor1 { public: Functor1(){} - [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {} + [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {} [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {} diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index 4396313eac485..cdbc5158535cc 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -7,17 +7,17 @@ queue q; class Functor16 { public: - [[intel::reqd_sub_group_size(16)]] void operator()() const {} + [[sycl::reqd_sub_group_size(16)]] void operator()() const {} }; template class Functor2 { public: - [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {} + [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {} }; template -[[intel::reqd_sub_group_size(N)]] void func() {} +[[sycl::reqd_sub_group_size(N)]] void func() {} int main() { q.submit([&](handler &h) { @@ -25,7 +25,7 @@ int main() { h.single_task(f16); h.single_task( - []() [[intel::reqd_sub_group_size(4)]]{}); + []() [[sycl::reqd_sub_group_size(4)]]{}); Functor2<2> f2; h.single_task(f2); diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index 45f258a01b67e..cde7f32ebe068 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -7,12 +7,12 @@ queue q; class Functor { public: - [[intel::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} + [[sycl::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} }; class Functor1 { public: - [[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {} + [[sycl::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {} }; template diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp index c4da05dfbf234..c38b40e183edc 100644 --- a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp +++ b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp @@ -6,7 +6,7 @@ template class Fobj { public: Fobj() {} void operator()() const { - auto L0 = []() [[intel::reqd_sub_group_size(4)]]{}; + auto L0 = []() [[sycl::reqd_sub_group_size(4)]]{}; L0(); } }; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp index 8386663bbd37a..b21ca48807622 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp @@ -9,14 +9,14 @@ queue q; class Functor16 { public: - [[intel::reqd_sub_group_size(16)]] void operator()() const {} + [[sycl::reqd_sub_group_size(16)]] void operator()() const {} }; // Test that checks template parameter support on member function of class template. template class KernelFunctor { public: - [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {} + [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {} }; // Test that checks template parameter support on function. @@ -35,7 +35,7 @@ class KernelFunctor { // CHECK-NEXT: NonTypeTemplateParmDecl // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12 template -[[intel::reqd_sub_group_size(N)]] void func() {} +[[sycl::reqd_sub_group_size(N)]] void func() {} int main() { q.submit([&](handler &h) { @@ -52,14 +52,14 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 2 // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - h.single_task([]() [[intel::reqd_sub_group_size(2)]] {}); + h.single_task([]() [[sycl::reqd_sub_group_size(2)]] {}); // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 6 // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}} - h.single_task([]() [[intel::reqd_sub_group_size(6)]] {}); + h.single_task([]() [[sycl::reqd_sub_group_size(6)]] {}); // CHECK: FunctionDecl {{.*}}kernel_name_6 // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size @@ -79,8 +79,8 @@ int main() { // CHECK-NEXT: value: Int 8 // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} // CHECK-NOT: IntelReqdSubGroupSizeAttr - []() [[intel::reqd_sub_group_size(8), - intel::reqd_sub_group_size(8)]] {}); + []() [[sycl::reqd_sub_group_size(8), + sycl::reqd_sub_group_size(8)]] {}); }); func<12>(); return 0; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp index ab11ef9a19ce9..0acbda529e3be 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp @@ -1,9 +1,9 @@ // RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s // expected-no-diagnostics -[[intel::reqd_sub_group_size(8)]] void fun() {} +[[sycl::reqd_sub_group_size(8)]] void fun() {} class Functor { public: - [[intel::reqd_sub_group_size(16)]] void operator()() {} + [[sycl::reqd_sub_group_size(16)]] void operator()() {} }; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp index e86b68db49cfa..b44892bf44176 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s -// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel. +// The test checks functionality of [[sycl::reqd_sub_group_size()]] attribute on SYCL kernel and [[intel::reqd_sub_group_size()]] is deprecated. #include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp @@ -32,44 +32,51 @@ int main() { }); return 0; } -[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); -[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} + +[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); +[[sycl::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() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} +[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); +[[sycl::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() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} +[[sycl::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@+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() { +[[sycl::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() { sg_size2(); } // Test that checks support and functionality of reqd_sub_group_size attribute support on function. // Tests for incorrect argument values for Intel reqd_sub_group_size attribute. -[[intel::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}} -[[intel::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}} -[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}} -[[intel::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} -[[intel::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}} +[[sycl::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}} +[[sycl::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}} +[[sycl::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}} +[[sycl::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +[[sycl::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}} // Diagnostic is emitted because the arguments mismatch. -[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}} -[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} +[[sycl::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}} +[[sycl::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} [[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} // Make sure there's at least one argument passed. [[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}} // No diagnostic is emitted because the arguments match. +[[sycl::reqd_sub_group_size(12)]] void same(); +[[sycl::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} + +// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}} +// expected-warning@+1{{attribute 'intel::reqd_sub_group_size' is deprecated}} [[intel::reqd_sub_group_size(12)]] void same(); -[[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 @@ -80,7 +87,7 @@ template // expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} -[[intel::reqd_sub_group_size(Ty{})]] void func() {} +[[sycl::reqd_sub_group_size(Ty{})]] void func() {} struct S {}; void test() { @@ -97,18 +104,18 @@ void test() { int foo1(); // expected-error@+2{{expression is not an integral constant expression}} // expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}} -[[intel::reqd_sub_group_size(foo1() + 12)]] void func1(); +[[sycl::reqd_sub_group_size(foo1() + 12)]] void func1(); // Test that checks expression is a constant expression. constexpr int bar1() { return 0; } -[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK +[[sycl::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK // Test that checks template parameter support on member function of class template. template class KernelFunctor { public: // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} - [[intel::reqd_sub_group_size(SIZE)]] void operator()() {} + [[sycl::reqd_sub_group_size(SIZE)]] void operator()() {} }; int check() { @@ -121,14 +128,14 @@ int check() { template // 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() {} +[[sycl::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}} +[[sycl::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}} template -[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} +[[sycl::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} int check1() { // no error expected diff --git a/clang/test/SemaSYCL/sub-group-size.cpp b/clang/test/SemaSYCL/sub-group-size.cpp index aa85dec244fb0..65f821c1026d8 100644 --- a/clang/test/SemaSYCL/sub-group-size.cpp +++ b/clang/test/SemaSYCL/sub-group-size.cpp @@ -3,29 +3,29 @@ // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s // Validate the semantic analysis checks for the interaction betwen the -// named_sub_group_size and sub_group_size attributes. These are not able to be +// named_sub_group_size and reqd_sub_group_size attributes. These are not able to be // combined, and require that they only be applied to non-sycl-kernel/ // non-sycl-device functions if they match the kernel they are being called // from. #include "Inputs/sycl.hpp" -// expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}} +// expected-error@+2 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}} // expected-note@+1 {{conflicting attribute is here}} -[[intel::sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1(); -// expected-error@+2 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}} +[[sycl::reqd_sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1(); +// expected-error@+2 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}} // expected-note@+1 {{conflicting attribute is here}} -[[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2(); +[[intel::named_sub_group_size(primary)]] [[sycl::reqd_sub_group_size(1)]] void f2(); // expected-note@+1 {{conflicting attribute is here}} -[[intel::sub_group_size(1)]] void f3(); -// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}} +[[sycl::reqd_sub_group_size(1)]] void f3(); +// expected-error@+1 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}} [[intel::named_sub_group_size(primary)]] void f3(); // expected-note@+1 {{conflicting attribute is here}} [[intel::named_sub_group_size(primary)]] void f4(); -// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}} -[[intel::sub_group_size(1)]] void f4(); +// expected-error@+1 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}} +[[sycl::reqd_sub_group_size(1)]] void f4(); // expected-note@+1 {{previous attribute is here}} [[intel::named_sub_group_size(automatic)]] void f5(); @@ -115,8 +115,14 @@ void calls_kernel_3() { }); } +// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}} +// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}} [[intel::sub_group_size(10)]] void AttrFunc2() {} // #AttrFunc2 +// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}} +// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}} [[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {} // #AttrExternalDefined2 +// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}} +// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}} [[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalNotDefined2(); // #AttrExternalNotDefined2 void calls_kernel_4() { @@ -153,7 +159,7 @@ void calls_kernel_5() { // Don't diag with the old sub-group-size. void calls_kernel_6() { - sycl::kernel_single_task([]() [[intel::reqd_sub_group_size(10)]] { // #Kernel6 + sycl::kernel_single_task([]() [[sycl::reqd_sub_group_size(10)]] { // #Kernel6 NoAttrExternalNotDefined(); }); } diff --git a/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp b/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp index c4ff35d6e02db..8115df801552e 100644 --- a/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp +++ b/clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp @@ -7,7 +7,7 @@ [[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_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}} @@ -41,7 +41,7 @@ class Functor16x16x16 { 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}} + [[sycl::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 { diff --git a/sycl/test-e2e/AOT/reqd-sg-size.cpp b/sycl/test-e2e/AOT/reqd-sg-size.cpp index df9eb918a4eac..fa66a64872420 100644 --- a/sycl/test-e2e/AOT/reqd-sg-size.cpp +++ b/sycl/test-e2e/AOT/reqd-sg-size.cpp @@ -43,7 +43,7 @@ template struct SubgroupDispatcher { accessor acc{buf, cgh}; cgh.parallel_for>( nd_range<1>(1, 1), - [=](auto item) [[intel::reqd_sub_group_size(size)]] { + [=](auto item) [[sycl::reqd_sub_group_size(size)]] { acc[0] = item.get_sub_group().get_max_local_range()[0]; }); }); diff --git a/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp b/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp index f734ae4e696bb..3f2a8818eaa40 100644 --- a/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp +++ b/sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp @@ -32,7 +32,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl; diff --git a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp index 50addd2898e6b..3adfc92ccf256 100644 --- a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp +++ b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp @@ -78,7 +78,7 @@ void RunSortOverGroupArray(sycl::queue &Q, const std::vector &DataToSort, CGH); CGH.parallel_for( - NDRange, [=](sycl::nd_item id) [[intel::reqd_sub_group_size( + NDRange, [=](sycl::nd_item id) [[sycl::reqd_sub_group_size( ReqSubGroupSize)]] { const size_t GlobalLinearID = id.get_global_linear_id(); using RadixSorterT = oneapi_exp::radix_sorters::group_sorter< diff --git a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp index 985a9fea71133..a5ca0006b79ed 100644 --- a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp +++ b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp @@ -115,7 +115,7 @@ void RunJointSort(sycl::queue &Q, const std::vector &DataToSort, CGH.parallel_for, UseGroupWrapper, T, Compare>>( - NDRange, [=](sycl::nd_item ID) [[intel::reqd_sub_group_size( + NDRange, [=](sycl::nd_item ID) [[sycl::reqd_sub_group_size( ReqSubGroupSize)]] { auto Group = [&]() { if constexpr (UseGroup == UseGroupT::SubGroup) @@ -282,7 +282,7 @@ void RunSortOVerGroup(sycl::queue &Q, const std::vector &DataToSort, CGH.parallel_for, UseGroupWrapper, T, Compare>>( - NDRange, [=](sycl::nd_item id) [[intel::reqd_sub_group_size( + NDRange, [=](sycl::nd_item id) [[sycl::reqd_sub_group_size( ReqSubGroupSize)]] { const size_t GlobalLinearID = id.get_global_linear_id(); diff --git a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_array_input_sort.cpp b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_array_input_sort.cpp index bcb5b3cbd5aa6..0b415f878e85f 100644 --- a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_array_input_sort.cpp +++ b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_array_input_sort.cpp @@ -94,7 +94,7 @@ void RunSortKeyValueOverGroupArray(sycl::queue &Q, CGH); CGH.parallel_for(NDRange, [=](sycl::nd_item - id) [[intel::reqd_sub_group_size( + id) [[sycl::reqd_sub_group_size( ReqSubGroupSize)]] { const size_t GlobalLinearID = id.get_global_linear_id(); diff --git a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_sort.cpp b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_sort.cpp index 95b1e07445150..b2347d9b6de6e 100644 --- a/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_sort.cpp +++ b/sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_sort.cpp @@ -120,7 +120,7 @@ void RunSortKeyValueOverGroup(sycl::queue &Q, CGH); auto KeyValueSortKernel = - [=](sycl::nd_item id) [[intel::reqd_sub_group_size( + [=](sycl::nd_item id) [[sycl::reqd_sub_group_size( ReqSubGroupSize)]] { const size_t GlobalLinearID = id.get_global_linear_id(); diff --git a/sycl/test-e2e/GroupAlgorithm/load_store/odd_wg_size.cpp b/sycl/test-e2e/GroupAlgorithm/load_store/odd_wg_size.cpp index adce5e9f588bc..c778631ccc05f 100644 --- a/sycl/test-e2e/GroupAlgorithm/load_store/odd_wg_size.cpp +++ b/sycl/test-e2e/GroupAlgorithm/load_store/odd_wg_size.cpp @@ -41,7 +41,7 @@ template void test(queue &q) { cgh.parallel_for( nd_range<1>{global_size, wg_size}, - [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(SG_SIZE)]] { + [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(SG_SIZE)]] { auto gid = ndi.get_global_id(0); auto g = ndi.get_group(); auto offset = g.get_group_id(0) * g.get_local_range(0) * elems_per_wi; diff --git a/sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp b/sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp index 72b87364f0ee9..d01fbf4e6b7ae 100644 --- a/sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp +++ b/sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp @@ -37,7 +37,7 @@ template void test(queue &q) { cgh.parallel_for( nd_range<1>{wg_size, wg_size}, - [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(SG_SIZE)]] { + [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(SG_SIZE)]] { auto gid = ndi.get_global_id(0); auto sg = ndi.get_sub_group(); auto offset = diff --git a/sycl/test-e2e/InvokeSimd/Feature/SPMD_invoke_ESIMD_external.cpp b/sycl/test-e2e/InvokeSimd/Feature/SPMD_invoke_ESIMD_external.cpp index 648df2f5c5489..906e2fa46114e 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/SPMD_invoke_ESIMD_external.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/SPMD_invoke_ESIMD_external.cpp @@ -25,7 +25,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct.cpp b/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct.cpp index 9c5b92645c9d7..75bd87d74d0f3 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct.cpp @@ -32,7 +32,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct_by_pointer.cpp b/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct_by_pointer.cpp index 27709232eaac1..d137545d96295 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct_by_pointer.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct_by_pointer.cpp @@ -34,7 +34,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/popcnt.cpp b/sycl/test-e2e/InvokeSimd/Feature/popcnt.cpp index c39983bba4ebe..f99c6fd0e097b 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/popcnt.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/popcnt.cpp @@ -48,7 +48,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp b/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp index 1301ccd0f0ce1..85ab93d7d43a9 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp @@ -49,7 +49,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/scale.cpp b/sycl/test-e2e/InvokeSimd/Feature/scale.cpp index e78871769aa41..5e4af520b43c3 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/scale.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/scale.cpp @@ -36,7 +36,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/split_module/SPMD_module.cpp b/sycl/test-e2e/InvokeSimd/Feature/split_module/SPMD_module.cpp index fc4168e2568aa..a159c685b3204 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/split_module/SPMD_module.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/split_module/SPMD_module.cpp @@ -25,7 +25,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp b/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp index 123c6f2dbe770..5989e49dfe41d 100644 --- a/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp +++ b/sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp @@ -35,7 +35,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/address_space_cast.cpp b/sycl/test-e2e/InvokeSimd/Regression/address_space_cast.cpp index 3890c41037fb9..8a97eabfe7e30 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/address_space_cast.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/address_space_cast.cpp @@ -65,7 +65,7 @@ bool test() { try { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(VL)]] { + Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(VL)]] { sub_group sg = ndi.get_sub_group(); group<1> g = ndi.get_group(); uint32_t i = sg.get_group_linear_id() * VL + diff --git a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop.cpp b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop.cpp index 2cd6dfc68bf0f..6ec1dbc05bc77 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop.cpp @@ -38,7 +38,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop_naive.cpp b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop_naive.cpp index f7b761f51d213..c1408c50f73d1 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop_naive.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop_naive.cpp @@ -41,7 +41,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp index 62368d1d93acd..068cce8002b21 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp @@ -30,7 +30,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp b/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp index 05d3d6b707052..11c118dad7615 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp @@ -25,7 +25,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/dp4a.cpp b/sycl/test-e2e/InvokeSimd/Regression/dp4a.cpp index 1904c394594e7..15be55f7dde64 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/dp4a.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/dp4a.cpp @@ -28,7 +28,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(SIZE)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(SIZE)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/matrix_add.cpp b/sycl/test-e2e/InvokeSimd/Regression/matrix_add.cpp index d9def83f25dc1..ebae8778eb5e8 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/matrix_add.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/matrix_add.cpp @@ -29,7 +29,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_USM.cpp b/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_USM.cpp index f96af712ce573..9c9c1f1b7d03f 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_USM.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_USM.cpp @@ -44,7 +44,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_accessor_get_pointer.cpp b/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_accessor_get_pointer.cpp index b1187b1f90821..7644c52c610e5 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_accessor_get_pointer.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_accessor_get_pointer.cpp @@ -40,7 +40,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_basic.cpp b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_basic.cpp index 2923023ee6fa2..60a58576d2883 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_basic.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_basic.cpp @@ -26,7 +26,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl; diff --git a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp index f8ca79827dae9..5ab6ac48ea87d 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp @@ -176,7 +176,7 @@ bool test(queue q) { cgh.parallel_for>( nd_range<1>(global_range, local_range), // This test requires an explicit specification of the subgroup size - [=](nd_item<1> item) [[intel::reqd_sub_group_size(VL)]] { + [=](nd_item<1> item) [[sycl::reqd_sub_group_size(VL)]] { sycl::group<1> g = item.get_group(); sycl::sub_group sg = item.get_sub_group(); diff --git a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp index ff2461b3aa4d2..c58b419b24feb 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp @@ -164,7 +164,7 @@ int main() { cgh.parallel_for( nd_range<1>(GlobalRange, LocalRange), // This test requires an explicit specification of the subgroup size - [=](nd_item<1> item) [[intel::reqd_sub_group_size(VL)]] { + [=](nd_item<1> item) [[sycl::reqd_sub_group_size(VL)]] { sycl::group<1> g = item.get_group(); sycl::sub_group sg = item.get_sub_group(); diff --git a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp index 77cc809abab70..8ec3c6ab0cf6c 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp @@ -138,7 +138,7 @@ bool test(queue q) { cgh.parallel_for>( nd_range<1>(global_range, local_range), // This test requires an explicit specification of the subgroup size - [=](nd_item<1> item) [[intel::reqd_sub_group_size(VL)]] { + [=](nd_item<1> item) [[sycl::reqd_sub_group_size(VL)]] { sycl::group<1> g = item.get_group(); sycl::sub_group sg = item.get_sub_group(); diff --git a/sycl/test-e2e/InvokeSimd/Regression/slm_gather_scatter.cpp b/sycl/test-e2e/InvokeSimd/Regression/slm_gather_scatter.cpp index f4cc97d3c34b6..e2a97031469bf 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/slm_gather_scatter.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/slm_gather_scatter.cpp @@ -34,7 +34,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl; diff --git a/sycl/test-e2e/InvokeSimd/Regression/slm_load_store.cpp b/sycl/test-e2e/InvokeSimd/Regression/slm_load_store.cpp index c5c92ff6ac4fc..fb981f36a3d16 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/slm_load_store.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/slm_load_store.cpp @@ -34,7 +34,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl; diff --git a/sycl/test-e2e/InvokeSimd/Regression/tiled_matrix_multiplication.cpp b/sycl/test-e2e/InvokeSimd/Regression/tiled_matrix_multiplication.cpp index d7e67d28c4c1b..7b0ca06812d5d 100644 --- a/sycl/test-e2e/InvokeSimd/Regression/tiled_matrix_multiplication.cpp +++ b/sycl/test-e2e/InvokeSimd/Regression/tiled_matrix_multiplication.cpp @@ -46,7 +46,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp b/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp index 75b8eef13f926..2e24ad430ffa1 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/ESIMD_to_unmarked_function.cpp @@ -43,7 +43,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl; diff --git a/sycl/test-e2e/InvokeSimd/Spec/clang_run_error/reference_argument.cpp b/sycl/test-e2e/InvokeSimd/Spec/clang_run_error/reference_argument.cpp index 8f7871bc1e26e..48d40b5067ad5 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/clang_run_error/reference_argument.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/clang_run_error/reference_argument.cpp @@ -67,7 +67,7 @@ int main(void) { try { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(VL)]] { + Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(VL)]] { sub_group sg = ndi.get_sub_group(); group<1> g = ndi.get_group(); uint32_t i = sg.get_group_linear_id() * VL + diff --git a/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp b/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp index 91bf5efe0386e..71be83ded52d1 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/function_overloads.cpp @@ -37,7 +37,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_multiple_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_multiple_ESIMD.cpp index 1a1c35c726367..4348b51f77eae 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_multiple_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_multiple_ESIMD.cpp @@ -33,7 +33,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_single_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_single_ESIMD.cpp index 0ab085c7dd121..3a760ab1824d5 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_single_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/multiple_SPMD_to_single_ESIMD.cpp @@ -33,7 +33,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp index da0c18e7c1996..762f98480bcd8 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/nested_ESIMD_to_ESIMD.cpp @@ -30,7 +30,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp b/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp index ece54a2dcf430..b0c062a2d72fe 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/nested_SPMD_to_ESIMD.cpp @@ -29,7 +29,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/simd_mask.cpp b/sycl/test-e2e/InvokeSimd/Spec/simd_mask.cpp index 1337efb63c3d4..eeabbe0f52774 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/simd_mask.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/simd_mask.cpp @@ -30,7 +30,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/simd_size/Inputs/common.hpp b/sycl/test-e2e/InvokeSimd/Spec/simd_size/Inputs/common.hpp index c2c7af2f678df..5f9ecb13d2d87 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/simd_size/Inputs/common.hpp +++ b/sycl/test-e2e/InvokeSimd/Spec/simd_size/Inputs/common.hpp @@ -69,7 +69,7 @@ template bool test(QueueTY q) { try { auto e = q.submit([&](handler &cgh) { cgh.parallel_for>( - Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(VL)]] { + Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(VL)]] { sub_group sg = ndi.get_sub_group(); group<1> g = ndi.get_group(); uint32_t i = sg.get_group_linear_id() * VL + diff --git a/sycl/test-e2e/InvokeSimd/Spec/tuple.cpp b/sycl/test-e2e/InvokeSimd/Spec/tuple.cpp index 5ff856e68e0c8..fd835a147c6f1 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/tuple.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/tuple.cpp @@ -39,7 +39,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/tuple_return.cpp b/sycl/test-e2e/InvokeSimd/Spec/tuple_return.cpp index f1d1d2c462bfe..14be37cf1c435 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/tuple_return.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/tuple_return.cpp @@ -40,7 +40,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/tuple_vadd.cpp b/sycl/test-e2e/InvokeSimd/Spec/tuple_vadd.cpp index e78b05d4c121b..5663aec9ff23c 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/tuple_vadd.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/tuple_vadd.cpp @@ -46,7 +46,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp b/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp index 771818933572f..bd2458ff278b8 100644 --- a/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp +++ b/sycl/test-e2e/InvokeSimd/Spec/uniform_retval.cpp @@ -62,7 +62,7 @@ #ifdef IMPL_SUBGROUP #define SUBGROUP_ATTR #else -#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]] +#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]] #endif using namespace sycl::ext::oneapi::experimental; diff --git a/sycl/test-e2e/InvokeSimd/invoke_simd_conv.cpp b/sycl/test-e2e/InvokeSimd/invoke_simd_conv.cpp index 023bc373611c6..00d59b8fe9f8f 100644 --- a/sycl/test-e2e/InvokeSimd/invoke_simd_conv.cpp +++ b/sycl/test-e2e/InvokeSimd/invoke_simd_conv.cpp @@ -88,7 +88,7 @@ template bool test(queue q) { try { auto e = q.submit([&](handler &cgh) { cgh.parallel_for>( - Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(VL)]] { + Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(VL)]] { sub_group sg = ndi.get_sub_group(); SpmdT val = (SpmdT)sg.get_group_linear_id(); // 0 .. GroupSize-1 SimdElemT res = 0; diff --git a/sycl/test-e2e/InvokeSimd/invoke_simd_smoke.cpp b/sycl/test-e2e/InvokeSimd/invoke_simd_smoke.cpp index e202e33a86df2..1db1a7c738be8 100644 --- a/sycl/test-e2e/InvokeSimd/invoke_simd_smoke.cpp +++ b/sycl/test-e2e/InvokeSimd/invoke_simd_smoke.cpp @@ -102,7 +102,7 @@ template bool test() { try { auto e = q.submit([&](handler &cgh) { - cgh.parallel_for(Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size( + cgh.parallel_for(Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size( VL)]] { sub_group sg = ndi.get_sub_group(); group<1> g = ndi.get_group(); diff --git a/sycl/test-e2e/Matrix/element_wise_abc_impl.hpp b/sycl/test-e2e/Matrix/element_wise_abc_impl.hpp index 5caf6d3e0a3e7..59a3b9656a606 100644 --- a/sycl/test-e2e/Matrix/element_wise_abc_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_abc_impl.hpp @@ -36,7 +36,7 @@ void matrix_elem_wise_ops(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp index d2bdcbcb2d04a..509bdf010b999 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_half_impl.hpp @@ -36,7 +36,7 @@ void matrix_verify_op(big_matrix &A, const R ref, OP op) { cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp index 765ab5c54f53f..ac6dc118ce8b2 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_impl.hpp @@ -40,7 +40,7 @@ void verify_op_ab(const T l, const T r, const float ref, OP op) { {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); @@ -80,7 +80,7 @@ void verify_op_c(const T l, const T r, const float ref, OP op) { {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp index 8a2f1f495e41d..b4b43d789eff9 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_int8_impl.hpp @@ -37,7 +37,7 @@ void matrix_verify_op(big_matrix &A, const R ref, OP op) { cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp index 7336bb8467fa5..68069f297fa64 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_int8_packed_impl.hpp @@ -39,7 +39,7 @@ void matrix_verify_op(big_matrix &B, const TResult ref, OP op) { cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp index 2d7b3a36d8296..6e64c8a0c2bc6 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_ops_tf32_impl.hpp @@ -35,7 +35,7 @@ void matrix_verify_op(big_matrix &A, const float ref, OP op) { cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp b/sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp index 5228a154e9f6f..6a62c06d1dc5a 100644 --- a/sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_all_sizes_impl.hpp @@ -46,7 +46,7 @@ void matrix_verify_add(const T1 val1, const T1 val2, const T1 result) { cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp b/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp index 5c89e16c35f5b..5a2cf801eb3ec 100644 --- a/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp +++ b/sycl/test-e2e/Matrix/element_wise_ops_impl.hpp @@ -28,7 +28,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp b/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp index eb0a22fa30566..8dfe68e24bef7 100644 --- a/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp +++ b/sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp @@ -50,7 +50,7 @@ void matrix_multiply(big_matrix &C, cgh.parallel_for( nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}), [accA, accB, accC, M, N, K](nd_item<2> spmd_item) - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/get_coord_float_matC_impl.hpp b/sycl/test-e2e/Matrix/get_coord_float_matC_impl.hpp index bb0b6336a0ec5..a6c7c5646a548 100644 --- a/sycl/test-e2e/Matrix/get_coord_float_matC_impl.hpp +++ b/sycl/test-e2e/Matrix/get_coord_float_matC_impl.hpp @@ -39,7 +39,7 @@ void matrix_sum_rows(big_matrix &C, T *sum_rows) { {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/get_coord_int8_matA_impl.hpp b/sycl/test-e2e/Matrix/get_coord_int8_matA_impl.hpp index f3d89cc717ef4..0bf9281ab9f45 100644 --- a/sycl/test-e2e/Matrix/get_coord_int8_matA_impl.hpp +++ b/sycl/test-e2e/Matrix/get_coord_int8_matA_impl.hpp @@ -90,7 +90,7 @@ void matrix_sum_rows(big_matrix &A) { cgh.parallel_for>( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/get_coord_int8_matB_impl.hpp b/sycl/test-e2e/Matrix/get_coord_int8_matB_impl.hpp index 8b63dadc029b3..08cb616cc6cc4 100644 --- a/sycl/test-e2e/Matrix/get_coord_int8_matB_impl.hpp +++ b/sycl/test-e2e/Matrix/get_coord_int8_matB_impl.hpp @@ -117,7 +117,7 @@ void matrix_sum_cols(big_matrix &B, cgh.parallel_for>( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp index b3001a68bb227..816a381617a7b 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_all_sizes_impl.hpp @@ -31,7 +31,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp index 5463ea040d1eb..6da2e74d5a359 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_annotated_ptr_impl.hpp @@ -24,7 +24,7 @@ void matrix_multiply(T1 *C, T2 *A, T2 *B, queue &q) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp index 3fc96f77e020a..b1237b0894fa7 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_apply_bf16_impl.hpp @@ -29,7 +29,7 @@ void matrix_verify_add(big_matrix &A, const TResult ref, cgh.parallel_for( r, [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp index 9751571bcbcf5..a88b0ca55416e 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_apply_two_matrices_impl.hpp @@ -39,7 +39,7 @@ bool apply_two_matrices(Tc *C, Tc *D, Ta *A, Ta *Ar, queue q) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp index 1b17dbee96a57..22ce2b3f0e16a 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_impl.hpp @@ -75,7 +75,7 @@ double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i // loop localrange [=](nd_item<2> it) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif // SG_SZ { // sg::load and sg::store expect decorations to be ON diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp index b9f474b9758dc..bd4a66d2c4dc9 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_load_store_impl.hpp @@ -18,7 +18,7 @@ void joint_B_rowmajor_load_store(Tb *B, Tb *OutB, queue &q) { h.parallel_for( nd_range<1>{global, local}, [=](nd_item<1> it) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { auto pB = diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp index 0cacda21b98e2..a7a502a3c16a1 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bf16_rowmajorB_pair_load_store_impl.hpp @@ -18,7 +18,7 @@ void joint_B_rowmajor_pair_load_store(Tb *B, Tb *OutB, queue &q) { h.parallel_for( nd_range<1>{global, local}, [=](nd_item<1> it) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { auto pB = diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp index f393eaa5e8436..de2d1d89deaf5 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_array_impl.hpp @@ -31,7 +31,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // Matrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp index e3234da2cd5d9..d8f5e45474a77 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_colmajorA_colmajorB_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp index ede4e795d0d69..00e804cef2fb5 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp index 6a7182c41985d..85d33f2c83173 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_bfloat16_packedB_impl.hpp @@ -27,7 +27,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp index 373ec652cc063..65b091477cae5 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_colA_rowB_colC_impl.hpp @@ -31,7 +31,7 @@ void matrix_multiply(T1 *C, T2 *A, T2 *B, queue q) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { auto pA = diff --git a/sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp index 8ac48511c7e10..8c93114cf019a 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_down_convert_impl.hpp @@ -32,7 +32,7 @@ void matrix_copy(big_matrix &C, big_matrix &A) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp index 69ee6d4da5464..e51e7c30fa810 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_half_impl.hpp @@ -28,7 +28,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp index 1390f8225406c..b7b55fa42a929 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB_impl.hpp @@ -35,7 +35,7 @@ void matrix_multiply(big_matrix &C, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp index 12f91f083def4..187e199d9b14d 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_int8_rowmajorA_rowmajorB_impl.hpp @@ -28,7 +28,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp index 7aba5911c8386..e07d151afe7af 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_opt_kernel_feature_impl.hpp @@ -32,7 +32,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp index 000984087eed1..d8c02b2dc36fb 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_out_bounds_impl.hpp @@ -33,7 +33,7 @@ void matrix_multiply(T1 *C, T2 *A, T2 *B, queue q) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { auto pA = diff --git a/sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp index 1e665f618860f..03dfa4649d6b3 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_prefetch_impl.hpp @@ -74,7 +74,7 @@ void matrix_multiply(T *C, T1 *A, T2 *B, queue q) { nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { const auto global_idx = spmd_item.get_global_id(0); diff --git a/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp b/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp index 2dd752b0a9c78..b0a09f05f6a1a 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp +++ b/sycl/test-e2e/Matrix/joint_matrix_query_default.cpp @@ -66,7 +66,7 @@ void matrix_multiply(big_matrix &C, cgh.parallel_for( nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}), [accA, accB, accC, M, N, K](nd_item<2> spmd_item) - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp index 44ad7b5910076..87fbc1e90a386 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_rowmajorA_rowmajorB_impl.hpp @@ -27,7 +27,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp index 974a489002b47..ef08722a3f8ac 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_ss_int8_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp index 17c9d47f61c36..5234af6b812ea 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_su_int8_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp index 9fb40a78f8b30..02f3df63b1efd 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_tf32_impl.hpp @@ -37,7 +37,7 @@ void matrix_multiply(big_matrix &C, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The matrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp index 278e5da5cf441..58505b6fd4fb6 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_transposeC_impl.hpp @@ -28,7 +28,7 @@ void matrix_load_and_store(T1 *input, T1 *out_col_major, T1 *out_row_major, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { auto p_input = diff --git a/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp index baef5c195a1e6..1a82d390cea8e 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_us_int8_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp b/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp index 090b552848f0b..d86ff267e3f64 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp +++ b/sycl/test-e2e/Matrix/joint_matrix_uu_int8_impl.hpp @@ -29,7 +29,7 @@ void matrix_multiply(big_matrix &C, big_matrix &A, nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}), [=](nd_item<2> spmd_item) #ifdef SG_SZ - [[intel::reqd_sub_group_size(SG_SZ)]] + [[sycl::reqd_sub_group_size(SG_SZ)]] #endif { // The submatrix API has to be accessed by all the workitems in a diff --git a/sycl/test-e2e/SubGroupMask/Basic.cpp b/sycl/test-e2e/SubGroupMask/Basic.cpp index d82a507b18dc9..40526c7ccad7e 100644 --- a/sycl/test-e2e/SubGroupMask/Basic.cpp +++ b/sycl/test-e2e/SubGroupMask/Basic.cpp @@ -35,7 +35,7 @@ int main() { auto resacc = resbuf.get_access(cgh); cgh.parallel_for( - NdRange, [=](nd_item<1> NdItem) [[intel::reqd_sub_group_size(32)]] { + NdRange, [=](nd_item<1> NdItem) [[sycl::reqd_sub_group_size(32)]] { size_t GID = NdItem.get_global_linear_id(); auto SG = NdItem.get_sub_group(); // AAAAAAAA diff --git a/sycl/test-e2e/SubGroupMask/GroupSize.cpp b/sycl/test-e2e/SubGroupMask/GroupSize.cpp index b61fa4aa0e69a..827dd6d351274 100644 --- a/sycl/test-e2e/SubGroupMask/GroupSize.cpp +++ b/sycl/test-e2e/SubGroupMask/GroupSize.cpp @@ -37,8 +37,8 @@ template void test(queue Queue) { auto resacc = resbuf.template get_access(cgh); cgh.parallel_for>( - NdRange, [= - ](nd_item<1> NdItem) [[intel::reqd_sub_group_size(SGSize)]] { + NdRange, + [=](nd_item<1> NdItem) [[sycl::reqd_sub_group_size(SGSize)]] { auto SG = NdItem.get_sub_group(); auto LID = SG.get_local_id(); auto SGID = SG.get_group_id(); diff --git a/sycl/test-e2e/syclcompat/util/util_logical_group.cpp b/sycl/test-e2e/syclcompat/util/util_logical_group.cpp index db387075fe2c9..8bf7ce8238379 100644 --- a/sycl/test-e2e/syclcompat/util/util_logical_group.cpp +++ b/sycl/test-e2e/syclcompat/util/util_logical_group.cpp @@ -68,7 +68,7 @@ void test_logical_group() { result_device = sycl::malloc_device(4, q_ct1); q_ct1.parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, 52), sycl::range<3>(1, 1, 52)), - [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { + [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(32)]] { kernel(result_device, item_ct1); }); q_ct1.memcpy(result_host, result_device, sizeof(unsigned int) * 4).wait(); diff --git a/sycl/test-e2e/syclcompat/util/util_match_all_over_group.cpp b/sycl/test-e2e/syclcompat/util/util_match_all_over_group.cpp index 7d72c0a5b39f8..a1abe3bae1ed0 100644 --- a/sycl/test-e2e/syclcompat/util/util_match_all_over_group.cpp +++ b/sycl/test-e2e/syclcompat/util/util_match_all_over_group.cpp @@ -92,7 +92,7 @@ void test_match_all_over_group() { sycl::queue q = syclcompat::get_default_queue(); q.parallel_for( sycl::nd_range<1>(threads.size(), threads.size()), - [=](sycl::nd_item<1> item) [[intel::reqd_sub_group_size(SUBGROUP_SIZE)]] { + [=](sycl::nd_item<1> item) [[sycl::reqd_sub_group_size(SUBGROUP_SIZE)]] { for (auto id = item.get_global_linear_id(); id < DATA_SIZE; id += SUBGROUP_SIZE) d_output[id] = syclcompat::match_all_over_sub_group( diff --git a/sycl/test-e2e/syclcompat/util/util_match_any_over_group.cpp b/sycl/test-e2e/syclcompat/util/util_match_any_over_group.cpp index bee9ecc21272f..3a78768599f65 100644 --- a/sycl/test-e2e/syclcompat/util/util_match_any_over_group.cpp +++ b/sycl/test-e2e/syclcompat/util/util_match_any_over_group.cpp @@ -79,7 +79,7 @@ void test_match_any_over_group() { sycl::queue q = syclcompat::get_default_queue(); q.parallel_for( sycl::nd_range<1>(grid.size() * threads.size(), threads.size()), - [=](sycl::nd_item<1> item) [[intel::reqd_sub_group_size(SUBGROUP_SIZE)]] { + [=](sycl::nd_item<1> item) [[sycl::reqd_sub_group_size(SUBGROUP_SIZE)]] { auto id = item.get_global_linear_id(); d_output[id] = syclcompat::match_any_over_sub_group( item.get_sub_group(), member_mask, d_input[id]); diff --git a/sycl/test-e2e/syclcompat/util/util_permute_sub_group_by_xor.cpp b/sycl/test-e2e/syclcompat/util/util_permute_sub_group_by_xor.cpp index 18656e205ba25..7b877d826f18b 100644 --- a/sycl/test-e2e/syclcompat/util/util_permute_sub_group_by_xor.cpp +++ b/sycl/test-e2e/syclcompat/util/util_permute_sub_group_by_xor.cpp @@ -112,7 +112,7 @@ void test_permute_sub_group_by_xor() { q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { permute_sub_group_by_xor1(dev_data_u, item_ct1); }); @@ -139,7 +139,7 @@ void test_permute_sub_group_by_xor() { .wait(); q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { permute_sub_group_by_xor2(dev_data_u, item_ct1); }); diff --git a/sycl/test-e2e/syclcompat/util/util_select_from_sub_group.cpp b/sycl/test-e2e/syclcompat/util/util_select_from_sub_group.cpp index 592a1932b7ea2..ffad55f257430 100644 --- a/sycl/test-e2e/syclcompat/util/util_select_from_sub_group.cpp +++ b/sycl/test-e2e/syclcompat/util/util_select_from_sub_group.cpp @@ -108,7 +108,7 @@ void test_select_from_sub_group() { .wait(); q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { select_from_sub_group1(dev_data_u, item_ct1); }); @@ -134,7 +134,7 @@ void test_select_from_sub_group() { .wait(); q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { select_from_sub_group2(dev_data_u, item_ct1); }); diff --git a/sycl/test-e2e/syclcompat/util/util_shift_sub_group_left.cpp b/sycl/test-e2e/syclcompat/util/util_shift_sub_group_left.cpp index 81c2bd2c9616d..0fac1ee013d06 100644 --- a/sycl/test-e2e/syclcompat/util/util_shift_sub_group_left.cpp +++ b/sycl/test-e2e/syclcompat/util/util_shift_sub_group_left.cpp @@ -107,7 +107,7 @@ void test_shift_sub_group_left() { .wait(); q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { shift_sub_group_left1(dev_data_u, item_ct1); }); dev_ct1.queues_wait_and_throw(); @@ -134,7 +134,7 @@ void test_shift_sub_group_left() { q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { shift_sub_group_left2(dev_data_u, item_ct1); }); diff --git a/sycl/test-e2e/syclcompat/util/util_shift_sub_group_right.cpp b/sycl/test-e2e/syclcompat/util/util_shift_sub_group_right.cpp index 5204586df79a1..0dbc985170f03 100644 --- a/sycl/test-e2e/syclcompat/util/util_shift_sub_group_right.cpp +++ b/sycl/test-e2e/syclcompat/util/util_shift_sub_group_right.cpp @@ -108,7 +108,7 @@ void test_shift_sub_group_right() { .wait(); q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { shift_sub_group_right1(dev_data_u, item_ct1); }); @@ -135,7 +135,7 @@ void test_shift_sub_group_right() { q_ct1->parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item_ct1) - [[intel::reqd_sub_group_size(32)]] { + [[sycl::reqd_sub_group_size(32)]] { shift_sub_group_right2(dev_data_u, item_ct1); }); diff --git a/sycl/test/check_device_code/matrix/matrix-int8-test.cpp b/sycl/test/check_device_code/matrix/matrix-int8-test.cpp index 30008cf5b99fb..ee56a0c73fe61 100644 --- a/sycl/test/check_device_code/matrix/matrix-int8-test.cpp +++ b/sycl/test/check_device_code/matrix/matrix-int8-test.cpp @@ -27,7 +27,7 @@ using namespace sycl::ext::oneapi::experimental::matrix; // int8_t B[MATRIX_K / 4][MATRIX_N * 4]; // int32_t C[MATRIX_M][MATRIX_N]; -SYCL_EXTERNAL [[intel::reqd_sub_group_size(SG_SZ)]] void +SYCL_EXTERNAL [[sycl::reqd_sub_group_size(SG_SZ)]] void matrix_multiply(size_t NUM_COLS_C, size_t NUM_COLS_A, sycl::accessor accA, sycl::accessor accB, diff --git a/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp b/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp index 34fae66a8f09a..771235690ac8d 100644 --- a/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp +++ b/sycl/test/check_device_code/matrix/matrix_load_store_as.cpp @@ -9,7 +9,7 @@ using namespace sycl; using namespace sycl::ext::oneapi::experimental::matrix; -SYCL_EXTERNAL [[intel::reqd_sub_group_size(16)]] void matrix_store_as( +SYCL_EXTERNAL [[sycl::reqd_sub_group_size(16)]] void matrix_store_as( multi_ptr pA, multi_ptr pB, multi_ptr pC, diff --git a/sycl/test/esimd/slm_init_invoke_simd.cpp b/sycl/test/esimd/slm_init_invoke_simd.cpp index bc57240644002..899e498de807c 100644 --- a/sycl/test/esimd/slm_init_invoke_simd.cpp +++ b/sycl/test/esimd/slm_init_invoke_simd.cpp @@ -25,7 +25,7 @@ SYCL_EXTERNAL int main() { queue Q; nd_range<1> NDR{range<1>{2}, range<1>{2}}; - Q.parallel_for(NDR, [=](nd_item<1> NDI) [[intel::reqd_sub_group_size(16)]] { + Q.parallel_for(NDR, [=](nd_item<1> NDI) [[sycl::reqd_sub_group_size(16)]] { sub_group sg = NDI.get_sub_group(); invoke_simd(sg, SIMD_CALLEE_VOID); }).wait(); diff --git a/sycl/test/extensions/inline_asm.cpp b/sycl/test/extensions/inline_asm.cpp index 613484cc0474c..0ce33c8d7dc7a 100644 --- a/sycl/test/extensions/inline_asm.cpp +++ b/sycl/test/extensions/inline_asm.cpp @@ -38,8 +38,8 @@ int main() { auto B = BufB.get_access(cgh); auto C = BufC.get_access(cgh); cgh.parallel_for( - sycl::range<1>{DEFAULT_PROBLEM_SIZE}, [= - ](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] { + sycl::range<1>{DEFAULT_PROBLEM_SIZE}, + [=](sycl::id<1> wiID) [[sycl::reqd_sub_group_size(8)]] { #if defined(__SYCL_DEVICE_ONLY__) asm volatile( ".decl P1 v_type=P num_elts=8\n" @@ -58,10 +58,10 @@ int main() { : "+rw"(C[wiID]) : "rw"(A[wiID]), "rw"(B[wiID])); #else - C[wiID] = 0; - for (int i = 0; i < A[wiID]; ++i) { - C[wiID] = C[wiID] + B[wiID]; - } + C[wiID] = 0; + for (int i = 0; i < A[wiID]; ++i) { + C[wiID] = C[wiID] + B[wiID]; + } #endif }); }); diff --git a/sycl/test/invoke_simd/invoke_simd.cpp b/sycl/test/invoke_simd/invoke_simd.cpp index 33017aa2d9337..a7d3e11983ebb 100644 --- a/sycl/test/invoke_simd/invoke_simd.cpp +++ b/sycl/test/invoke_simd/invoke_simd.cpp @@ -110,7 +110,7 @@ int main(void) { try { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - Range, [=](nd_item<1> ndi) [[intel::reqd_sub_group_size(VL)]] { + Range, [=](nd_item<1> ndi) [[sycl::reqd_sub_group_size(VL)]] { sub_group sg = ndi.get_sub_group(); group<1> g = ndi.get_group(); uint32_t i = diff --git a/sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp b/sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp index 43de05867ec03..15bc6f3466ef1 100644 --- a/sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp +++ b/sycl/test/invoke_simd/invoke_simd_address_space_inferral.cpp @@ -50,7 +50,7 @@ bool test() { try { auto e = q.submit([&](handler &cgh) { local_accessor LocalAcc(Size, cgh); - cgh.parallel_for(Range, [=](nd_item<1> item) [[intel::reqd_sub_group_size( + cgh.parallel_for(Range, [=](nd_item<1> item) [[sycl::reqd_sub_group_size( VL)]] { sycl::group<1> g = item.get_group(); sycl::sub_group sg = item.get_sub_group(); diff --git a/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp b/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp index 02cfbc0f8b904..baefb8d602487 100644 --- a/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp +++ b/sycl/test/matrix/matrix-bfloat16-test-coord-basicB.cpp @@ -147,7 +147,7 @@ void matrix_sum_cols(queue q, big_matrix &B, nd_range<2> &r) { auto os = sycl::stream(100000, 6144, cgh); cgh.parallel_for( - r, [=](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] { + r, [=](nd_item<2> spmd_item) [[sycl::reqd_sub_group_size(SG_SZ)]] { const auto global_idx = spmd_item.get_global_id(0); const auto global_idy = spmd_item.get_global_id(1); const auto sg_startx = global_idx - spmd_item.get_local_id(0); diff --git a/sycl/test/matrix/matrix-tf32-test.cpp b/sycl/test/matrix/matrix-tf32-test.cpp index fb2e7ab66492d..5e8413346c3d0 100644 --- a/sycl/test/matrix/matrix-tf32-test.cpp +++ b/sycl/test/matrix/matrix-tf32-test.cpp @@ -47,7 +47,7 @@ void matrix_multiply(big_matrix &C, cgh.parallel_for( nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}), - [=](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] + [=](nd_item<2> spmd_item) [[sycl::reqd_sub_group_size(SG_SZ)]] { // The matrix API has to be accessed by all the workitems in a diff --git a/sycl/test/warnings/warnings.cpp b/sycl/test/warnings/warnings.cpp index 2fb70f50ddb23..86bc5789e44c2 100644 --- a/sycl/test/warnings/warnings.cpp +++ b/sycl/test/warnings/warnings.cpp @@ -4,6 +4,7 @@ #include using namespace sycl; + int main() { vec newVec; queue myQueue; @@ -39,4 +40,4 @@ template class device_image; template class device_image; template class device_image; -} +} // namespace sycl