diff --git a/ptx/src/test/spirv_run/bra.ptx b/ptx/src/test/spirv_run/bra.ptx new file mode 100644 index 0000000..fd5a0a3 --- /dev/null +++ b/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; +} diff --git a/ptx/src/test/spirv_run/bra.spvtxt b/ptx/src/test/spirv_run/bra.spvtxt new file mode 100644 index 0000000..81fedc5 --- /dev/null +++ b/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 + \ No newline at end of file diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index e0b7d74..c90e487 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/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 { err: T, diff --git a/ptx/src/test/spirv_run/setp.ptx b/ptx/src/test/spirv_run/setp.ptx new file mode 100644 index 0000000..8032505 --- /dev/null +++ b/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; +} diff --git a/ptx/src/test/spirv_run/setp.spvtxt b/ptx/src/test/spirv_run/setp.spvtxt new file mode 100644 index 0000000..22e7b54 --- /dev/null +++ b/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 + \ No newline at end of file diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs index 12b9aae..c40e554 100644 --- a/ptx/src/translate.rs +++ b/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))),