[SYCLomatic] Add asm test case in test_feature.py (#359)

Signed-off-by: Wang, Yihan <[email protected]>

Author: Yihan Wang

features/config/TEMPLATE_asm.xml

@@ -5,4 +5,7 @@
     <files>
         <file path="feature_case/asm/${testName}.cu" />
     </files>
+    <rules>
+        <optlevelRule GPUFeature="NOT double" excludeOptlevelNameString="gpu" />
+    </rules>
 </test>

features/feature_case/asm/asm.cu

@@ -8,15 +8,15 @@
 // ===--------------------------------------------------------------------===//
 //
 // This file defines the test cases to check inline device asm parser.
-// 
+//
 // (1) Floating point/integer constant.
 // (2) Binary/Unary/Conditional operators and paren expressions.
 // (3) Compound statements.
 // (4) Conditional instructions.
 // (5) Instructions(mov, setp, and lop3).
 //
-// Usually, we check the result of inline asm statement to ensure that the migrated 
-// programe has the same behavior with the inline asmstatement.
+// Usually, we check the result of inline asm statement to ensure that the
+// migrated programe has the same behavior with the inline asmstatement.
 //
 //===----------------------------------------------------------------------===//
 
@@ -1386,36 +1386,34 @@ __global__ void lop3(int *ec) {
 
 int main() {
   int ret = 0;
-  int *ec = nullptr;
-  cudaMalloc(&ec, sizeof(int));
-
+  int *d_ec = nullptr;
+  cudaMalloc(&d_ec, sizeof(int));
 
   auto wait_and_check = [&](const char *case_name) {
     cudaDeviceSynchronize();
-    int res = 0;
-    cudaMemcpy(&res, ec, sizeof(int), cudaMemcpyDeviceToHost);
-    if (res != 0)
-      printf("Test %s failed: return code = %d\n", case_name, res);
+    int ec = 0;
+    cudaMemcpy(&ec, d_ec, sizeof(int), cudaMemcpyDeviceToHost);
+    if (ec != 0)
+      printf("Test %s failed: return code = %d\n", case_name, ec);
     ret = ret || ec;
   };
 
-
-  floating_point<<<1, 1>>>(ec);
+  floating_point<<<1, 1>>>(d_ec);
   wait_and_check("floating point");
 
-  integer_literal<<<1, 1>>>(ec);
+  integer_literal<<<1, 1>>>(d_ec);
   wait_and_check("integer literal");
 
-  expression<<<1, 1>>>(ec);
+  expression<<<1, 1>>>(d_ec);
   wait_and_check("expression");
-  
-  declaration<<<1, 1>>>(ec);
+
+  declaration<<<1, 1>>>(d_ec);
   wait_and_check("declaration");
 
-  setp<<<1, 1>>>(ec);
+  setp<<<1, 1>>>(d_ec);
   wait_and_check("setp");
 
-  lop3<<<1, 1>>>(ec);
+  lop3<<<1, 1>>>(d_ec);
   wait_and_check("lop3");
 
   return ret;

features/test_feature.py

@@ -17,7 +17,7 @@
 
 from test_utils import *
 
-exec_tests = ['thrust-vector-2', 'thrust-binary-search', 'thrust-count', 'thrust-copy',
+exec_tests = ['asm', 'thrust-vector-2', 'thrust-binary-search', 'thrust-count', 'thrust-copy',
               'thrust-qmc', 'thrust-transform-if', 'thrust-policy', 'thrust-list', 'module-kernel',
               'kernel-launch', 'thrust-gather', 'thrust-gather_if',
               'thrust-scatter', 'thrust-unique_by_key_copy', 'thrust-for-hypre', 'thrust-merge_by_key',