Skip to content

Commit 2172d9e

Browse files
dklochkov-embdklochkov-emb
authored
[SYCL] Deprecate [[intel::reqd_sub_group_size]] (#15798)
That particular spelling had been introduced as part of an Intel extension to support sub-groups in SYCL 1.2.1 before they became a core feature of SYCL 2020. The extension has been deprecated for a while and no one should use that legacy spelling anymore. The official SYCL 2020 spelling should be used instead (with `sycl::` namespace).
1 parent b7ef830 commit 2172d9e

File tree

115 files changed

+188
-167
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

115 files changed

+188
-167
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3167,7 +3167,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
31673167
// // code
31683168
// }
31693169
//
3170-
// [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
3170+
// [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
31713171
// {
31723172
// // code
31733173
// }

‎clang/lib/Sema/SemaSYCLDeclAttr.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
132132
return;
133133
}
134134

135+
// Additionally, diagnose deprecated [[intel::reqd_sub_group_size]] spelling
136+
if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize && A.getScopeName() &&
137+
A.getScopeName()->isStr("intel")) {
138+
diagnoseDeprecatedAttribute(A, "sycl", "reqd_sub_group_size");
139+
return;
140+
}
141+
135142
// Diagnose SYCL 2020 spellings in later SYCL modes.
136143
if (getLangOpts().getSYCLVersion() >= LangOptions::SYCL_2020) {
137144
// All attributes in the cl vendor namespace are deprecated in favor of a

‎clang/test/CodeGenSYCL/kernel-op-calls.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@ class Functor1 {
1111
public:
1212
Functor1(){}
1313

14-
[[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
14+
[[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
1515

1616
[[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}
1717

‎clang/test/CodeGenSYCL/reqd-sub-group-size.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7,25 +7,25 @@ queue q;
77

88
class Functor16 {
99
public:
10-
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
10+
[[sycl::reqd_sub_group_size(16)]] void operator()() const {}
1111
};
1212

1313
template <int SIZE>
1414
class Functor2 {
1515
public:
16-
[[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
16+
[[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
1717
};
1818

1919
template <int N>
20-
[[intel::reqd_sub_group_size(N)]] void func() {}
20+
[[sycl::reqd_sub_group_size(N)]] void func() {}
2121

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

2727
h.single_task<class kernel_name3>(
28-
[]() [[intel::reqd_sub_group_size(4)]]{});
28+
[]() [[sycl::reqd_sub_group_size(4)]]{});
2929

3030
Functor2<2> f2;
3131
h.single_task<class kernel_name4>(f2);

‎clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,12 +7,12 @@ queue q;
77

88
class Functor {
99
public:
10-
[[intel::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
10+
[[sycl::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
1111
};
1212

1313
class Functor1 {
1414
public:
15-
[[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
15+
[[sycl::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
1616
};
1717

1818
template <int SIZE, int SIZE1, int SIZE2>

‎clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@ template <typename T> class Fobj {
66
public:
77
Fobj() {}
88
void operator()() const {
9-
auto L0 = []() [[intel::reqd_sub_group_size(4)]]{};
9+
auto L0 = []() [[sycl::reqd_sub_group_size(4)]]{};
1010
L0();
1111
}
1212
};

‎clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -9,14 +9,14 @@ queue q;
99

1010
class Functor16 {
1111
public:
12-
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
12+
[[sycl::reqd_sub_group_size(16)]] void operator()() const {}
1313
};
1414

1515
// Test that checks template parameter support on member function of class template.
1616
template <int SIZE>
1717
class KernelFunctor {
1818
public:
19-
[[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
19+
[[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
2020
};
2121

2222
// Test that checks template parameter support on function.
@@ -35,7 +35,7 @@ class KernelFunctor {
3535
// CHECK-NEXT: NonTypeTemplateParmDecl
3636
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12
3737
template <int N>
38-
[[intel::reqd_sub_group_size(N)]] void func() {}
38+
[[sycl::reqd_sub_group_size(N)]] void func() {}
3939

4040
int main() {
4141
q.submit([&](handler &h) {
@@ -52,14 +52,14 @@ int main() {
5252
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
5353
// CHECK-NEXT: value: Int 2
5454
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
55-
h.single_task<class kernel_name3>([]() [[intel::reqd_sub_group_size(2)]] {});
55+
h.single_task<class kernel_name3>([]() [[sycl::reqd_sub_group_size(2)]] {});
5656

5757
// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
5858
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
5959
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
6060
// CHECK-NEXT: value: Int 6
6161
// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
62-
h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(6)]] {});
62+
h.single_task<class kernel_name5>([]() [[sycl::reqd_sub_group_size(6)]] {});
6363

6464
// CHECK: FunctionDecl {{.*}}kernel_name_6
6565
// CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
@@ -79,8 +79,8 @@ int main() {
7979
// CHECK-NEXT: value: Int 8
8080
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
8181
// CHECK-NOT: IntelReqdSubGroupSizeAttr
82-
[]() [[intel::reqd_sub_group_size(8),
83-
intel::reqd_sub_group_size(8)]] {});
82+
[]() [[sycl::reqd_sub_group_size(8),
83+
sycl::reqd_sub_group_size(8)]] {});
8484
});
8585
func<12>();
8686
return 0;

‎clang/test/SemaSYCL/reqd-sub-group-size-host.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,9 @@
11
// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
22
// expected-no-diagnostics
33

4-
[[intel::reqd_sub_group_size(8)]] void fun() {}
4+
[[sycl::reqd_sub_group_size(8)]] void fun() {}
55

66
class Functor {
77
public:
8-
[[intel::reqd_sub_group_size(16)]] void operator()() {}
8+
[[sycl::reqd_sub_group_size(16)]] void operator()() {}
99
};

‎clang/test/SemaSYCL/reqd-sub-group-size.cpp

Lines changed: 28 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s
22

3-
// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel.
3+
// The test checks functionality of [[sycl::reqd_sub_group_size()]] attribute on SYCL kernel and [[intel::reqd_sub_group_size()]] is deprecated.
44

55

66
#include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -32,44 +32,51 @@ int main() {
3232
});
3333
return 0;
3434
}
35-
[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
36-
[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
35+
36+
[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
37+
[[sycl::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
3738
{
3839
}
3940

40-
[[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}}
41+
[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
42+
[[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}}
4143
A();
4244
}
45+
4346
// expected-note@+1 {{conflicting attribute is here}}
44-
[[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}}
47+
[[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}}
4548

4649
// expected-note@+3 {{conflicting attribute is here}}
4750
// expected-error@+2 {{conflicting attributes applied to a SYCL kernel}}
4851
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
49-
[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
52+
[[sycl::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
5053
sg_size2();
5154
}
5255

5356
// Test that checks support and functionality of reqd_sub_group_size attribute support on function.
5457

5558
// Tests for incorrect argument values for Intel reqd_sub_group_size attribute.
56-
[[intel::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
57-
[[intel::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
58-
[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
59-
[[intel::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
60-
[[intel::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
59+
[[sycl::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
60+
[[sycl::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
61+
[[sycl::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
62+
[[sycl::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
63+
[[sycl::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
6164

6265
// Diagnostic is emitted because the arguments mismatch.
63-
[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}}
64-
[[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}}
66+
[[sycl::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}}
67+
[[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}}
6568
[[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
6669

6770
// Make sure there's at least one argument passed.
6871
[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
6972

7073
// No diagnostic is emitted because the arguments match.
74+
[[sycl::reqd_sub_group_size(12)]] void same();
75+
[[sycl::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
76+
77+
// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
78+
// expected-warning@+1{{attribute 'intel::reqd_sub_group_size' is deprecated}}
7179
[[intel::reqd_sub_group_size(12)]] void same();
72-
[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
7380

7481
// No diagnostic because the attributes are synonyms with identical behavior.
7582
[[sycl::reqd_sub_group_size(12)]] void same(); // OK
@@ -80,7 +87,7 @@ template <typename Ty>
8087
// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
8188
// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
8289
// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
83-
[[intel::reqd_sub_group_size(Ty{})]] void func() {}
90+
[[sycl::reqd_sub_group_size(Ty{})]] void func() {}
8491

8592
struct S {};
8693
void test() {
@@ -97,18 +104,18 @@ void test() {
97104
int foo1();
98105
// expected-error@+2{{expression is not an integral constant expression}}
99106
// expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}}
100-
[[intel::reqd_sub_group_size(foo1() + 12)]] void func1();
107+
[[sycl::reqd_sub_group_size(foo1() + 12)]] void func1();
101108

102109
// Test that checks expression is a constant expression.
103110
constexpr int bar1() { return 0; }
104-
[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
111+
[[sycl::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
105112

106113
// Test that checks template parameter support on member function of class template.
107114
template <int SIZE>
108115
class KernelFunctor {
109116
public:
110117
// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
111-
[[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
118+
[[sycl::reqd_sub_group_size(SIZE)]] void operator()() {}
112119
};
113120

114121
int check() {
@@ -121,14 +128,14 @@ int check() {
121128
template <int N>
122129
// expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
123130
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
124-
[[intel::reqd_sub_group_size(N)]] void func3() {}
131+
[[sycl::reqd_sub_group_size(N)]] void func3() {}
125132

126133
template <int N>
127134
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
128-
[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
135+
[[sycl::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
129136

130137
template <int N>
131-
[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
138+
[[sycl::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
132139

133140
int check1() {
134141
// no error expected

‎clang/test/SemaSYCL/sub-group-size.cpp

Lines changed: 16 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -3,29 +3,29 @@
33
// 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
44

55
// Validate the semantic analysis checks for the interaction betwen the
6-
// named_sub_group_size and sub_group_size attributes. These are not able to be
6+
// named_sub_group_size and reqd_sub_group_size attributes. These are not able to be
77
// combined, and require that they only be applied to non-sycl-kernel/
88
// non-sycl-device functions if they match the kernel they are being called
99
// from.
1010

1111
#include "Inputs/sycl.hpp"
1212

13-
// expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
13+
// expected-error@+2 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}}
1414
// expected-note@+1 {{conflicting attribute is here}}
15-
[[intel::sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1();
16-
// expected-error@+2 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
15+
[[sycl::reqd_sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1();
16+
// expected-error@+2 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
1717
// expected-note@+1 {{conflicting attribute is here}}
18-
[[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2();
18+
[[intel::named_sub_group_size(primary)]] [[sycl::reqd_sub_group_size(1)]] void f2();
1919

2020
// expected-note@+1 {{conflicting attribute is here}}
21-
[[intel::sub_group_size(1)]] void f3();
22-
// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
21+
[[sycl::reqd_sub_group_size(1)]] void f3();
22+
// expected-error@+1 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}}
2323
[[intel::named_sub_group_size(primary)]] void f3();
2424

2525
// expected-note@+1 {{conflicting attribute is here}}
2626
[[intel::named_sub_group_size(primary)]] void f4();
27-
// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
28-
[[intel::sub_group_size(1)]] void f4();
27+
// expected-error@+1 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
28+
[[sycl::reqd_sub_group_size(1)]] void f4();
2929

3030
// expected-note@+1 {{previous attribute is here}}
3131
[[intel::named_sub_group_size(automatic)]] void f5();
@@ -115,8 +115,14 @@ void calls_kernel_3() {
115115
});
116116
}
117117

118+
// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
119+
// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
118120
[[intel::sub_group_size(10)]] void AttrFunc2() {} // #AttrFunc2
121+
// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
122+
// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
119123
[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {} // #AttrExternalDefined2
124+
// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
125+
// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
120126
[[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalNotDefined2(); // #AttrExternalNotDefined2
121127

122128
void calls_kernel_4() {
@@ -153,7 +159,7 @@ void calls_kernel_5() {
153159

154160
// Don't diag with the old sub-group-size.
155161
void calls_kernel_6() {
156-
sycl::kernel_single_task<class Kernel6>([]() [[intel::reqd_sub_group_size(10)]] { // #Kernel6
162+
sycl::kernel_single_task<class Kernel6>([]() [[sycl::reqd_sub_group_size(10)]] { // #Kernel6
157163
NoAttrExternalNotDefined();
158164
});
159165
}

‎clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
[[sycl::reqd_work_group_size(16)]] void f1(){ // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
88
}
99

10-
[[intel::reqd_sub_group_size(12)]] void f3(){ // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
10+
[[sycl::reqd_sub_group_size(12)]] void f3(){ // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
1111
}
1212

1313
[[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 {
4141

4242
class FunctorSubGroupSize4 {
4343
public:
44-
[[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}}
44+
[[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}}
4545
};
4646

4747
class Functor8 {

‎sycl/test-e2e/AOT/reqd-sg-size.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ template <size_t... Ns> struct SubgroupDispatcher {
4242
accessor acc{buf, cgh};
4343
cgh.parallel_for<kernel_name<size>>(
4444
nd_range<1>(1, 1),
45-
[=](auto item) [[intel::reqd_sub_group_size(size)]] {
45+
[=](auto item) [[sycl::reqd_sub_group_size(size)]] {
4646
acc[0] = item.get_sub_group().get_max_local_range()[0];
4747
});
4848
});

‎sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@
3232
#ifdef IMPL_SUBGROUP
3333
#define SUBGROUP_ATTR
3434
#else
35-
#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]]
35+
#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]]
3636
#endif
3737

3838
using namespace sycl;

‎sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -78,7 +78,7 @@ void RunSortOverGroupArray(sycl::queue &Q, const std::vector<T> &DataToSort,
7878
CGH);
7979

8080
CGH.parallel_for(
81-
NDRange, [=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
81+
NDRange, [=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
8282
ReqSubGroupSize)]] {
8383
const size_t GlobalLinearID = id.get_global_linear_id();
8484
using RadixSorterT = oneapi_exp::radix_sorters::group_sorter<

‎sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -115,7 +115,7 @@ void RunJointSort(sycl::queue &Q, const std::vector<T> &DataToSort,
115115

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

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

0 commit comments

Comments
 (0)