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

stdale-intel · web-flow · commit 417b60f355ae · 2023-01-04T11:25:12.000-08:00
Subgroup size 8 is not supported on all gpu platforms. Instead use 16 as default subgroup size unless specifically testing a given size (8 or 32 for example in some tests).
diff --git a/SYCL/InlineAsm/Negative/asm_bad_opcode.cpp b/SYCL/InlineAsm/Negative/asm_bad_opcode.cpp
@@ -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
         });
   }
diff --git a/SYCL/InlineAsm/Negative/asm_bad_operand_syntax.cpp b/SYCL/InlineAsm/Negative/asm_bad_operand_syntax.cpp
@@ -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
         });
   }
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,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
         });
   }
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,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
         });
   }
diff --git a/SYCL/InlineAsm/Negative/asm_undefined_pred.cpp b/SYCL/InlineAsm/Negative/asm_undefined_pred.cpp
@@ -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
         });
   }
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_arbitrary_ops_order.cpp b/SYCL/InlineAsm/asm_arbitrary_ops_order.cpp
@@ -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]));
diff --git a/SYCL/InlineAsm/asm_float_add.cpp b/SYCL/InlineAsm/asm_float_add.cpp
@@ -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
diff --git a/SYCL/InlineAsm/asm_float_imm_arg.cpp b/SYCL/InlineAsm/asm_float_imm_arg.cpp
@@ -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
diff --git a/SYCL/InlineAsm/asm_float_neg.cpp b/SYCL/InlineAsm/asm_float_neg.cpp
@@ -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
diff --git a/SYCL/InlineAsm/asm_if.cpp b/SYCL/InlineAsm/asm_if.cpp
@@ -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));
diff --git a/SYCL/InlineAsm/asm_imm_arg.cpp b/SYCL/InlineAsm/asm_imm_arg.cpp
@@ -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
diff --git a/SYCL/InlineAsm/asm_loop.cpp b/SYCL/InlineAsm/asm_loop.cpp
@@ -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])
diff --git a/SYCL/InlineAsm/asm_mul.cpp b/SYCL/InlineAsm/asm_mul.cpp
@@ -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
diff --git a/SYCL/InlineAsm/asm_multiple_instructions.cpp b/SYCL/InlineAsm/asm_multiple_instructions.cpp
@@ -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]));
@@ -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]));
diff --git a/SYCL/InlineAsm/asm_no_operands.cpp b/SYCL/InlineAsm/asm_no_operands.cpp
@@ -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
-                                             });
+        });
   });
 }
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,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)
