Add more tests

7 files changed, 295 insertions, 0 deletions

diff --git a/ptx/src/test/spirv_run/cvt_rzi.ptx b/ptx/src/test/spirv_run/cvt_rzi.ptx
new file mode 100644
index 0000000..ba5cc0e
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_rzi.ptx
@@ -0,0 +1,25 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry cvt_rzi(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .f32 	    temp1;

+    .reg .f32 	    temp2;

+

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

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

+

+    ld.f32          temp1, [in_addr];

+    ld.f32          temp2, [in_addr+4];

+    cvt.rzi.f32.f32 temp1, temp1;

+    cvt.rzi.f32.f32 temp2, temp2;

+    st.f32          [out_addr], temp1;

+    st.f32          [out_addr+4], temp2;

+	ret;

+}

diff --git a/ptx/src/test/spirv_run/cvt_rzi.spvtxt b/ptx/src/test/spirv_run/cvt_rzi.spvtxt
new file mode 100644
index 0000000..68c12c6
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_rzi.spvtxt
@@ -0,0 +1,63 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %34 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "cvt_rzi"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+         %37 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+      %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_Generic_float = OpTypePointer Generic %float
+    %ulong_4 = OpConstant %ulong 4
+  %ulong_4_0 = OpConstant %ulong 4
+          %1 = OpFunction %void None %37
+          %8 = OpFunctionParameter %ulong
+          %9 = OpFunctionParameter %ulong
+         %32 = 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_float Function
+          %7 = OpVariable %_ptr_Function_float Function
+               OpStore %2 %8
+               OpStore %3 %9
+         %10 = OpLoad %ulong %2 Aligned 8
+               OpStore %4 %10
+         %11 = OpLoad %ulong %3 Aligned 8
+               OpStore %5 %11
+         %13 = OpLoad %ulong %4
+         %28 = OpConvertUToPtr %_ptr_Generic_float %13
+         %12 = OpLoad %float %28 Aligned 4
+               OpStore %6 %12
+         %15 = OpLoad %ulong %4
+         %25 = OpIAdd %ulong %15 %ulong_4
+         %29 = OpConvertUToPtr %_ptr_Generic_float %25
+         %14 = OpLoad %float %29 Aligned 4
+               OpStore %7 %14
+         %17 = OpLoad %float %6
+         %16 = OpExtInst %float %34 trunc %17
+               OpStore %6 %16
+         %19 = OpLoad %float %7
+         %18 = OpExtInst %float %34 trunc %19
+               OpStore %7 %18
+         %20 = OpLoad %ulong %5
+         %21 = OpLoad %float %6
+         %30 = OpConvertUToPtr %_ptr_Generic_float %20
+               OpStore %30 %21 Aligned 4
+         %22 = OpLoad %ulong %5
+         %23 = OpLoad %float %7
+         %27 = OpIAdd %ulong %22 %ulong_4_0
+         %31 = OpConvertUToPtr %_ptr_Generic_float %27
+               OpStore %31 %23 Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 86f9c16..5435b5f 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -49,6 +49,8 @@ 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!(setp_gt, [f32::NAN, 1f32], [1f32]);

+test_ptx!(setp_leu, [1f32, f32::NAN], [1f32]);

 test_ptx!(bra, [10u64], [11u64]);

 test_ptx!(not, [0u64], [u64::max_value()]);

 test_ptx!(shl, [11u64], [44u64]);

@@ -114,6 +116,7 @@ test_ptx!(cos, [std::f32::consts::PI], [-1f32]);
 test_ptx!(lg2, [512f32], [9f32]);

 test_ptx!(ex2, [10f32], [1024f32]);

 test_ptx!(cvt_rni, [9.5f32, 10.5f32], [10f32, 10f32]);

