Skip to content

Commit d4eba22

Browse files
smanna12smanna12
committed
address @alexbatashev comments
Signed-off-by: Soumi Manna <[email protected]>
1 parent 8b7a71c commit d4eba22

18 files changed

+71
-38
lines changed

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
C[wiID] = 43;
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
asm volatile("");

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
volatile int output = 0;
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
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 & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
volatile int output = 0;
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
asm volatile("mov (M1,16) %0(0,0)<1> 0x7:d"

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
for (int i = 0; i < 10; ++i) {
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
asm("fence_sw");

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
C[wiID] = 43;
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
asm volatile("");

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
volatile int output = 0;
2527
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2628
asm volatile("mov (M1,8) %0(0,0)<1> 0x7:d"

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

Lines changed: 4 additions & 2 deletions
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-
cl::sycl::range<1>{this->getOutputBufferSize()}, [=
26-
](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
25+
// clang-format off
26+
cl::sycl::range<1>{this->getOutputBufferSize()},
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 & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,8 +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()}, [=
27-
](cl::sycl::id<1> wiID) [[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
2830
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2931
asm("add (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
3032
: "=rw"(C[wiID])

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +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()}, [=
25-
](cl::sycl::id<1> wiID) [[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
2628
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2729
asm("mov (M1, 8) %0(0, 0)<1> (-)%1(0, 0)<1;1,0>"
2830
: "=rw"(B[wiID])

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +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()}, [=
25-
](cl::sycl::id<1> wiID) [[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
2628
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2729
asm("mul (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
2830
: "=rw"(C[wiID])

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,8 +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()}, [=
30-
](cl::sycl::id<1> wiID) [[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
3133
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3234
asm("{\n"
3335
"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: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -28,8 +28,8 @@ int main() {
2828
NumOfWorkItems, [=](cl::sycl::id<1> WIid)
2929
[[intel::reqd_sub_group_size(8)]] {
3030
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
31-
asm("barrier");
31+
asm("barrier");
3232
#endif
33-
});
33+
});
3434
});
3535
}

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +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()}, [=
23-
](cl::sycl::id<1> wiID) [[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
2426
volatile int local_var = 47;
2527
local_var += C[0];
2628
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +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()}, [=
25-
](cl::sycl::id<1> wiID) [[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
2628
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
2729
asm("add (M1, 16) %0(0, 0)<1> %0(0, 0)<1;1,0> %1(0, 0)<1;1,0>"
2830
: "+rw"(B[wiID])

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,9 @@ int main() {
2929
q.submit([&](cl::sycl::handler &cgh) {
3030
cgh.parallel_for<kernel_name>(
3131
cl::sycl::range<1>(problem_size), [=](cl::sycl::id<1> idx)
32-
[[intel::reqd_sub_group_size(
33-
16)]] {
32+
// clang-format off
33+
[[intel::reqd_sub_group_size(16)]] {
34+
// clang-format on
3435
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3536
int i = idx[0];
3637
asm volatile("{\n.decl V52 v_type=G type=d num_elts=16 align=GRF\n"

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

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -35,8 +35,10 @@ int main() {
3535

3636
q.submit([&](cl::sycl::handler &cgh) {
3737
cgh.parallel_for<kernel_name>(
38-
cl::sycl::range<1>(problem_size), [=
39-
](cl::sycl::id<1> idx) [[intel::reqd_sub_group_size(32)]] {
38+
// clang-format off
39+
cl::sycl::range<1>(problem_size),
40+
[=](cl::sycl::id<1> idx) [[intel::reqd_sub_group_size(32)]] {
41+
// clang-format on
4042
int i = idx[0];
4143
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
4244
asm volatile(R"a(
@@ -57,8 +59,8 @@ int main() {
5759
svm_scatter.4.1 (M1, 16) %1.0 V53.0
5860
}
5961
)a" ::"rw"(&b[i]),
60-
"rw"(&b[i] + 16), "rw"(&a[i]), "rw"(&a[i] + 16),
61-
"rw"(&c[i]), "rw"(&c[i] + 16));
62+
"rw"(&b[i] + 16), "rw"(&a[i]), "rw"(&a[i] + 16), "rw"(&c[i]),
63+
"rw"(&c[i] + 16));
6264
#else
6365
b[i] = a[i] * c[i];
6466
#endif

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,10 @@ int main() {
3131

3232
q.submit([&](cl::sycl::handler &cgh) {
3333
cgh.parallel_for<kernel_name>(
34-
cl::sycl::range<1>(problem_size), [=
35-
](cl::sycl::id<1> idx) [[intel::reqd_sub_group_size(16)]] {
34+
// clang-format off
35+
cl::sycl::range<1>(problem_size),
36+
[=](cl::sycl::id<1> idx) [[intel::reqd_sub_group_size(16)]] {
37+
// clang-format on
3638
int i = idx[0];
3739
volatile int tmp = a[i];
3840
tmp += 1;

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

Lines changed: 4 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), [=
32-
](cl::sycl::id<1> idx) [[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
int i = idx[0];
3436
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
3537
asm volatile("mov (M1, 16) %0(0,0)<1> 0x7:d"

0 commit comments

Comments
 (0)