diff options
Diffstat (limited to 'ptx/src/test/spirv_run')
| -rw-r--r-- | ptx/src/test/spirv_run/brev.ptx | 21 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/brev.spvtxt | 47 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/clz.ptx | 21 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/clz.spvtxt | 47 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 9 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/popc.ptx | 21 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/popc.spvtxt | 47 |
7 files changed, 212 insertions, 1 deletions
diff --git a/ptx/src/test/spirv_run/brev.ptx b/ptx/src/test/spirv_run/brev.ptx new file mode 100644 index 0000000..1d9dd75 --- /dev/null +++ b/ptx/src/test/spirv_run/brev.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry brev(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .b32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.b32 temp, [in_addr];
+ brev.b32 temp, temp;
+ st.b32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/brev.spvtxt b/ptx/src/test/spirv_run/brev.spvtxt new file mode 100644 index 0000000..df5df53 --- /dev/null +++ b/ptx/src/test/spirv_run/brev.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "brev" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpBitReverse %uint %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/clz.ptx b/ptx/src/test/spirv_run/clz.ptx new file mode 100644 index 0000000..b475b90 --- /dev/null +++ b/ptx/src/test/spirv_run/clz.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry clz(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .b32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.b32 temp, [in_addr];
+ clz.b32 temp, temp;
+ st.b32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/clz.spvtxt b/ptx/src/test/spirv_run/clz.spvtxt new file mode 100644 index 0000000..5d1ebc8 --- /dev/null +++ b/ptx/src/test/spirv_run/clz.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "clz" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpExtInst %uint %21 clz %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 163caac..a7ef75b 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -104,11 +104,18 @@ test_ptx!(div_approx, [1f32, 2f32], [0.5f32]); test_ptx!(sqrt, [0.25f32], [0.5f32]);
test_ptx!(rsqrt, [0.25f64], [2f64]);
test_ptx!(neg, [181i32], [-181i32]);
-test_ptx!(sin, [std::f32::consts::PI/2f32], [1f32]);
+test_ptx!(sin, [std::f32::consts::PI / 2f32], [1f32]);
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!(clz, [0b00000101_00101101_00010011_10101011u32], [5u32]);
+test_ptx!(popc, [0b10111100_10010010_01001001_10001010u32], [14u32]);
+test_ptx!(
+ brev,
+ [0b11000111_01011100_10101110_11111011u32],
+ [0b11011111_01110101_00111010_11100011u32]
+);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/test/spirv_run/popc.ptx b/ptx/src/test/spirv_run/popc.ptx new file mode 100644 index 0000000..7106422 --- /dev/null +++ b/ptx/src/test/spirv_run/popc.ptx @@ -0,0 +1,21 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry popc(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .b32 temp;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.b32 temp, [in_addr];
+ popc.b32 temp, temp;
+ st.b32 [out_addr], temp;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/popc.spvtxt b/ptx/src/test/spirv_run/popc.spvtxt new file mode 100644 index 0000000..bb4968f --- /dev/null +++ b/ptx/src/test/spirv_run/popc.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "popc" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 + OpStore %4 %9 + %10 = OpLoad %ulong %3 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpBitCount %uint %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 + OpReturn + OpFunctionEnd |