Merge pull request intel#1353 from WangYutao1995/shliclel725-20221026-topic

WangYutao1995 · web-flow · commit 89464b4d8d26 · 2022-10-27T13:18:51.000+08:00
[SYCL] Update inlineASM tests to use subgroup size 16 instead of 8.
diff --git a/SYCL/InlineAsm/Negative/asm_bad_opcode.cpp b/SYCL/InlineAsm/Negative/asm_bad_opcode.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_bad_operand_syntax.cpp b/SYCL/InlineAsm/Negative/asm_bad_operand_syntax.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_duplicate_label.cpp b/SYCL/InlineAsm/Negative/asm_duplicate_label.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_illegal_exec_size.cpp b/SYCL/InlineAsm/Negative/asm_illegal_exec_size.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_missing_label.cpp b/SYCL/InlineAsm/Negative/asm_missing_label.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_missing_region.cpp b/SYCL/InlineAsm/Negative/asm_missing_region.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_simple.cpp b/SYCL/InlineAsm/Negative/asm_simple.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_undefined_decl.cpp b/SYCL/InlineAsm/Negative/asm_undefined_decl.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_undefined_pred.cpp b/SYCL/InlineAsm/Negative/asm_undefined_pred.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/Negative/asm_wrong_declare.cpp b/SYCL/InlineAsm/Negative/asm_wrong_declare.cpp
@@ -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"
diff --git a/SYCL/InlineAsm/asm_8_empty.cpp b/SYCL/InlineAsm/asm_8_empty.cpp
@@ -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)]] {
           C[wiID] = 43;
 #if defined(__SYCL_DEVICE_ONLY__)
           asm volatile("");
diff --git a/SYCL/InlineAsm/asm_8_no_input_int.cpp b/SYCL/InlineAsm/asm_8_no_input_int.cpp
@@ -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 output = 0;
 #if defined(__SYCL_DEVICE_ONLY__)
           asm volatile("mov (M1,16) %0(0,0)<1> 0x7:d" : "=rw"(output));
diff --git a/SYCL/InlineAsm/asm_arbitrary_ops_order.cpp b/SYCL/InlineAsm/asm_arbitrary_ops_order.cpp
@@ -33,7 +33,7 @@ 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, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0> %3(0, "
               "0)<1;1,0>"
diff --git a/SYCL/InlineAsm/asm_float_add.cpp b/SYCL/InlineAsm/asm_float_add.cpp
@@ -30,7 +30,7 @@ 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, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
               : "=rw"(C[wiID])
diff --git a/SYCL/InlineAsm/asm_float_imm_arg.cpp b/SYCL/InlineAsm/asm_float_imm_arg.cpp
@@ -27,7 +27,7 @@ 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, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
               : "=rw"(B[wiID])
diff --git a/SYCL/InlineAsm/asm_float_neg.cpp b/SYCL/InlineAsm/asm_float_neg.cpp
@@ -25,7 +25,7 @@ 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, 16) %0(0, 0)<1> (-)%1(0, 0)<1;1,0>"
               : "=rw"(B[wiID])
diff --git a/SYCL/InlineAsm/asm_if.cpp b/SYCL/InlineAsm/asm_if.cpp
@@ -18,7 +18,7 @@ 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"
diff --git a/SYCL/InlineAsm/asm_imm_arg.cpp b/SYCL/InlineAsm/asm_imm_arg.cpp
@@ -26,7 +26,7 @@ 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, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2"
               : "=rw"(B[wiID])
diff --git a/SYCL/InlineAsm/asm_loop.cpp b/SYCL/InlineAsm/asm_loop.cpp
@@ -29,7 +29,7 @@ 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=16\n"
diff --git a/SYCL/InlineAsm/asm_mul.cpp b/SYCL/InlineAsm/asm_mul.cpp
@@ -28,7 +28,7 @@ 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, 16) %0(0, 0)<1> %1(0, 0)<1;1,0> %2(0, 0)<1;1,0>"
               : "=rw"(C[wiID])
diff --git a/SYCL/InlineAsm/asm_multiple_instructions.cpp b/SYCL/InlineAsm/asm_multiple_instructions.cpp
@@ -36,7 +36,7 @@ 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;
diff --git a/SYCL/InlineAsm/asm_no_operands.cpp b/SYCL/InlineAsm/asm_no_operands.cpp
@@ -27,7 +27,7 @@ int main() {
     // Executing kernel
     cgh.parallel_for<no_operands_kernel>(NumOfWorkItems,
                                          [=](sycl::id<1> WIid)
-                                             [[intel::reqd_sub_group_size(8)]] {
+                                             [[intel::reqd_sub_group_size(16)]] {
 #if defined(__SYCL_DEVICE_ONLY__)
                                                asm("barrier");
 #endif
diff --git a/SYCL/InlineAsm/asm_no_output.cpp b/SYCL/InlineAsm/asm_no_output.cpp
@@ -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__)
diff --git a/SYCL/InlineAsm/asm_switch.cpp b/SYCL/InlineAsm/asm_switch.cpp
@@ -18,7 +18,7 @@ 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"