+test_ptx!(cvt_rzi, [-13.8f32, 12.9f32], [-13f32, 12f32]);

 test_ptx!(clz, [0b00000101_00101101_00010011_10101011u32], [5u32]);

 test_ptx!(popc, [0b10111100_10010010_01001001_10001010u32], [14u32]);

 test_ptx!(

diff --git a/ptx/src/test/spirv_run/setp_gt.ptx b/ptx/src/test/spirv_run/setp_gt.ptx
new file mode 100644
index 0000000..5f45300
--- /dev/null
+++ b/ptx/src/test/spirv_run/setp_gt.ptx
@@ -0,0 +1,27 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry setp_gt(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .f32 	    r1;

+    .reg .f32 	    r2;

+    .reg .f32 	    r3;

+    .reg .pred 	    pred;

+

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

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

+

+    ld.f32              r1, [in_addr];

+    ld.f32              r2, [in_addr + 4];

+    setp.gt.ftz.f32     pred, r1, r2;

+    @pred mov.f32       r3, r1;

+    @!pred mov.f32      r3, r2;

+    st.f32              [out_addr], r3;

+    ret;

+}

diff --git a/ptx/src/test/spirv_run/setp_gt.spvtxt b/ptx/src/test/spirv_run/setp_gt.spvtxt
new file mode 100644
index 0000000..77f6546
--- /dev/null
+++ b/ptx/src/test/spirv_run/setp_gt.spvtxt
@@ -0,0 +1,75 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %40 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "setp_gt"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+         %43 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+      %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+       %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+%_ptr_Generic_float = OpTypePointer Generic %float
+    %ulong_4 = OpConstant %ulong 4
+          %1 = OpFunction %void None %43
+         %14 = OpFunctionParameter %ulong
+         %15 = OpFunctionParameter %ulong
+         %38 = 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_float Function
+          %7 = OpVariable %_ptr_Function_float Function
+          %8 = OpVariable %_ptr_Function_float Function
+          %9 = OpVariable %_ptr_Function_bool Function
+               OpStore %2 %14
+               OpStore %3 %15
+         %16 = OpLoad %ulong %2 Aligned 8
+               OpStore %4 %16
+         %17 = OpLoad %ulong %3 Aligned 8
+               OpStore %5 %17
+         %19 = OpLoad %ulong %4
+         %35 = OpConvertUToPtr %_ptr_Generic_float %19
+         %18 = OpLoad %float %35 Aligned 4
+               OpStore %6 %18
+         %21 = OpLoad %ulong %4
+         %34 = OpIAdd %ulong %21 %ulong_4
+         %36 = OpConvertUToPtr %_ptr_Generic_float %34
+         %20 = OpLoad %float %36 Aligned 4
+               OpStore %7 %20
+         %23 = OpLoad %float %6
+         %24 = OpLoad %float %7
+         %22 = OpFOrdGreaterThan %bool %23 %24
+               OpStore %9 %22
+         %25 = OpLoad %bool %9
+               OpBranchConditional %25 %10 %11
+         %10 = OpLabel
+         %27 = OpLoad %float %6
+         %26 = OpCopyObject %float %27
+               OpStore %8 %26
+               OpBranch %11
+         %11 = OpLabel
+         %28 = OpLoad %bool %9
+               OpBranchConditional %28 %13 %12
+         %12 = OpLabel
+         %30 = OpLoad %float %7
+         %29 = OpCopyObject %float %30
+               OpStore %8 %29
+               OpBranch %13
+         %13 = OpLabel
+         %31 = OpLoad %ulong %5
+         %32 = OpLoad %float %8
+         %37 = OpConvertUToPtr %_ptr_Generic_float %31
+               OpStore %37 %32 Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/setp_leu.ptx b/ptx/src/test/spirv_run/setp_leu.ptx
new file mode 100644
index 0000000..be7538a
--- /dev/null
+++ b/ptx/src/test/spirv_run/setp_leu.ptx
@@ -0,0 +1,27 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry setp_leu(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .f32 	    r1;

+    .reg .f32 	    r2;

+    .reg .f32 	    r3;

+    .reg .pred 	    pred;

+

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

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

+

+    ld.f32              r1, [in_addr];

+    ld.f32              r2, [in_addr + 4];

+	setp.leu.ftz.f32    pred, r1, r2;

+    @pred mov.f32       r3, r1;

+    @!pred mov.f32      r3, r2;

+    st.f32              [out_addr], r3;

+	ret;

+}

diff --git a/ptx/src/test/spirv_run/setp_leu.spvtxt b/ptx/src/test/spirv_run/setp_leu.spvtxt
new file mode 100644
index 0000000..f80880a
--- /dev/null
+++ b/ptx/src/test/spirv_run/setp_leu.spvtxt
@@ -0,0 +1,75 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %40 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "setp_leu"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+         %43 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+      %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+       %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+%_ptr_Generic_float = OpTypePointer Generic %float
+    %ulong_4 = OpConstant %ulong 4
+          %1 = OpFunction %void None %43
+         %14 = OpFunctionParameter %ulong
+         %15 = OpFunctionParameter %ulong
+         %38 = 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_float Function
+          %7 = OpVariable %_ptr_Function_float Function
+          %8 = OpVariable %_ptr_Function_float Function
+          %9 = OpVariable %_ptr_Function_bool Function
+               OpStore %2 %14
+               OpStore %3 %15
+         %16 = OpLoad %ulong %2 Aligned 8
+               OpStore %4 %16
+         %17 = OpLoad %ulong %3 Aligned 8
+               OpStore %5 %17
+         %19 = OpLoad %ulong %4
+         %35 = OpConvertUToPtr %_ptr_Generic_float %19
+         %18 = OpLoad %float %35 Aligned 4
+               OpStore %6 %18
+         %21 = OpLoad %ulong %4
+         %34 = OpIAdd %ulong %21 %ulong_4
+         %36 = OpConvertUToPtr %_ptr_Generic_float %34
+         %20 = OpLoad %float %36 Aligned 4
+               OpStore %7 %20
+         %23 = OpLoad %float %6
+         %24 = OpLoad %float %7
+         %22 = OpFUnordLessThanEqual %bool %23 %24
+               OpStore %9 %22
+         %25 = OpLoad %bool %9
+               OpBranchConditional %25 %10 %11
+         %10 = OpLabel
+         %27 = OpLoad %float %6
+         %26 = OpCopyObject %float %27
+               OpStore %8 %26
+               OpBranch %11
+         %11 = OpLabel
+         %28 = OpLoad %bool %9
+               OpBranchConditional %28 %13 %12
+         %12 = OpLabel
+         %30 = OpLoad %float %7
+         %29 = OpCopyObject %float %30
+               OpStore %8 %29
+               OpBranch %13
+         %13 = OpLabel
+         %31 = OpLoad %ulong %5
+         %32 = OpLoad %float %8
+         %37 = OpConvertUToPtr %_ptr_Generic_float %31
+               OpStore %37 %32 Aligned 4
+               OpReturn
+               OpFunctionEnd