[SYCL] Localize variables declared in inline asms. (#105)

Signed-off-by: rdeodhar <[email protected]>

Author: rdeodhar

SYCL/InlineAsm/asm_if.cpp

@@ -21,9 +21,11 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
     ](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
           int Output = 0;
 #if defined(__SYCL_DEVICE_ONLY__)
-          asm volatile(".decl P1 v_type=P num_elts=1\n"
+          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"
+                       "}\n"
                        : "=rw"(Output)
                        : "rw"(switchField));
 

SYCL/InlineAsm/asm_loop.cpp

@@ -29,7 +29,8 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
         cl::sycl::range<1>{this->getOutputBufferSize()}, [=
     ](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
 #if defined(__SYCL_DEVICE_ONLY__)
-          asm volatile(".decl P1 v_type=P num_elts=8\n"
+          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"
@@ -42,6 +43,7 @@ struct KernelFunctor : WithInputBuffers<T, 2>, WithOutputBuffer<T> {
                        "cmp.lt (M1, 8) P2 temp(0,0)<0;8,1> %1(0,0)<0;8,1>\n"
                        "(P2) goto (M1, 8) label1\n"
                        "label0:"
+                       "}\n"
                        : "+rw"(C[wiID])
                        : "rw"(A[wiID]), "rw"(B[wiID]));
 #else

SYCL/InlineAsm/asm_switch.cpp

@@ -21,7 +21,8 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
     ](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
           int Output = 0;
 #if defined(__SYCL_DEVICE_ONLY__)
-          asm volatile(".decl P1 v_type=P num_elts=1\n"
+          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"
@@ -37,6 +38,7 @@ template <typename T = DataType> struct KernelFunctor : WithOutputBuffer<T> {
                        "(P3) goto (M1, 1) label2\n"
                        "mov (M1, 8) %0(0,0)<1> 0x7:d\n"
                        "label2:"
+                       "}\n"
                        : "=rw"(Output)
                        : "rw"(switchField));