Fix typo in selp

5 files changed, 83 insertions, 2 deletions

diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index d6a8038..3148120 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -98,6 +98,7 @@ test_ptx!(constant_f32, [10f32], [5f32]);
 test_ptx!(constant_negative, [-101i32], [101i32]);

 test_ptx!(and, [6u32, 3u32], [2u32]);

 test_ptx!(selp, [100u16, 200u16], [200u16]);

+test_ptx!(selp_true, [100u16, 200u16], [100u16]);

 test_ptx!(fma, [2f32, 3f32, 5f32], [11f32]);

 test_ptx!(shared_variable, [513u64], [513u64]);

 test_ptx!(shared_ptr_32, [513u64], [513u64]);

diff --git a/ptx/src/test/spirv_run/selp.spvtxt b/ptx/src/test/spirv_run/selp.spvtxt
index 44e87e4..84976ac 100644
--- a/ptx/src/test/spirv_run/selp.spvtxt
+++ b/ptx/src/test/spirv_run/selp.spvtxt
@@ -47,7 +47,7 @@
                OpStore %7 %14
          %17 = OpLoad %ushort %6
          %18 = OpLoad %ushort %7
-         %16 = OpSelect %ushort %false %18 %18
+         %16 = OpSelect %ushort %false %17 %18
                OpStore %6 %16
          %19 = OpLoad %ulong %5
          %20 = OpLoad %ushort %6
diff --git a/ptx/src/test/spirv_run/selp_true.ptx b/ptx/src/test/spirv_run/selp_true.ptx
new file mode 100644
index 0000000..7abad81
--- /dev/null
+++ b/ptx/src/test/spirv_run/selp_true.ptx
@@ -0,0 +1,23 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry selp_true(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .u16 	    temp1;

+    .reg .u16 	    temp2;

+

+	ld.param.u64 	in_addr, [input];

+    ld.param.u64 	out_addr, [output];

+

+    ld.u16          temp1, [in_addr];

+    ld.u16          temp2, [in_addr + 2];

+    selp.u16   		temp1, temp1, temp2, 1;

+    st.u16          [out_addr], temp1;

+	ret;

+}

diff --git a/ptx/src/test/spirv_run/selp_true.spvtxt b/ptx/src/test/spirv_run/selp_true.spvtxt
new file mode 100644
index 0000000..1880de1
--- /dev/null
+++ b/ptx/src/test/spirv_run/selp_true.spvtxt
@@ -0,0 +1,57 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %29 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "selp_true"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+         %32 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+     %ushort = OpTypeInt 16 0
+%_ptr_Function_ushort = OpTypePointer Function %ushort
+%_ptr_Generic_ushort = OpTypePointer Generic %ushort
+    %ulong_2 = OpConstant %ulong 2
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+          %1 = OpFunction %void None %32
+          %8 = OpFunctionParameter %ulong
+          %9 = OpFunctionParameter %ulong
+         %27 = OpLabel
+          %2 = OpVariable %_ptr_Function_ulong Function
+          %3 = OpVariable %_ptr_Function_ulong Function
+          %4 = OpVariable %_ptr_Function_ulong Function
+          %5 = OpVariable %_ptr_Function_ulong Function
+          %6 = OpVariable %_ptr_Function_ushort Function
+          %7 = OpVariable %_ptr_Function_ushort Function
+               OpStore %2 %8
+               OpStore %3 %9
+         %10 = OpLoad %ulong %2
+               OpStore %4 %10
+         %11 = OpLoad %ulong %3
+               OpStore %5 %11
+         %13 = OpLoad %ulong %4
+         %24 = OpConvertUToPtr %_ptr_Generic_ushort %13
+         %12 = OpLoad %ushort %24
+               OpStore %6 %12
+         %15 = OpLoad %ulong %4
+         %22 = OpIAdd %ulong %15 %ulong_2
+         %25 = OpConvertUToPtr %_ptr_Generic_ushort %22
+         %14 = OpLoad %ushort %25
+               OpStore %7 %14
+         %17 = OpLoad %ushort %6
+         %18 = OpLoad %ushort %7
+         %16 = OpSelect %ushort %true %17 %18
+               OpStore %6 %16
+         %19 = OpLoad %ulong %5
+         %20 = OpLoad %ushort %6
+         %26 = OpConvertUToPtr %_ptr_Generic_ushort %19
+               OpStore %26 %20
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index ed752e9..3133fdc 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -3025,7 +3025,7 @@ fn emit_function_body_ops(
                 }

                 ast::Instruction::Selp(t, a) => {

                     let result_type = map.get_or_add_scalar(builder, ast::ScalarType::from(*t));

-                    builder.select(result_type, Some(a.dst), a.src3, a.src2, a.src2)?;

+                    builder.select(result_type, Some(a.dst), a.src3, a.src1, a.src2)?;

                 }

                 // TODO: implement named barriers

                 ast::Instruction::Bar(d, _) => {