Add missing setp tests and fix bugs in bra support

ptx/src/test/spirv_run/bra.ptx

@@ -0,0 +1,28 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry bra(
+	.param .u64 input,
+	.param .u64 output
+)
+{
+	.reg .u64 	    in_addr;
+    .reg .u64 	    out_addr;
+    .reg .u64 	    temp;
+    .reg .u64 	    temp2;
+
+	ld.param.u64 	in_addr, [input];
+    ld.param.u64 	out_addr, [output];
+
+    ld.u64          temp, [in_addr];
+	bra case1;
+case1:
+    add.u64         temp2, temp, 1;
+    bra case3;
+case2:
+    add.u64         temp2, temp, 2;
+case3:
+    st.u64          [out_addr], temp2;
+	ret;
+}

ptx/src/test/spirv_run/bra.spvtxt

@@ -0,0 +1,49 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability Int8
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %5 "bra"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+          %4 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+    %ulong_1 = OpConstant %ulong 1
+    %ulong_2 = OpConstant %ulong 2
+          %5 = OpFunction %void None %4
+          %6 = OpFunctionParameter %ulong
+          %7 = OpFunctionParameter %ulong
+         %27 = OpLabel
+         %11 = OpVariable %_ptr_Function_ulong Function
+         %12 = OpVariable %_ptr_Function_ulong Function
+         %13 = OpVariable %_ptr_Function_ulong Function
+         %14 = OpVariable %_ptr_Function_ulong Function
+               OpStore %11 %6
+               OpStore %12 %7
+         %16 = OpLoad %ulong %11
+         %25 = OpConvertUToPtr %_ptr_Generic_ulong %16
+         %15 = OpLoad %ulong %25
+               OpStore %13 %15
+               OpBranch %8
+          %8 = OpLabel
+         %18 = OpLoad %ulong %13
+         %17 = OpIAdd %ulong %18 %ulong_1
+               OpStore %14 %17
+               OpBranch %10
+         %30 = OpLabel
+         %20 = OpLoad %ulong %13
+         %19 = OpIAdd %ulong %20 %ulong_2
+               OpStore %14 %19
+               OpBranch %10
+         %10 = OpLabel
+         %21 = OpLoad %ulong %12
+         %22 = OpLoad %ulong %14
+         %26 = OpConvertUToPtr %_ptr_Generic_ulong %21
+               OpStore %26 %22
+               OpReturn
+               OpFunctionEnd
+               

ptx/src/test/spirv_run/mod.rs

@@ -45,6 +45,7 @@ test_ptx!(mul_lo, [1u64], [2u64]);
 test_ptx!(mul_hi, [u64::max_value()], [1u64]);
 test_ptx!(add, [1u64], [2u64]);
 test_ptx!(setp, [10u64, 11u64], [1u64, 0u64]);
+test_ptx!(bra, [10u64], [11u64]);
 
 struct DisplayError<T: Display + Debug> {
     err: T,

ptx/src/test/spirv_run/setp.ptx

@@ -0,0 +1,27 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry setp(
+	.param .u64 input,
+	.param .u64 output
+)
+{
+	.reg .u64 	    in_addr;
+    .reg .u64 	    out_addr;
+    .reg .u64 	    temp;
+    .reg .u64 	    temp2;
+    .reg .u64 	    temp3;
+    .reg .pred 	    pred;
+
+	ld.param.u64 	in_addr, [input];
+    ld.param.u64 	out_addr, [output];
+
+    ld.u64          temp, [in_addr];
+    ld.u64          temp2, [in_addr + 8];
+	setp.lt.u64		pred, temp, temp2;
+    @pred mov.u64   temp3, 1;
+    @!pred mov.u64  temp3, 2;
+    st.u64          [out_addr], temp3;
+	ret;
+}

ptx/src/test/spirv_run/setp.spvtxt

@@ -0,0 +1,65 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability Int8
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %5 "setp"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+          %4 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+       %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+    %ulong_8 = OpConstant %ulong 8
+    %ulong_1 = OpConstant %ulong 1
+    %ulong_2 = OpConstant %ulong 2
+          %5 = OpFunction %void None %4
+          %6 = OpFunctionParameter %ulong
+          %7 = OpFunctionParameter %ulong
+         %38 = OpLabel
+          %8 = OpVariable %_ptr_Function_ulong Function
+          %9 = OpVariable %_ptr_Function_ulong Function
+         %10 = OpVariable %_ptr_Function_ulong Function
+         %11 = OpVariable %_ptr_Function_ulong Function
+         %12 = OpVariable %_ptr_Function_ulong Function
+         %13 = OpVariable %_ptr_Function_bool Function
+               OpStore %8 %6
+               OpStore %9 %7
+         %19 = OpLoad %ulong %8
+         %35 = OpConvertUToPtr %_ptr_Generic_ulong %19
+         %18 = OpLoad %ulong %35
+               OpStore %10 %18
+         %21 = OpLoad %ulong %8
+         %32 = OpIAdd %ulong %21 %ulong_8
+         %36 = OpConvertUToPtr %_ptr_Generic_ulong %32
+         %20 = OpLoad %ulong %36
+               OpStore %11 %20
+         %23 = OpLoad %ulong %10
+         %24 = OpLoad %ulong %11
+         %22 = OpULessThan %bool %23 %24
+               OpStore %13 %22
+         %25 = OpLoad %bool %13
+               OpBranchConditional %25 %14 %15
+         %14 = OpLabel
+         %26 = OpCopyObject %ulong %ulong_1
+               OpStore %12 %26
+               OpBranch %15
+         %15 = OpLabel
+         %27 = OpLoad %bool %13
+               OpBranchConditional %27 %17 %16
+         %16 = OpLabel
+         %28 = OpCopyObject %ulong %ulong_2
+               OpStore %12 %28
+               OpBranch %17
+         %17 = OpLabel
+         %29 = OpLoad %ulong %9
+         %30 = OpLoad %ulong %12
+         %37 = OpConvertUToPtr %_ptr_Generic_ulong %29
+               OpStore %37 %30
+               OpReturn
+               OpFunctionEnd
+               

ptx/src/translate.rs

@@ -554,6 +554,14 @@ fn emit_function_body_ops(
                 }
                 builder.begin_block(Some(*id))?;
             }
+            _ => {
+                if builder.block.is_none() {
+                    builder.begin_block(None)?;
+                }
+            }
+        }
+        match s {
+            Statement::Label(_) => (),
             Statement::Variable(id, typ, ss) => {
                 let type_id = map.get_or_add(
                     builder,
@@ -858,6 +866,14 @@ fn normalize_identifiers<'a, 'b>(
     for arg in args {
         id_defs.add_def(arg.name, Some(ast::Type::Scalar(arg.a_type)));
     }
+    for s in func.iter() {
+        match s {
+            ast::Statement::Label(id) => {
+                id_defs.add_def(*id, None);
+            }
+            _ => (),
+        }
+    }
     let mut result = Vec::new();
     for s in func {
         expand_map_variables(&mut id_defs, &mut result, s);
@@ -872,7 +888,7 @@ fn expand_map_variables<'a>(
 ) {
     match s {
         ast::Statement::Label(name) => {
-            result.push(ast::Statement::Label(id_defs.add_def(name, None)))
+            result.push(ast::Statement::Label(id_defs.get_id(name)))
         }
         ast::Statement::Instruction(p, i) => result.push(ast::Statement::Instruction(
             p.map(|p| p.map_variable(&mut |id| id_defs.get_id(id))),