diff options
Diffstat (limited to 'ptx/src/test/spirv_run')
| -rw-r--r-- | ptx/src/test/spirv_run/local_align.ptx | 21 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/local_align.spvtxt | 38 | ||||
| -rw-r--r-- | ptx/src/test/spirv_run/mod.rs | 1 |
3 files changed, 60 insertions, 0 deletions
diff --git a/ptx/src/test/spirv_run/local_align.ptx b/ptx/src/test/spirv_run/local_align.ptx new file mode 100644 index 0000000..6e10de3 --- /dev/null +++ b/ptx/src/test/spirv_run/local_align.ptx @@ -0,0 +1,21 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.visible .entry local_align( + .param .u64 input, + .param .u64 output +) +{ + .local .align 8 .b8 __local_depot0[8]; + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .u64 temp; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.u64 temp, [in_addr]; + st.u64 [out_addr], temp; + ret; +} diff --git a/ptx/src/test/spirv_run/local_align.spvtxt b/ptx/src/test/spirv_run/local_align.spvtxt new file mode 100644 index 0000000..beefb76 --- /dev/null +++ b/ptx/src/test/spirv_run/local_align.spvtxt @@ -0,0 +1,38 @@ + OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int64
+ OpCapability Int8
+ %1 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %5 "local_align"
+ OpDecorate %8 Alignment 8
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+ %4 = OpTypeFunction %void %ulong %ulong
+ %uchar = OpTypeInt 8 0
+%_arr_uchar_8 = OpTypeArray %uchar %8
+%_ptr_Function__arr_uchar_8 = OpTypePointer Function %_arr_uchar_8
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+ %5 = OpFunction %void None %4
+ %6 = OpFunctionParameter %ulong
+ %7 = OpFunctionParameter %ulong
+ %18 = OpLabel
+ %8 = OpVariable %_ptr_Function__arr_uchar_8 Workgroup
+ %9 = OpVariable %_ptr_Function_ulong Function
+ %10 = OpVariable %_ptr_Function_ulong Function
+ %11 = OpVariable %_ptr_Function_ulong Function
+ OpStore %9 %6
+ OpStore %10 %7
+ %13 = OpLoad %ulong %9
+ %16 = OpConvertUToPtr %_ptr_Generic_ulong %13
+ %12 = OpLoad %ulong %16
+ OpStore %11 %12
+ %14 = OpLoad %ulong %10
+ %15 = OpLoad %ulong %11
+ %17 = OpConvertUToPtr %_ptr_Generic_ulong %14
+ OpStore %17 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 23852a1..8883669 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -51,6 +51,7 @@ test_ptx!(shl, [11u64], [44u64]); test_ptx!(cvt_sat_s_u, [-1i32], [0i32]);
test_ptx!(cvta, [3.0f32], [3.0f32]);
test_ptx!(block, [1u64], [2u64]);
+test_ptx!(local_align, [1u64], [1u64]);
struct DisplayError<T: Debug> {
err: T,
|