aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx/src/test
diff options
context:
space:
mode:
Diffstat (limited to 'ptx/src/test')
-rw-r--r--ptx/src/test/mod.rs7
-rw-r--r--ptx/src/test/spirv_run/mad_s32.ptx28
-rw-r--r--ptx/src/test/spirv_run/mad_s32.spvtxt77
-rw-r--r--ptx/src/test/spirv_run/mod.rs10
-rw-r--r--ptx/src/test/spirv_run/mul_wide.ptx24
-rw-r--r--ptx/src/test/spirv_run/mul_wide.spvtxt64
-rw-r--r--ptx/src/test/vectorAdd_11.ptx55
7 files changed, 261 insertions, 4 deletions
diff --git a/ptx/src/test/mod.rs b/ptx/src/test/mod.rs
index d251884..0339141 100644
--- a/ptx/src/test/mod.rs
+++ b/ptx/src/test/mod.rs
@@ -40,3 +40,10 @@ fn _Z9vectorAddPKfS0_Pfi_ptx() -> Result<(), TranslateError> {
let vector_add = include_str!("_Z9vectorAddPKfS0_Pfi.ptx");
compile_and_assert(vector_add)
}
+
+#[test]
+#[allow(non_snake_case)]
+fn vectorAdd_11_ptx() -> Result<(), TranslateError> {
+ let vector_add = include_str!("vectorAdd_11.ptx");
+ compile_and_assert(vector_add)
+}
diff --git a/ptx/src/test/spirv_run/mad_s32.ptx b/ptx/src/test/spirv_run/mad_s32.ptx
new file mode 100644
index 0000000..a864266
--- /dev/null
+++ b/ptx/src/test/spirv_run/mad_s32.ptx
@@ -0,0 +1,28 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry mad_s32(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .s32 dst;
+ .reg .s32 src1;
+ .reg .s32 src2;
+ .reg .s32 src3;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.s32 src1, [in_addr];
+ ld.s32 src2, [in_addr+4];
+ ld.s32 src3, [in_addr+8];
+ mad.lo.s32 dst, src1, src2, src3;
+ st.s32 [out_addr], dst;
+ st.s32 [out_addr+4], dst;
+ st.s32 [out_addr+8], dst;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/mad_s32.spvtxt b/ptx/src/test/spirv_run/mad_s32.spvtxt
new file mode 100644
index 0000000..3a7153d
--- /dev/null
+++ b/ptx/src/test/spirv_run/mad_s32.spvtxt
@@ -0,0 +1,77 @@
+ OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int64
+ OpCapability Int8
+ OpCapability Float64
+ %48 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %1 "mad_s32"
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+ %51 = 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
+ %ulong_4 = OpConstant %ulong 4
+ %ulong_8 = OpConstant %ulong 8
+ %ulong_4_0 = OpConstant %ulong 4
+ %ulong_8_0 = OpConstant %ulong 8
+ %1 = OpFunction %void None %51
+ %10 = OpFunctionParameter %ulong
+ %11 = OpFunctionParameter %ulong
+ %46 = 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
+ %7 = OpVariable %_ptr_Function_uint Function
+ %8 = OpVariable %_ptr_Function_uint Function
+ %9 = OpVariable %_ptr_Function_uint Function
+ OpStore %2 %10
+ OpStore %3 %11
+ %13 = OpLoad %ulong %2
+ %12 = OpCopyObject %ulong %13
+ OpStore %4 %12
+ %15 = OpLoad %ulong %3
+ %14 = OpCopyObject %ulong %15
+ OpStore %5 %14
+ %17 = OpLoad %ulong %4
+ %40 = OpConvertUToPtr %_ptr_Generic_uint %17
+ %16 = OpLoad %uint %40
+ OpStore %7 %16
+ %19 = OpLoad %ulong %4
+ %33 = OpIAdd %ulong %19 %ulong_4
+ %41 = OpConvertUToPtr %_ptr_Generic_uint %33
+ %18 = OpLoad %uint %41
+ OpStore %8 %18
+ %21 = OpLoad %ulong %4
+ %35 = OpIAdd %ulong %21 %ulong_8
+ %42 = OpConvertUToPtr %_ptr_Generic_uint %35
+ %20 = OpLoad %uint %42
+ OpStore %9 %20
+ %23 = OpLoad %uint %7
+ %24 = OpLoad %uint %8
+ %25 = OpLoad %uint %9
+ %56 = OpIMul %uint %23 %24
+ %22 = OpIAdd %uint %25 %56
+ OpStore %6 %22
+ %26 = OpLoad %ulong %5
+ %27 = OpLoad %uint %6
+ %43 = OpConvertUToPtr %_ptr_Generic_uint %26
+ OpStore %43 %27
+ %28 = OpLoad %ulong %5
+ %29 = OpLoad %uint %6
+ %37 = OpIAdd %ulong %28 %ulong_4_0
+ %44 = OpConvertUToPtr %_ptr_Generic_uint %37
+ OpStore %44 %29
+ %30 = OpLoad %ulong %5
+ %31 = OpLoad %uint %6
+ %39 = OpIAdd %ulong %30 %ulong_8_0
+ %45 = OpConvertUToPtr %_ptr_Generic_uint %39
+ OpStore %45 %31
+ OpReturn
+ OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 78c3375..27dc063 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -8,7 +8,6 @@ use spirv_headers::Word;
use spirv_tools_sys::{
spv_binary, spv_endianness_t, spv_parsed_instruction_t, spv_result_t, spv_target_env,
};
-use std::{collections::hash_map::Entry, cmp};
use std::error;
use std::ffi::{c_void, CStr, CString};
use std::fmt;
@@ -17,6 +16,7 @@ use std::hash::Hash;
use std::mem;
use std::slice;
use std::{borrow::Cow, collections::HashMap, env, fs, path::PathBuf, ptr, str};
+use std::{cmp, collections::hash_map::Entry};
macro_rules! test_ptx {
($fn_name:ident, $input:expr, $output:expr) => {
@@ -65,6 +65,8 @@ test_ptx!(mov_address, [0xDEADu64], [0u64]);
test_ptx!(b64tof64, [111u64], [111u64]);
test_ptx!(implicit_param, [34u32], [34u32]);
test_ptx!(pred_not, [10u64, 11u64], [2u64, 0u64]);
+test_ptx!(mad_s32, [2i32, 3i32, 4i32], [10i32, 10i32, 10i32]);
+test_ptx!(mul_wide, [0x01_00_00_00__01_00_00_00i64], [0x1_00_00_00_00_00_00i64]);
struct DisplayError<T: Debug> {
err: T,
@@ -93,7 +95,7 @@ fn test_ptx_assert<'a, T: From<u8> + ze::SafeRepr + Debug + Copy + PartialEq>(
let mut errors = Vec::new();
let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_text)?;
assert!(errors.len() == 0);
- let spirv = translate::to_spirv(ast)?;
+ let (spirv, _) = translate::to_spirv(ast)?;
let name = CString::new(name)?;
let result =
run_spirv(name.as_c_str(), &spirv, input, output).map_err(|err| DisplayError { err })?;
@@ -127,7 +129,7 @@ fn run_spirv<T: From<u8> + ze::SafeRepr + Copy + Debug>(
kernel.set_indirect_access(
ze::sys::ze_kernel_indirect_access_flags_t::ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE,
)?;
- let mut inp_b = ze::DeviceBuffer::<T>::new(&mut ctx, &dev, cmp::max(input.len(),1))?;
+ let mut inp_b = ze::DeviceBuffer::<T>::new(&mut ctx, &dev, cmp::max(input.len(), 1))?;
let mut out_b = ze::DeviceBuffer::<T>::new(&mut ctx, &dev, cmp::max(output.len(), 1))?;
let inp_b_ptr_mut: ze::BufferPtrMut<T> = (&mut inp_b).into();
let event_pool = ze::EventPool::new(&mut ctx, 3, Some(&[&dev]))?;
@@ -157,7 +159,7 @@ fn test_spvtxt_assert<'a>(
let mut errors = Vec::new();
let ast = ptx::ModuleParser::new().parse(&mut errors, ptx_txt)?;
assert!(errors.len() == 0);
- let ptx_mod = translate::to_spirv_module(ast)?;
+ let (ptx_mod, _) = translate::to_spirv_module(ast)?;
let spv_context =
unsafe { spirv_tools::spvContextCreate(spv_target_env::SPV_ENV_UNIVERSAL_1_3) };
assert!(spv_context != ptr::null_mut());
diff --git a/ptx/src/test/spirv_run/mul_wide.ptx b/ptx/src/test/spirv_run/mul_wide.ptx
new file mode 100644
index 0000000..2d6f8a5
--- /dev/null
+++ b/ptx/src/test/spirv_run/mul_wide.ptx
@@ -0,0 +1,24 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry mul_wide(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .s32 inp1;
+ .reg .s32 inp2;
+ .reg .s64 result;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.global.s32 inp1, [in_addr];
+ ld.global.s32 inp2, [in_addr+4];
+ mul.wide.s32 result, inp1, inp2;
+ st.u64 [out_addr], result;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/mul_wide.spvtxt b/ptx/src/test/spirv_run/mul_wide.spvtxt
new file mode 100644
index 0000000..274612c
--- /dev/null
+++ b/ptx/src/test/spirv_run/mul_wide.spvtxt
@@ -0,0 +1,64 @@
+ OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int64
+ OpCapability Int8
+ OpCapability Float64
+ %32 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %1 "mul_wide"
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+ %35 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+ %uint = OpTypeInt 32 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
+ %ulong_4 = OpConstant %ulong 4
+ %_struct_40 = OpTypeStruct %uint %uint
+ %v2uint = OpTypeVector %uint 2
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+ %1 = OpFunction %void None %35
+ %9 = OpFunctionParameter %ulong
+ %10 = OpFunctionParameter %ulong
+ %30 = 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
+ %7 = OpVariable %_ptr_Function_uint Function
+ %8 = OpVariable %_ptr_Function_ulong Function
+ OpStore %2 %9
+ OpStore %3 %10
+ %12 = OpLoad %ulong %2
+ %11 = OpCopyObject %ulong %12
+ OpStore %4 %11
+ %14 = OpLoad %ulong %3
+ %13 = OpCopyObject %ulong %14
+ OpStore %5 %13
+ %16 = OpLoad %ulong %4
+ %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %16
+ %15 = OpLoad %uint %26
+ OpStore %6 %15
+ %18 = OpLoad %ulong %4
+ %25 = OpIAdd %ulong %18 %ulong_4
+ %27 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %25
+ %17 = OpLoad %uint %27
+ OpStore %7 %17
+ %20 = OpLoad %uint %6
+ %21 = OpLoad %uint %7
+ %41 = OpSMulExtended %_struct_40 %20 %21
+ %42 = OpCompositeExtract %uint %41 0
+ %43 = OpCompositeExtract %uint %41 1
+ %45 = OpCompositeConstruct %v2uint %42 %43
+ %19 = OpBitcast %ulong %45
+ OpStore %8 %19
+ %22 = OpLoad %ulong %5
+ %23 = OpLoad %ulong %8
+ %28 = OpCopyObject %ulong %23
+ %29 = OpConvertUToPtr %_ptr_Generic_ulong %22
+ OpStore %29 %28
+ OpReturn
+ OpFunctionEnd
diff --git a/ptx/src/test/vectorAdd_11.ptx b/ptx/src/test/vectorAdd_11.ptx
new file mode 100644
index 0000000..ba0381e
--- /dev/null
+++ b/ptx/src/test/vectorAdd_11.ptx
@@ -0,0 +1,55 @@
+
+
+
+
+
+
+
+
+.version 7.0
+.target sm_80
+.address_size 64
+
+
+
+.visible .entry _Z9vectorAddPKfS0_Pfi(
+.param .u64 _Z9vectorAddPKfS0_Pfi_param_0,
+.param .u64 _Z9vectorAddPKfS0_Pfi_param_1,
+.param .u64 _Z9vectorAddPKfS0_Pfi_param_2,
+.param .u32 _Z9vectorAddPKfS0_Pfi_param_3
+)
+{
+.reg .pred %p<2>;
+.reg .f32 %f<4>;
+.reg .b32 %r<6>;
+.reg .b64 %rd<11>;
+
+
+ld.param.u64 %rd1, [_Z9vectorAddPKfS0_Pfi_param_0];
+ld.param.u64 %rd2, [_Z9vectorAddPKfS0_Pfi_param_1];
+ld.param.u64 %rd3, [_Z9vectorAddPKfS0_Pfi_param_2];
+ld.param.u32 %r2, [_Z9vectorAddPKfS0_Pfi_param_3];
+mov.u32 %r3, %ntid.x;
+mov.u32 %r4, %ctaid.x;
+mov.u32 %r5, %tid.x;
+mad.lo.s32 %r1, %r4, %r3, %r5;
+setp.ge.s32 %p1, %r1, %r2;
+@%p1 bra BB0_2;
+
+cvta.to.global.u64 %rd4, %rd1;
+mul.wide.s32 %rd5, %r1, 4;
+add.s64 %rd6, %rd4, %rd5;
+cvta.to.global.u64 %rd7, %rd2;
+add.s64 %rd8, %rd7, %rd5;
+ld.global.f32 %f1, [%rd8];
+ld.global.f32 %f2, [%rd6];
+add.f32 %f3, %f2, %f1;
+cvta.to.global.u64 %rd9, %rd3;
+add.s64 %rd10, %rd9, %rd5;
+st.global.f32 [%rd10], %f3;
+
+BB0_2:
+ret;
+}
+
+
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-26 12:03:04 +0000