Archive/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-04-28 21:47:57 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Add test for vector extract

Browse source
This commit is contained in:
Andrzej Janik 2020-09-27 13:51:52 +02:00
parent e0190fcbe1
commit 7c26568cbf
3 changed files with 122 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

1
ptx/src/test/spirv_run/mod.rs
View file

@ -67,6 +67,7 @@ test_ptx!(implicit_param, [34u32], [34u32]);
test_ptx!(pred_not, [10u64, 11u64], [2u64, 0u64]); test_ptx!(pred_not, [10u64, 11u64], [2u64, 0u64]);
test_ptx!(mad_s32, [2i32, 3i32, 4i32], [10i32, 10i32, 10i32]); 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]); test_ptx!(mul_wide, [0x01_00_00_00__01_00_00_00i64], [0x1_00_00_00_00_00_00i64]);
test_ptx!(vector_extract, [1u8, 2u8, 3u8, 4u8], [3u8, 4u8, 1u8, 2u8]);
struct DisplayError<T: Debug> { struct DisplayError<T: Debug> {
err: T, err: T,

24
ptx/src/test/spirv_run/vector_extract.ptx Normal file
View file

@ -0,0 +1,24 @@
.version 6.5
.target sm_30
.address_size 64
.visible .entry vector_extract(
.param .u64 input_p,
.param .u64 output_p
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .u16 temp1;
.reg .u16 temp2;
.reg .u16 temp3;
.reg .u16 temp4;
.reg .v4.u16 foo;
ld.global.v4.u8 {temp1, temp2, temp3, temp4}, [in_addr];
mov.v4.u16 foo, {temp2, temp3, temp4, temp1};
mov.v4.u16 {temp3, temp4, temp1, temp2}, foo;
mov.v4.u16 {temp4, temp1, temp2, temp3}, {temp3, temp4, temp1, temp2};
st.global.v4.u8 [out_addr], {temp1, temp2, temp3, temp4};
ret;
}

97
ptx/src/test/spirv_run/vector_extract.spvtxt Normal file
View file

@ -0,0 +1,97 @@
OpCapability GenericPointer
OpCapability Linkage
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
OpCapability Int8
%60 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %31 "vector"
%void = OpTypeVoid
%uint = OpTypeInt 32 0
%v2uint = OpTypeVector %uint 2
%64 = OpTypeFunction %v2uint %v2uint
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%_ptr_Function_uint = OpTypePointer Function %uint
%ulong = OpTypeInt 64 0
%68 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_v2uint = OpTypePointer Generic %v2uint
%1 = OpFunction %v2uint None %64
%7 = OpFunctionParameter %v2uint
%30 = OpLabel
%3 = OpVariable %_ptr_Function_v2uint Function
%2 = OpVariable %_ptr_Function_v2uint Function
%4 = OpVariable %_ptr_Function_v2uint Function
%5 = OpVariable %_ptr_Function_uint Function
%6 = OpVariable %_ptr_Function_uint Function
OpStore %3 %7
%9 = OpLoad %v2uint %3
%27 = OpCompositeExtract %uint %9 0
%8 = OpCopyObject %uint %27
OpStore %5 %8
%11 = OpLoad %v2uint %3
%28 = OpCompositeExtract %uint %11 1
%10 = OpCopyObject %uint %28
OpStore %6 %10
%13 = OpLoad %uint %5
%14 = OpLoad %uint %6
%12 = OpIAdd %uint %13 %14
OpStore %6 %12
%16 = OpLoad %v2uint %4
%17 = OpLoad %uint %6
%15 = OpCompositeInsert %v2uint %17 %16 0
OpStore %4 %15
%19 = OpLoad %v2uint %4
%20 = OpLoad %uint %6
%18 = OpCompositeInsert %v2uint %20 %19 1
OpStore %4 %18
%22 = OpLoad %v2uint %4
%23 = OpLoad %v2uint %4
%29 = OpCompositeExtract %uint %23 1
%21 = OpCompositeInsert %v2uint %29 %22 0
OpStore %4 %21
%25 = OpLoad %v2uint %4
%24 = OpCopyObject %v2uint %25
OpStore %2 %24
%26 = OpLoad %v2uint %2
OpReturnValue %26
OpFunctionEnd
%31 = OpFunction %void None %68
%40 = OpFunctionParameter %ulong
%41 = OpFunctionParameter %ulong
%58 = OpLabel
%32 = OpVariable %_ptr_Function_ulong Function
%33 = OpVariable %_ptr_Function_ulong Function
%34 = OpVariable %_ptr_Function_ulong Function
%35 = OpVariable %_ptr_Function_ulong Function
%36 = OpVariable %_ptr_Function_v2uint Function
%37 = OpVariable %_ptr_Function_uint Function
%38 = OpVariable %_ptr_Function_uint Function
%39 = OpVariable %_ptr_Function_ulong Function
OpStore %32 %40
OpStore %33 %41
%43 = OpLoad %ulong %32
%42 = OpCopyObject %ulong %43
OpStore %34 %42
%45 = OpLoad %ulong %33
%44 = OpCopyObject %ulong %45
OpStore %35 %44
%47 = OpLoad %ulong %34
%54 = OpConvertUToPtr %_ptr_Generic_v2uint %47
%46 = OpLoad %v2uint %54
OpStore %36 %46
%49 = OpLoad %v2uint %36
%48 = OpFunctionCall %v2uint %1 %49
OpStore %36 %48
%51 = OpLoad %v2uint %36
%55 = OpBitcast %ulong %51
%56 = OpCopyObject %ulong %55
%50 = OpCopyObject %ulong %56
OpStore %39 %50
%52 = OpLoad %ulong %35
%53 = OpLoad %v2uint %36
%57 = OpConvertUToPtr %_ptr_Generic_v2uint %52
OpStore %57 %53
OpReturn
OpFunctionEnd