Fix typo in selp

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]);

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

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;
+}

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

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, _) => {