aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx/src/test
diff options
context:
space:
mode:
Diffstat (limited to 'ptx/src/test')
-rw-r--r--ptx/src/test/spirv_run/div_approx.ptx23
-rw-r--r--ptx/src/test/spirv_run/div_approx.spvtxt65
-rw-r--r--ptx/src/test/spirv_run/mod.rs4
-rw-r--r--ptx/src/test/spirv_run/rsqrt.ptx21
-rw-r--r--ptx/src/test/spirv_run/rsqrt.spvtxt56
-rw-r--r--ptx/src/test/spirv_run/shared_ptr_32.ptx29
-rw-r--r--ptx/src/test/spirv_run/shared_ptr_32.spvtxt74
-rw-r--r--ptx/src/test/spirv_run/sqrt.ptx21
-rw-r--r--ptx/src/test/spirv_run/sqrt.spvtxt56
9 files changed, 349 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/div_approx.ptx b/ptx/src/test/spirv_run/div_approx.ptx
new file mode 100644
index 0000000..b25e320
--- /dev/null
+++ b/ptx/src/test/spirv_run/div_approx.ptx
@@ -0,0 +1,23 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry div_approx(
+ .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];
+ div.approx.f32 temp1, temp1, temp2;
+ st.f32 [out_addr], temp1;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/div_approx.spvtxt b/ptx/src/test/spirv_run/div_approx.spvtxt
new file mode 100644
index 0000000..40cc152
--- /dev/null
+++ b/ptx/src/test/spirv_run/div_approx.spvtxt
@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: rspirv
+; Bound: 38
+OpCapability GenericPointer
+OpCapability Linkage
+OpCapability Addresses
+OpCapability Kernel
+OpCapability Int8
+OpCapability Int16
+OpCapability Int64
+OpCapability Float16
+OpCapability Float64
+; OpCapability FunctionFloatControlINTEL
+; OpExtension "SPV_INTEL_float_controls2"
+%30 = OpExtInstImport "OpenCL.std"
+OpMemoryModel Physical64 OpenCL
+OpEntryPoint Kernel %1 "div_approx"
+OpDecorate %1 FunctionDenormModeINTEL 32 Preserve
+OpDecorate %18 FPFastMathMode AllowRecip
+%31 = OpTypeVoid
+%32 = OpTypeInt 64 0
+%33 = OpTypeFunction %31 %32 %32
+%34 = OpTypePointer Function %32
+%35 = OpTypeFloat 32
+%36 = OpTypePointer Function %35
+%37 = OpTypePointer Generic %35
+%23 = OpConstant %32 4
+%1 = OpFunction %31 None %33
+%8 = OpFunctionParameter %32
+%9 = OpFunctionParameter %32
+%28 = OpLabel
+%2 = OpVariable %34 Function
+%3 = OpVariable %34 Function
+%4 = OpVariable %34 Function
+%5 = OpVariable %34 Function
+%6 = OpVariable %36 Function
+%7 = OpVariable %36 Function
+OpStore %2 %8
+OpStore %3 %9
+%11 = OpLoad %32 %2
+%10 = OpCopyObject %32 %11
+OpStore %4 %10
+%13 = OpLoad %32 %3
+%12 = OpCopyObject %32 %13
+OpStore %5 %12
+%15 = OpLoad %32 %4
+%25 = OpConvertUToPtr %37 %15
+%14 = OpLoad %35 %25
+OpStore %6 %14
+%17 = OpLoad %32 %4
+%24 = OpIAdd %32 %17 %23
+%26 = OpConvertUToPtr %37 %24
+%16 = OpLoad %35 %26
+OpStore %7 %16
+%19 = OpLoad %35 %6
+%20 = OpLoad %35 %7
+%18 = OpFDiv %35 %19 %20
+OpStore %6 %18
+%21 = OpLoad %32 %5
+%22 = OpLoad %35 %6
+%27 = OpConvertUToPtr %37 %21
+OpStore %27 %22
+OpReturn
+OpFunctionEnd \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 40a9d64..4e9d39f 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -97,9 +97,13 @@ test_ptx!(and, [6u32, 3u32], [2u32]);
test_ptx!(selp, [100u16, 200u16], [200u16]);
test_ptx!(fma, [2f32, 3f32, 5f32], [11f32]);
test_ptx!(shared_variable, [513u64], [513u64]);
+test_ptx!(shared_ptr_32, [513u64], [513u64]);
test_ptx!(atom_cas, [91u32, 91u32], [91u32, 100u32]);
test_ptx!(atom_inc, [100u32], [100u32, 101u32, 0u32]);
test_ptx!(atom_add, [2u32, 4u32], [2u32, 6u32]);
+test_ptx!(div_approx, [1f32, 2f32], [0.5f32]);
+test_ptx!(sqrt, [0.25f32], [0.5f32]);
+test_ptx!(rsqrt, [0.25f64], [2f64]);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/test/spirv_run/rsqrt.ptx b/ptx/src/test/spirv_run/rsqrt.ptx
new file mode 100644
index 0000000..5821501
--- /dev/null
+++ b/ptx/src/test/spirv_run/rsqrt.ptx
@@ -0,0 +1,21 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry rsqrt(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f64 temp1;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f64 temp1, [in_addr];
+ rsqrt.approx.f64 temp1, temp1;
+ st.f64 [out_addr], temp1;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/rsqrt.spvtxt b/ptx/src/test/spirv_run/rsqrt.spvtxt
new file mode 100644
index 0000000..5c3ba97
--- /dev/null
+++ b/ptx/src/test/spirv_run/rsqrt.spvtxt
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: rspirv
+; Bound: 31
+OpCapability GenericPointer
+OpCapability Linkage
+OpCapability Addresses
+OpCapability Kernel
+OpCapability Int8
+OpCapability Int16
+OpCapability Int64
+OpCapability Float16
+OpCapability Float64
+; OpCapability FunctionFloatControlINTEL
+; OpExtension "SPV_INTEL_float_controls2"
+%23 = OpExtInstImport "OpenCL.std"
+OpMemoryModel Physical64 OpenCL
+OpEntryPoint Kernel %1 "rsqrt"
+OpDecorate %1 FunctionDenormModeINTEL 64 Preserve
+%24 = OpTypeVoid
+%25 = OpTypeInt 64 0
+%26 = OpTypeFunction %24 %25 %25
+%27 = OpTypePointer Function %25
+%28 = OpTypeFloat 64
+%29 = OpTypePointer Function %28
+%30 = OpTypePointer Generic %28
+%1 = OpFunction %24 None %26
+%7 = OpFunctionParameter %25
+%8 = OpFunctionParameter %25
+%21 = OpLabel
+%2 = OpVariable %27 Function
+%3 = OpVariable %27 Function
+%4 = OpVariable %27 Function
+%5 = OpVariable %27 Function
+%6 = OpVariable %29 Function
+OpStore %2 %7
+OpStore %3 %8
+%10 = OpLoad %25 %2
+%9 = OpCopyObject %25 %10
+OpStore %4 %9
+%12 = OpLoad %25 %3
+%11 = OpCopyObject %25 %12
+OpStore %5 %11
+%14 = OpLoad %25 %4
+%19 = OpConvertUToPtr %30 %14
+%13 = OpLoad %28 %19
+OpStore %6 %13
+%16 = OpLoad %28 %6
+%15 = OpExtInst %28 %23 native_rsqrt %16
+OpStore %6 %15
+%17 = OpLoad %25 %5
+%18 = OpLoad %28 %6
+%20 = OpConvertUToPtr %30 %17
+OpStore %20 %18
+OpReturn
+OpFunctionEnd \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/shared_ptr_32.ptx b/ptx/src/test/spirv_run/shared_ptr_32.ptx
new file mode 100644
index 0000000..0334aa0
--- /dev/null
+++ b/ptx/src/test/spirv_run/shared_ptr_32.ptx
@@ -0,0 +1,29 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+
+.visible .entry shared_ptr_32(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .shared .align 4 .b8 shared_mem1[128];
+
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .u32 shared_addr;
+
+ .reg .u64 temp1;
+ .reg .u64 temp2;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+ mov.u32 shared_addr, shared_mem1;
+
+ ld.global.u64 temp1, [in_addr];
+ st.shared.u64 [shared_addr], temp1;
+ ld.shared.u64 temp2, [shared_addr+0];
+ st.global.u64 [out_addr], temp2;
+ ret;
+} \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/shared_ptr_32.spvtxt b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
new file mode 100644
index 0000000..609cc0e
--- /dev/null
+++ b/ptx/src/test/spirv_run/shared_ptr_32.spvtxt
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: rspirv
+; Bound: 47
+OpCapability GenericPointer
+OpCapability Linkage
+OpCapability Addresses
+OpCapability Kernel
+OpCapability Int8
+OpCapability Int16
+OpCapability Int64
+OpCapability Float16
+OpCapability Float64
+; OpCapability FunctionFloatControlINTEL
+; OpExtension "SPV_INTEL_float_controls2"
+%34 = OpExtInstImport "OpenCL.std"
+OpMemoryModel Physical64 OpenCL
+OpEntryPoint Kernel %1 "shared_ptr_32" %4
+OpDecorate %4 Alignment 4
+%35 = OpTypeVoid
+%36 = OpTypeInt 32 0
+%37 = OpTypeInt 8 0
+%38 = OpConstant %36 128
+%39 = OpTypeArray %37 %38
+%40 = OpTypePointer Workgroup %39
+%4 = OpVariable %40 Workgroup
+%41 = OpTypeInt 64 0
+%42 = OpTypeFunction %35 %41 %41
+%43 = OpTypePointer Function %41
+%44 = OpTypePointer Function %36
+%45 = OpTypePointer CrossWorkgroup %41
+%46 = OpTypePointer Workgroup %41
+%25 = OpConstant %36 0
+%1 = OpFunction %35 None %42
+%10 = OpFunctionParameter %41
+%11 = OpFunctionParameter %41
+%32 = OpLabel
+%2 = OpVariable %43 Function
+%3 = OpVariable %43 Function
+%5 = OpVariable %43 Function
+%6 = OpVariable %43 Function
+%7 = OpVariable %44 Function
+%8 = OpVariable %43 Function
+%9 = OpVariable %43 Function
+OpStore %2 %10
+OpStore %3 %11
+%13 = OpLoad %41 %2
+%12 = OpCopyObject %41 %13
+OpStore %5 %12
+%15 = OpLoad %41 %3
+%14 = OpCopyObject %41 %15
+OpStore %6 %14
+%27 = OpConvertPtrToU %36 %4
+%16 = OpCopyObject %36 %27
+OpStore %7 %16
+%18 = OpLoad %41 %5
+%28 = OpConvertUToPtr %45 %18
+%17 = OpLoad %41 %28
+OpStore %8 %17
+%19 = OpLoad %36 %7
+%20 = OpLoad %41 %8
+%29 = OpConvertUToPtr %46 %19
+OpStore %29 %20
+%22 = OpLoad %36 %7
+%26 = OpIAdd %36 %22 %25
+%30 = OpConvertUToPtr %46 %26
+%21 = OpLoad %41 %30
+OpStore %9 %21
+%23 = OpLoad %41 %6
+%24 = OpLoad %41 %9
+%31 = OpConvertUToPtr %45 %23
+OpStore %31 %24
+OpReturn
+OpFunctionEnd \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/sqrt.ptx b/ptx/src/test/spirv_run/sqrt.ptx
new file mode 100644
index 0000000..8b42f34
--- /dev/null
+++ b/ptx/src/test/spirv_run/sqrt.ptx
@@ -0,0 +1,21 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry sqrt(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp1;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.f32 temp1, [in_addr];
+ sqrt.approx.f32 temp1, temp1;
+ st.f32 [out_addr], temp1;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/sqrt.spvtxt b/ptx/src/test/spirv_run/sqrt.spvtxt
new file mode 100644
index 0000000..d2c5b20
--- /dev/null
+++ b/ptx/src/test/spirv_run/sqrt.spvtxt
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: rspirv
+; Bound: 31
+OpCapability GenericPointer
+OpCapability Linkage
+OpCapability Addresses
+OpCapability Kernel
+OpCapability Int8
+OpCapability Int16
+OpCapability Int64
+OpCapability Float16
+OpCapability Float64
+; OpCapability FunctionFloatControlINTEL
+; OpExtension "SPV_INTEL_float_controls2"
+%23 = OpExtInstImport "OpenCL.std"
+OpMemoryModel Physical64 OpenCL
+OpEntryPoint Kernel %1 "sqrt"
+OpDecorate %1 FunctionDenormModeINTEL 32 Preserve
+%24 = OpTypeVoid
+%25 = OpTypeInt 64 0
+%26 = OpTypeFunction %24 %25 %25
+%27 = OpTypePointer Function %25
+%28 = OpTypeFloat 32
+%29 = OpTypePointer Function %28
+%30 = OpTypePointer Generic %28
+%1 = OpFunction %24 None %26
+%7 = OpFunctionParameter %25
+%8 = OpFunctionParameter %25
+%21 = OpLabel
+%2 = OpVariable %27 Function
+%3 = OpVariable %27 Function
+%4 = OpVariable %27 Function
+%5 = OpVariable %27 Function
+%6 = OpVariable %29 Function
+OpStore %2 %7
+OpStore %3 %8
+%10 = OpLoad %25 %2
+%9 = OpCopyObject %25 %10
+OpStore %4 %9
+%12 = OpLoad %25 %3
+%11 = OpCopyObject %25 %12
+OpStore %5 %11
+%14 = OpLoad %25 %4
+%19 = OpConvertUToPtr %30 %14
+%13 = OpLoad %28 %19
+OpStore %6 %13
+%16 = OpLoad %28 %6
+%15 = OpExtInst %28 %23 native_sqrt %16
+OpStore %6 %15
+%17 = OpLoad %25 %5
+%18 = OpLoad %28 %6
+%20 = OpConvertUToPtr %30 %17
+OpStore %20 %18
+OpReturn
+OpFunctionEnd \ No newline at end of file
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-26 01:02:47 +0000