Skip to content

Deprecate intel reqd sub group size #15798

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 11 commits into from
Oct 29, 2024
Merged

Deprecate intel reqd sub group size #15798

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
1d6b9f9
[SYCL] make intel::reqd_sub_group_size deprecat`ed
dklochkov-emb Oct 21, 2024
695c112
[SYCL] fix tests due to intel::reqd_sub_group_size deprecated
dklochkov-emb Oct 22, 2024
f8bb79c
[SYCL] revert back intel deprecated attribute
dklochkov-emb Oct 22, 2024
5ff233f
[SYCL] fix cuda test
dklochkov-emb Oct 22, 2024
916fcba
[SYCL] update tests after deprecating old attribute
dklochkov-emb Oct 22, 2024
0db7d55
[SYCL] do not modify tests of clang subgroup
dklochkov-emb Oct 23, 2024
a86f9da
[SYCL] handle group size if sycl or not
dklochkov-emb Oct 23, 2024
aa97bd5
[SYCL] fix null pointer crash
dklochkov-emb Oct 25, 2024
7957b00
fix formatting
dklochkov-emb Oct 25, 2024
1e3a3ef
[SYCL] fix tests
dklochkov-emb Oct 25, 2024
d4d2d26
[SYCL] remove unnecessary comment
dklochkov-emb Oct 28, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
// }
Expand Down
7 changes: 7 additions & 0 deletions clang/lib/Sema/SemaSYCLDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
return;
}

// Additionally, diagnose deprecated [[sycl::reqd_sub_group_size]] spelling
if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize &&
A.getAttrName()->isStr("reqd_sub_group_size") && 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
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-op-calls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {}

Expand Down
8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,25 +7,25 @@ queue q;

class Functor16 {
public:
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
[[sycl::reqd_sub_group_size(16)]] void operator()() const {}
};

template <int SIZE>
class Functor2 {
public:
[[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
[[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
};

template <int N>
[[intel::reqd_sub_group_size(N)]] void func() {}
[[sycl::reqd_sub_group_size(N)]] void func() {}

int main() {
q.submit([&](handler &h) {
Functor16 f16;
h.single_task<class kernel_name1>(f16);

h.single_task<class kernel_name3>(
[]() [[intel::reqd_sub_group_size(4)]]{});
[]() [[sycl::reqd_sub_group_size(4)]]{});

Functor2<2> f2;
h.single_task<class kernel_name4>(f2);
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <int SIZE, int SIZE1, int SIZE2>
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ template <typename T> class Fobj {
public:
Fobj() {}
void operator()() const {
auto L0 = []() [[intel::reqd_sub_group_size(4)]]{};
auto L0 = []() [[sycl::reqd_sub_group_size(4)]]{};
L0();
}
};
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s
// RUN: %clang_cc1 -fsycl-is-device -Wno-error=deprecated-declarations -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s

// The test checks AST of [[intel::reqd_sub_group_size()]] attribute.

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsycl-is-host -Wno-error=deprecated-declarations -fsyntax-only -verify %s
// expected-no-diagnostics

[[intel::reqd_sub_group_size(8)]] void fun() {}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/reqd-sub-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s
// RUN: %clang_cc1 -fsycl-is-device -Wno-error=deprecated-declarations -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s

// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel.

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/sub-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary,integer %s
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,integer %s
// 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
// RUN: %clang_cc1 -internal-isystem %S/Inputs -Wno-error=deprecated-declarations -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
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s
// RUN: %clang_cc1 -fsycl-is-device -Wno-error=deprecated-declarations -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

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/AOT/reqd-sg-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ template <size_t... Ns> struct SubgroupDispatcher {
accessor acc{buf, cgh};
cgh.parallel_for<kernel_name<size>>(
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];
});
});
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,7 +78,7 @@ void RunSortOverGroupArray(sycl::queue &Q, const std::vector<T> &DataToSort,
CGH);

CGH.parallel_for(
NDRange, [=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
NDRange, [=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
ReqSubGroupSize)]] {
const size_t GlobalLinearID = id.get_global_linear_id();
using RadixSorterT = oneapi_exp::radix_sorters::group_sorter<
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ void RunJointSort(sycl::queue &Q, const std::vector<T> &DataToSort,

CGH.parallel_for<KernelNameJoint<IntWrapper<Dims>,
UseGroupWrapper<UseGroup>, T, Compare>>(
NDRange, [=](sycl::nd_item<Dims> ID) [[intel::reqd_sub_group_size(
NDRange, [=](sycl::nd_item<Dims> ID) [[sycl::reqd_sub_group_size(
ReqSubGroupSize)]] {
auto Group = [&]() {
if constexpr (UseGroup == UseGroupT::SubGroup)
Expand Down Expand Up @@ -282,7 +282,7 @@ void RunSortOVerGroup(sycl::queue &Q, const std::vector<T> &DataToSort,

CGH.parallel_for<KernelNameOverGroup<
IntWrapper<Dims>, UseGroupWrapper<UseGroup>, T, Compare>>(
NDRange, [=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
NDRange, [=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
ReqSubGroupSize)]] {
const size_t GlobalLinearID = id.get_global_linear_id();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_array_input_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,7 @@ void RunSortKeyValueOverGroupArray(sycl::queue &Q,
CGH);

CGH.parallel_for(NDRange, [=](sycl::nd_item<Dims>
id) [[intel::reqd_sub_group_size(
id) [[sycl::reqd_sub_group_size(
ReqSubGroupSize)]] {
const size_t GlobalLinearID = id.get_global_linear_id();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/key_value_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,7 +120,7 @@ void RunSortKeyValueOverGroup(sycl::queue &Q,
CGH);

auto KeyValueSortKernel =
[=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
[=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
ReqSubGroupSize)]] {
const size_t GlobalLinearID = id.get_global_linear_id();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/GroupAlgorithm/load_store/odd_wg_size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ template <int SG_SIZE, int WG_SIZE> 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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/GroupAlgorithm/load_store/partial_sg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ template <int SG_SIZE> 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 =
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/SPMD_invoke_ESIMD_external.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/invoke_simd_struct_by_pointer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/popcnt.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/scale.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/split_module/SPMD_module.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/address_space_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 +
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_loop_naive.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/dp4a.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/matrix_add.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_USM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/matrix_multiply_accessor_get_pointer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/nbarrier_basic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -176,7 +176,7 @@ bool test(queue q) {
cgh.parallel_for<KernelID<CaseNum>>(
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();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,7 +164,7 @@ int main() {
cgh.parallel_for<KernelID>(
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();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -138,7 +138,7 @@ bool test(queue q) {
cgh.parallel_for<KernelID<CaseNum>>(
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();

Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/InvokeSimd/Regression/slm_gather_scatter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
Loading
Loading