diff options
Diffstat (limited to 'ptx/src/test')
| -rw-r--r-- | ptx/src/test/spirv_run/local_align.spvtxt | 8 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 11 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mov_address.ptx | 15 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mov_address.spvtxt (renamed from ptx/src/test/spirv_run/reg_slm.spvtxt) | 0 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/reg_local.ptx (renamed from ptx/src/test/spirv_run/reg_slm.ptx) | 10 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/reg_local.spvtxt | 46 |
6 files changed, 76 insertions, 14 deletions
diff --git a/ptx/src/test/spirv_run/local_align.spvtxt b/ptx/src/test/spirv_run/local_align.spvtxt index 09a3f92..2482a75 100644 --- a/ptx/src/test/spirv_run/local_align.spvtxt +++ b/ptx/src/test/spirv_run/local_align.spvtxt @@ -13,8 +13,10 @@ %25 = OpTypeFunction %void %ulong %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong %uchar = OpTypeInt 8 0 -%_arr_uchar_8 = OpTypeArray %uchar %8 -%_ptr_Function__arr_uchar_8 = OpTypePointer Function %_arr_uchar_8 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 +%_arr_uchar_uint_8 = OpTypeArray %uchar %uint_8 +%_ptr_Function__arr_uchar_uint_8 = OpTypePointer Function %_arr_uchar_uint_8 %_ptr_Generic_ulong = OpTypePointer Generic %ulong %1 = OpFunction %void None %25 %8 = OpFunctionParameter %ulong @@ -22,7 +24,7 @@ %20 = OpLabel %2 = OpVariable %_ptr_Function_ulong Function %3 = OpVariable %_ptr_Function_ulong Function - %4 = OpVariable %_ptr_Function__arr_uchar_8 Workgroup + %4 = OpVariable %_ptr_Function__arr_uchar_uint_8 Function %5 = OpVariable %_ptr_Function_ulong Function %6 = OpVariable %_ptr_Function_ulong Function %7 = OpVariable %_ptr_Function_ulong Function diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 4a793c4..4c9d779 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -8,7 +8,7 @@ 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;
+use std::{collections::hash_map::Entry, cmp};
use std::error;
use std::ffi::{c_void, CStr, CString};
use std::fmt;
@@ -59,8 +59,9 @@ test_ptx!(local_align, [1u64], [1u64]); test_ptx!(call, [1u64], [2u64]);
test_ptx!(vector, [1u32, 2u32], [3u32, 3u32]);
test_ptx!(ld_st_offset, [1u32, 2u32], [2u32, 1u32]);
-//test_ptx!(ntid, [3u32], [4u32]);
-//test_ptx!(reg_slm, [12u64], [12u64]);
+test_ptx!(ntid, [3u32], [4u32]);
+test_ptx!(reg_local, [12u64], [12u64]);
+test_ptx!(mov_address, [0xDEADu64], [0u64]);
struct DisplayError<T: Debug> {
err: T,
@@ -123,8 +124,8 @@ 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, input.len())?;
- let mut out_b = ze::DeviceBuffer::<T>::new(&mut ctx, &dev, output.len())?;
+ 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]))?;
let ev0 = ze::Event::new(&event_pool, 0)?;
diff --git a/ptx/src/test/spirv_run/mov_address.ptx b/ptx/src/test/spirv_run/mov_address.ptx new file mode 100644 index 0000000..433fc0e --- /dev/null +++ b/ptx/src/test/spirv_run/mov_address.ptx @@ -0,0 +1,15 @@ +.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry mov_address(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .local .b8 __local_depot0[8];
+ .reg .u64 temp;
+
+ mov.u64 temp, __local_depot0;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/reg_slm.spvtxt b/ptx/src/test/spirv_run/mov_address.spvtxt index 6810fec..6810fec 100644 --- a/ptx/src/test/spirv_run/reg_slm.spvtxt +++ b/ptx/src/test/spirv_run/mov_address.spvtxt diff --git a/ptx/src/test/spirv_run/reg_slm.ptx b/ptx/src/test/spirv_run/reg_local.ptx index 929d116..fb234d8 100644 --- a/ptx/src/test/spirv_run/reg_slm.ptx +++ b/ptx/src/test/spirv_run/reg_local.ptx @@ -2,12 +2,12 @@ .target sm_30
.address_size 64
-.visible .entry reg_slm(
+.visible .entry reg_local(
.param .u64 input,
.param .u64 output
)
{
- .local .align 8 .b8 slm[8];
+ .local .align 8 .b8 local_x[8];
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .b64 temp;
@@ -16,11 +16,9 @@ ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
- mov.s64 unused, slm;
-
ld.global.u64 temp, [in_addr];
- st.u64 [slm], temp;
- ld.u64 temp, [slm];
+ st.u64 [local_x], temp;
+ ld.u64 temp, [local_x];
st.global.u64 [out_addr], temp;
ret;
}
\ No newline at end of file diff --git a/ptx/src/test/spirv_run/reg_local.spvtxt b/ptx/src/test/spirv_run/reg_local.spvtxt new file mode 100644 index 0000000..6810fec --- /dev/null +++ b/ptx/src/test/spirv_run/reg_local.spvtxt @@ -0,0 +1,46 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int64 + OpCapability Int8 + %25 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "add" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %28 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong +%_ptr_Generic_ulong = OpTypePointer Generic %ulong + %ulong_1 = OpConstant %ulong 1 + %1 = OpFunction %void None %28 + %8 = OpFunctionParameter %ulong + %9 = OpFunctionParameter %ulong + %23 = 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_ulong Function + %7 = OpVariable %_ptr_Function_ulong Function + OpStore %2 %8 + OpStore %3 %9 + %11 = OpLoad %ulong %2 + %10 = OpCopyObject %ulong %11 + OpStore %4 %10 + %13 = OpLoad %ulong %3 + %12 = OpCopyObject %ulong %13 + OpStore %5 %12 + %15 = OpLoad %ulong %4 + %21 = OpConvertUToPtr %_ptr_Generic_ulong %15 + %14 = OpLoad %ulong %21 + OpStore %6 %14 + %17 = OpLoad %ulong %6 + %16 = OpIAdd %ulong %17 %ulong_1 + OpStore %7 %16 + %18 = OpLoad %ulong %5 + %19 = OpLoad %ulong %7 + %22 = OpConvertUToPtr %_ptr_Generic_ulong %18 + OpStore %22 %19 + OpReturn + OpFunctionEnd |