aboutsummaryrefslogtreecommitdiffhomepage
diff options
context:
space:
mode:
Diffstat
-rw-r--r--ptx/src/test/spirv_run/cvt_rzi.ptx25
-rw-r--r--ptx/src/test/spirv_run/cvt_rzi.spvtxt63
-rw-r--r--ptx/src/test/spirv_run/mod.rs3
-rw-r--r--ptx/src/test/spirv_run/setp_gt.ptx27
-rw-r--r--ptx/src/test/spirv_run/setp_gt.spvtxt75
-rw-r--r--ptx/src/test/spirv_run/setp_leu.ptx27
-rw-r--r--ptx/src/test/spirv_run/setp_leu.spvtxt75
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
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-26 02:10:39 +0000