Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[SYCL]Use Subgroup size 16 as default for InlineASM tests #1476

Merged
merged 6 commits into from
Jan 4, 2023
Merged

[SYCL]Use Subgroup size 16 as default for InlineASM tests #1476

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
b3abba4
Subgroup size 8 is not supported on all gpu platforms. Instead use 16…
stdale-intel Dec 29, 2022
103fcaa
align to clang-format
stdale-intel Dec 29, 2022
b53e5c7
more clang-format fixes
stdale-intel Dec 29, 2022
f6d8ac2
more clang-format fixes
stdale-intel Jan 4, 2023
47ea92b
more clang-format fix for weird spacing
stdale-intel Jan 4, 2023
ce4525e
fix asm_loop still using :8 in some places
stdale-intel Jan 4, 2023
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
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/Negative/asm_bad_opcode.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,11 +13,11 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
"movi (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
"movi (M1_NM, 16) tmp1(0,1)<1> tmp2(0,0)\n");
#endif
});
}
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/Negative/asm_bad_operand_syntax.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,11 +13,11 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
"mov (M1_NM, 8) tmp1(0,1)<1>:f tmp2(0,0)<1;1,0>\n");
"mov (M1_NM, 16) tmp1(0,1)<1>:f tmp2(0,0)<1;1,0>\n");
#endif
});
}
Expand Down
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/Negative/asm_duplicate_label.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
Expand Down
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/Negative/asm_illegal_exec_size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
Expand Down
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/Negative/asm_missing_label.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/Negative/asm_missing_region.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,11 +13,11 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
"mov (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
"mov (M1_NM, 16) tmp1(0,1)<1> tmp2(0,0)\n");
#endif
});
}
Expand Down
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/Negative/asm_simple.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/Negative/asm_undefined_decl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,11 +13,11 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
"mov (M1_NM, 8) tmp1(0,1)<1> my_super_var(0,0)\n");
"mov (M1_NM, 16) tmp1(0,1)<1> my_super_var(0,0)\n");
#endif
});
}
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/Negative/asm_undefined_pred.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,11 +13,11 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
"cmp.lt (M1_NM, 8) P3 tmp1(0,0)<0;1,0> 0x3:ud\n");
"cmp.lt (M1_NM, 16) P3 tmp1(0,0)<0;1,0> 0x3:ud\n");
#endif
});
}
Expand Down
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/Negative/asm_wrong_declare.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ struct KernelFunctor {
void operator()(sycl::handler &cgh) {
cgh.parallel_for<KernelFunctor>(
sycl::range<1>{16},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_arbitrary_ops_order.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,9 +33,9 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("mad (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0> %3(0, "
asm("mad (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0> %3(0, "
"0)<1;1,0>"
: "=rw"(D[wiID])
: "rw"(A[wiID]), "rw"(B[wiID]), "rw"(C[wiID]));
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_float_add.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,9 +30,9 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("add (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
asm("add (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
: "=rw"(C[wiID])
: "rw"(A[wiID]), "rw"(B[wiID]));
#else
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_float_imm_arg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,9 +27,9 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("mul (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
asm("mul (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
: "=rw"(B[wiID])
: "rw"(A[wiID]), "i"(IMM_ARGUMENT));
#else
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_float_neg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,9 +25,9 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("mov (M1, 8) %0(0, 0)<1> (-)%1(0, 0)<1;1,0>"
asm("mov (M1, 16) %0(0, 0)<1> (-)%1(0, 0)<1;1,0>"
: "=rw"(B[wiID])
: "rw"(A[wiID]));
#else
Expand Down
6 changes: 3 additions & 3 deletions SYCL/InlineAsm/asm_if.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,13 +18,13 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
bool switchField = false;
CGH.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
int Output = 0;
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile("{\n"
".decl P1 v_type=P num_elts=1\n"
"cmp.eq (M1_NM, 8) P1 %1(0,0)<0;1,0> 0x0:b\n"
"(P1) sel (M1_NM, 8) %0(0,0)<1> 0x7:d 0x8:d"
"cmp.eq (M1_NM, 16) P1 %1(0,0)<0;1,0> 0x0:b\n"
"(P1) sel (M1_NM, 16) %0(0,0)<1> 0x7:d 0x8:d"
"}\n"
: "=rw"(Output)
: "rw"(switchField));
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_imm_arg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,9 +26,9 @@ struct KernelFunctor : WithInputBuffers<T, 1>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("add (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
asm("add (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
: "=rw"(B[wiID])
: "rw"(A[wiID]), "i"(CONST_ARGUMENT));
#else
Expand Down
24 changes: 12 additions & 12 deletions SYCL/InlineAsm/asm_loop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,21 +29,21 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
CGH);
CGH.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile("{\n"
".decl P1 v_type=P num_elts=8\n"
".decl P2 v_type=P num_elts=8\n"
".decl temp v_type=G type=d num_elts=8 align=dword\n"
"mov (M1, 8) %0(0, 0)<1> 0x0:d\n"
"cmp.le (M1, 8) P1 %1(0,0)<1;1,0> 0x0:d\n"
"(P1) goto (M1, 8) label0%=\n"
"mov (M1, 8) temp(0,0)<1> 0x0:d\n"
".decl P1 v_type=P num_elts=16\n"
".decl P2 v_type=P num_elts=16\n"
".decl temp v_type=G type=d num_elts=16 align=dword\n"
"mov (M1, 16) %0(0, 0)<1> 0x0:d\n"
"cmp.le (M1, 16) P1 %1(0,0)<1;1,0> 0x0:d\n"
"(P1) goto (M1, 16) label0%=\n"
"mov (M1, 16) temp(0,0)<1> 0x0:d\n"
"label1%=:\n"
"add (M1, 8) temp(0,0)<1> temp(0,0)<1;1,0> 0x1:w\n"
"add (M1, 8) %0(0,0)<1> %0(0,0)<1;1,0> %2(0,0)<1;1,0>\n"
"cmp.lt (M1, 8) P2 temp(0,0)<0;8,1> %1(0,0)<0;8,1>\n"
"(P2) goto (M1, 8) label1%=\n"
"add (M1, 16) temp(0,0)<1> temp(0,0)<1;1,0> 0x1:w\n"
"add (M1, 16) %0(0,0)<1> %0(0,0)<1;1,0> %2(0,0)<1;1,0>\n"
"cmp.lt (M1, 16) P2 temp(0,0)<0;16,1> %1(0,0)<0;16,1>\n"
"(P2) goto (M1, 16) label1%=\n"
"label0%=:"
"}\n"
: "+rw"(C[wiID])
Expand Down
4 changes: 2 additions & 2 deletions SYCL/InlineAsm/asm_mul.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("mul (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
asm("mul (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
: "=rw"(C[wiID])
: "rw"(A[wiID]), "rw"(B[wiID]));
#else
Expand Down
17 changes: 9 additions & 8 deletions SYCL/InlineAsm/asm_multiple_instructions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,17 +36,17 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {

cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
#if defined(TO_PASS)
// The code below passing verification
volatile int output = -1;

#if defined(__SYCL_DEVICE_ONLY__)
asm volatile(
"{\n"
"add (M1, 8) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"
"add (M1, 8) %1(0, 0)<1> %1(0, 0)<1;1,0> %3(0, 0)<1;1,0>\n"
"mov (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0>\n"
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %3(0, 0)<1;1,0>\n"
"mov (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0>\n"
"}\n"
: "=rw"(output), "+rw"(A[wiID])
: "rw"(B[wiID]), "rw"(C[wiID]));
Expand All @@ -58,10 +58,11 @@ struct KernelFunctor : WithInputBuffers<T, 3>, WithOutputBuffer<T> {
D[wiID] = output;
#else
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile("{\n"
"add (M1, 8) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"
"add (M1, 8) %1(0, 0)<1> %1(0, 0)<1;1,0> %3(0, 0)<1;1,0>\n"
"mov (M1, 8) %0(0, 0)<1> %1(0, 0)<1;1,0>\n"
asm volatile(
"{\n"
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>\n"
"add (M1, 16) %1(0, 0)<1> %1(0, 0)<1;1,0> %3(0, 0)<1;1,0>\n"
"mov (M1, 16) %0(0, 0)<1> %1(0, 0)<1;1,0>\n"
"}\n"
: "=rw"(D[wiID]), "+rw"(A[wiID])
: "rw"(B[wiID]), "rw"(C[wiID]));
Expand Down
10 changes: 5 additions & 5 deletions SYCL/InlineAsm/asm_no_operands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,12 +25,12 @@ int main() {
// Submitting command group(work) to queue
Queue.submit([&](sycl::handler &cgh) {
// Executing kernel
cgh.parallel_for<no_operands_kernel>(NumOfWorkItems,
[=](sycl::id<1> WIid)
[[intel::reqd_sub_group_size(8)]] {
cgh.parallel_for<no_operands_kernel>(
NumOfWorkItems,
[=](sycl::id<1> WIid) [[intel::reqd_sub_group_size(16)]] {
#if defined(__SYCL_DEVICE_ONLY__)
asm("barrier");
asm("barrier");
#endif
});
});
});
}
2 changes: 1 addition & 1 deletion SYCL/InlineAsm/asm_no_output.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ template <typename T = dataType> struct KernelFunctor : WithOutputBuffer<T> {
cgh);
cgh.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
volatile int local_var = 47;
local_var += C[0];
#if defined(__SYCL_DEVICE_ONLY__)
Expand Down
10 changes: 5 additions & 5 deletions SYCL/InlineAsm/asm_switch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,25 +18,25 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
int switchField = 2;
CGH.parallel_for<KernelFunctor<T>>(
sycl::range<1>{this->getOutputBufferSize()},
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
[=](sycl::id<1> wiID) [[intel::reqd_sub_group_size(16)]] {
int Output = 0;
#if defined(__SYCL_DEVICE_ONLY__)
asm volatile("{\n"
".decl P1 v_type=P num_elts=1\n"
".decl P2 v_type=P num_elts=1\n"
".decl P3 v_type=P num_elts=1\n"
"cmp.ne (M1_NM, 8) P1 %1(0,0)<0;1,0> 0x0:d\n"
"cmp.ne (M1_NM, 16) P1 %1(0,0)<0;1,0> 0x0:d\n"
"(P1) goto (M1, 1) label0%=\n"
"mov (M1, 8) %0(0,0)<1> 0x9:d\n"
"(P1) goto (M1, 1) label0%=\n"
"label0%=:\n"
"cmp.ne (M1_NM, 8) P2 %1(0,0)<0;1,0> 0x1:d\n"
"cmp.ne (M1_NM, 16) P2 %1(0,0)<0;1,0> 0x1:d\n"
"(P2) goto (M1, 1) label1%=\n"
"mov (M1, 8) %0(0,0)<1> 0x8:d\n"
"label1%=:\n"
"cmp.ne (M1_NM, 8) P3 %1(0,0)<0;1,0> 0x2:d\n"
"cmp.ne (M1_NM, 16) P3 %1(0,0)<0;1,0> 0x2:d\n"
"(P3) goto (M1, 1) label2%=\n"
"mov (M1, 8) %0(0,0)<1> 0x7:d\n"
"mov (M1, 16) %0(0,0)<1> 0x7:d\n"
"label2%=:"
"}\n"
: "=rw"(Output)
Expand Down