Skip to content

Commit 9dda36f

Browse files
smanna12smanna12
authored
[NFC][SYCL] Remove deprecated spelling of IntelReqdSubGroupSize attribute (#2399)
Commit b2da2c8 added support for new spelling [[intel::reqd_sub_group_size()]] and deprecated previous spelling [[cl::intel_reqd_sub_group_size()]] of the IntelReqdSubGroupSize attribute to match with spec. This patch removes that deprecated spelling and updates several tests of IntelReqdSubGroupSize attribute with new spelling. Signed-off-by: Soumi Manna <[email protected]>
1 parent 66484f2 commit 9dda36f

33 files changed

+95
-72
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1299,7 +1299,6 @@ def LoopUnrollHint : InheritableAttr {
12991299

13001300
def IntelReqdSubGroupSize: InheritableAttr {
13011301
let Spellings = [GNU<"intel_reqd_sub_group_size">,
1302-
CXX11<"cl", "intel_reqd_sub_group_size">,
13031302
CXX11<"intel", "reqd_sub_group_size">];
13041303
let Args = [ExprArgument<"SubGroupSize">];
13051304
let Subjects = SubjectList<[Function, CXXMethod], ErrorDiag>;

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 3 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -3476,30 +3476,16 @@ code. See `cl_intel_required_subgroup_size
34763476
for details.
34773477

34783478
SYCL documentation:
3479-
The [[cl::intel_reqd_sub_group_size(n)]] and [[intel::reqd_sub_group_size(n)]]
3480-
attribute indicates that the kernel must be compiled and executed with a
3481-
sub-group of size n. The value of n must be set to a sub-group size supported
3482-
by the device, or device compilation will fail.
3479+
The [[intel::reqd_sub_group_size(n)]] attribute indicates that the kernel must
3480+
be compiled and executed with a sub-group of size n. The value of n must be set
3481+
to a sub-group size supported by the device, or device compilation will fail.
34833482

34843483
In addition to device functions, the required sub-group size attribute may also
34853484
be specified in the definition of a named functor object and lambda functions,
34863485
as in the examples below:
34873486

34883487
.. code-block:: c++
34893488

3490-
class Functor
3491-
{
3492-
void operator()(item<1> item) [[cl::intel_reqd_sub_group_size(16)]]
3493-
{
3494-
/* kernel code */
3495-
}
3496-
}
3497-
3498-
kernel<class kernel_name>(
3499-
[]() [[cl::intel_reqd_sub_group_size(n)]] {
3500-
/* kernel code */
3501-
});
3502-
35033489
class Functor
35043490
{
35053491
[[intel::reqd_sub_group_size(16)]] void operator()(item<1> item)

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -11070,11 +11070,6 @@ def err_ivdep_declrefexpr_arg : Error<
1107011070
def warn_ivdep_redundant : Warning <"ignoring redundant Intel FPGA loop "
1107111071
"attribute 'ivdep': safelen %select{INF|%1}0 >= safelen %select{INF|%3}2">,
1107211072
InGroup<IgnoredAttributes>;
11073-
def warn_attribute_spelling_deprecated : Warning<
11074-
"attribute %0 is deprecated">,
11075-
InGroup<DeprecatedAttributes>;
11076-
def note_spelling_suggestion : Note<
11077-
"did you mean to use %0 instead?">;
1107811073

1107911074
// errors of expect.with.probability
1108011075
def err_probability_not_constant_float : Error<

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -3013,13 +3013,6 @@ static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
30133013
if (D->getAttr<IntelReqdSubGroupSizeAttr>())
30143014
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
30153015

3016-
if (AL.getAttributeSpellingListIndex() ==
3017-
IntelReqdSubGroupSizeAttr::CXX11_cl_intel_reqd_sub_group_size) {
3018-
S.Diag(AL.getLoc(), diag::warn_attribute_spelling_deprecated) << AL;
3019-
S.Diag(AL.getLoc(), diag::note_spelling_suggestion)
3020-
<< "'intel::reqd_sub_group_size'";
3021-
}
3022-
30233016
S.addIntelReqdSubGroupSizeAttr(D, AL, E);
30243017
}
30253018

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,10 +2,10 @@
22

33
class Functor16 {
44
public:
5-
[[cl::intel_reqd_sub_group_size(16)]] void operator()() const {}
5+
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
66
};
77

8-
[[cl::intel_reqd_sub_group_size(8)]] void foo() {}
8+
[[intel::reqd_sub_group_size(8)]] void foo() {}
99

1010
class Functor {
1111
public:
@@ -17,7 +17,7 @@ class Functor {
1717
template <int SIZE>
1818
class Functor5 {
1919
public:
20-
[[cl::intel_reqd_sub_group_size(SIZE)]] void operator()() const {}
20+
[[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
2121
};
2222

2323
template <typename name, typename Func>
@@ -33,7 +33,7 @@ void bar() {
3333
kernel<class kernel_name2>(f);
3434

3535
kernel<class kernel_name3>(
36-
[]() [[cl::intel_reqd_sub_group_size(4)]] {});
36+
[]() [[intel::reqd_sub_group_size(4)]]{});
3737

3838
Functor5<2> f5;
3939
kernel<class kernel_name4>(f5);

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22

33
class Functor {
44
public:
5-
[[cl::intel_reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
5+
[[intel::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
66
};
77

88
template <typename Name, typename Func>

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

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7,9 +7,7 @@
77

88
class Functor16 {
99
public:
10-
// expected-warning@+2 {{attribute 'intel_reqd_sub_group_size' is deprecated}}
11-
// expected-note@+1 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
12-
[[cl::intel_reqd_sub_group_size(16)]] void operator()() const {}
10+
[[intel::reqd_sub_group_size(16)]] void operator()() const {}
1311
};
1412

1513
class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
@@ -55,9 +53,7 @@ void bar() {
5553

5654
kernel<class kernel_name5>([]() [[intel::reqd_sub_group_size(2)]]{});
5755
kernel<class kernel_name6>([]() [[intel::reqd_sub_group_size(4)]] { foo(); });
58-
// expected-warning@+2 {{attribute 'intel_reqd_sub_group_size' is deprecated}}
59-
// expected-note@+1 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
60-
kernel<class kernel_name7>([]() [[cl::intel_reqd_sub_group_size(6)]]{});
56+
kernel<class kernel_name7>([]() [[intel::reqd_sub_group_size(6)]]{});
6157

6258
Functor4 f4;
6359
kernel<class kernel_name8>(f4);

‎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 -fsycl-is-host -fsyntax-only -verify %s
22
// expected-no-diagnostics
33

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

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

‎sycl/test/inline-asm/asm_16_empty.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
25+
// clang-format on
2326
C[wiID] = 43;
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm volatile("");

‎sycl/test/inline-asm/asm_16_matrix_mult.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
25+
// clang-format on
2326
volatile int output = 0;
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm volatile("mov (M1,16) %0(0,0)<1> 0x7:d"

‎sycl/test/inline-asm/asm_16_no_input_int.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
25+
// clang-format on
2326
volatile int output = 0;
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm volatile("mov (M1,16) %0(0,0)<1> 0x7:d"

‎sycl/test/inline-asm/asm_16_no_opts.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
25+
// clang-format on
2326
for (int i = 0; i < 10; ++i) {
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm("fence_sw");

‎sycl/test/inline-asm/asm_8_empty.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
25+
// clang-format on
2326
C[wiID] = 43;
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm volatile("");

‎sycl/test/inline-asm/asm_8_no_input_int.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
25+
// clang-format on
2326
volatile int output = 0;
2427
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2528
asm volatile("mov (M1,8) %0(0,0)<1> 0x7:d"

‎sycl/test/inline-asm/asm_arbitrary_ops_order.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {
2727

2828
cgh.parallel_for<KernelFunctor<T>>(
2929
cl::sycl::range<1>{this->getOutputBufferSize()}, [=
30-
](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
30+
](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
3131
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3232
asm("mad (M1, 8) %0(0, 0)<1> %3(0, 0)<1;1,0> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
3333
: "=rw"(D[wiID])

‎sycl/test/inline-asm/asm_decl_in_scope.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,10 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
2222
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2323

2424
cgh.parallel_for<KernelFunctor<T>>(
25+
// clang-format off
2526
cl::sycl::range<1>{this->getOutputBufferSize()},
26-
[=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
27+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
28+
// clang-format on
2729
// declaration of temp within and outside the scope
2830
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2931
asm("{\n"

‎sycl/test/inline-asm/asm_float_add.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,10 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
2323
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2424

2525
cgh.parallel_for<KernelFunctor<T>>(
26-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
26+
// clang-format off
27+
cl::sycl::range<1>{this->getOutputBufferSize()},
28+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
29+
// clang-format on
2730
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2831
asm("add (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
2932
: "=rw"(C[wiID])

‎sycl/test/inline-asm/asm_float_imm_arg.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {
2424

2525
cgh.parallel_for<KernelFunctor<T>>(
2626
cl::sycl::range<1>{this->getOutputBufferSize()}, [=
27-
](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
27+
](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
2828
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2929
asm("mul (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
3030
: "=rw"(B[wiID])

‎sycl/test/inline-asm/asm_float_neg.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,10 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {
2121
auto B = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2222

2323
cgh.parallel_for<KernelFunctor<T>>(
24-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
24+
// clang-format off
25+
cl::sycl::range<1>{this->getOutputBufferSize()},
26+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
27+
// clang-format on
2528
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2629
asm("mov (M1, 8) %0(0, 0)<1> (-)%1(0, 0)<1;1,0>"
2730
: "=rw"(B[wiID])

‎sycl/test/inline-asm/asm_if.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
2020
// clang-format off
2121
CGH.parallel_for<KernelFunctor<T>>(
2222
cl::sycl::range<1>{this->getOutputBufferSize()},
23-
[=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
23+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
2424
// clang-format on
2525
int Output = 0;
2626
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)

‎sycl/test/inline-asm/asm_imm_arg.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {
2323

2424
cgh.parallel_for<KernelFunctor<T>>(
2525
cl::sycl::range<1>{this->getOutputBufferSize()}, [=
26-
](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
26+
](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
2727
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2828
asm("add (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
2929
: "=rw"(B[wiID])

‎sycl/test/inline-asm/asm_loop.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
2929
// clang-format off
3030
CGH.parallel_for<KernelFunctor<T>>(
3131
cl::sycl::range<1>{this->getOutputBufferSize()},
32-
[=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
32+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
3333
// clang-format on
3434
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3535
asm volatile(".decl P1 v_type=P num_elts=8\n"

‎sycl/test/inline-asm/asm_mul.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,10 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
2121
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2222

2323
cgh.parallel_for<KernelFunctor<T>>(
24-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
24+
// clang-format off
25+
cl::sycl::range<1>{this->getOutputBufferSize()},
26+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
27+
// clang-format on
2528
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2629
asm("mul (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
2730
: "=rw"(C[wiID])

‎sycl/test/inline-asm/asm_multiple_instructions.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,10 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {
2626
auto D = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2727

2828
cgh.parallel_for<KernelFunctor<T>>(
29-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
29+
// clang-format off
30+
cl::sycl::range<1>{this->getOutputBufferSize()},
31+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
32+
// clang-format on
3033
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3134
asm("{\n"
3235
"add (M1, 8) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"

‎sycl/test/inline-asm/asm_no_operands.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,10 +25,13 @@ int main() {
2525
Queue.submit([&](cl::sycl::handler &cgh) {
2626
// Executing kernel
2727
cgh.parallel_for<no_operands_kernel>(
28-
NumOfWorkItems, [=](cl::sycl::id<1> WIid) [[cl::intel_reqd_sub_group_size(8)]] {
28+
NumOfWorkItems, [=](cl::sycl::id<1> WIid)
29+
[[intel::reqd_sub_group_size(8)]] {
30+
// clang-format off
2931
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3032
asm("barrier");
3133
#endif
3234
});
35+
// clang-format on
3336
});
3437
}

‎sycl/test/inline-asm/asm_no_output.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,10 @@ struct KernelFunctor : WithOutputBuffer<T> {
1919
void operator()(cl::sycl::handler &cgh) {
2020
auto C = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2121
cgh.parallel_for<KernelFunctor<T>>(
22-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
22+
// clang-format off
23+
cl::sycl::range<1>{this->getOutputBufferSize()},
24+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
25+
// clang-format on
2326
volatile int local_var = 47;
2427
local_var += C[0];
2528
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)

‎sycl/test/inline-asm/asm_plus_mod.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,10 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {
2121
auto B = this->getOutputBuffer().template get_access<cl::sycl::access::mode::write>(cgh);
2222

2323
cgh.parallel_for<KernelFunctor<T>>(
24-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(16)]] {
24+
// clang-format off
25+
cl::sycl::range<1>{this->getOutputBufferSize()},
26+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
27+
// clang-format on
2528
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2629
asm("add (M1, 16) %0(0, 0)<1> %0(0, 0)<1;1,0> %1(0, 0)<1;1,0>"
2730
: "+rw"(B[wiID])

‎sycl/test/inline-asm/asm_switch.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
2020
// clang-format off
2121
CGH.parallel_for<KernelFunctor<T>>(
2222
cl::sycl::range<1>{this->getOutputBufferSize()},
23-
[=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
23+
[=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
2424
// clang-format on
2525
int Output = 0;
2626
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)

‎sycl/test/inline-asm/letter_example.cpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -28,8 +28,10 @@ int main() {
2828
}
2929
q.submit([&](cl::sycl::handler &cgh) {
3030
cgh.parallel_for<kernel_name>(
31-
cl::sycl::range<1>(problem_size), [=](cl::sycl::id<1> idx)
32-
[[cl::intel_reqd_sub_group_size(16)]] {
31+
// clang-format off
32+
cl::sycl::range<1>(problem_size),
33+
[=](cl::sycl::id<1> idx) [[intel::reqd_sub_group_size(16)]] {
34+
// clang-format on
3335
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3436
int i = idx[0];
3537
asm volatile("{\n.decl V52 v_type=G type=d num_elts=16 align=GRF\n"
@@ -39,7 +41,9 @@ int main() {
3941
:
4042
: "rw"(&a[i]));
4143
#else
44+
// clang-format off
4245
a[idx[0]]++;
46+
// clang-format on
4347
#endif
4448
});
4549
}).wait();

0 commit comments

Comments
 (0)