From 7c26568cbf017c55b27b72a7fcfe7761ce31e33c Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Sun, 27 Sep 2020 13:51:52 +0200 Subject: Add test for vector extract --- ptx/src/test/spirv_run/mod.rs | 1 + ptx/src/test/spirv_run/vector_extract.ptx | 24 +++++++ ptx/src/test/spirv_run/vector_extract.spvtxt | 97 ++++++++++++++++++++++++++++ 3 files changed, 122 insertions(+) create mode 100644 ptx/src/test/spirv_run/vector_extract.ptx create mode 100644 ptx/src/test/spirv_run/vector_extract.spvtxt (limited to 'ptx/src') diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 27dc063..5a16755 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -67,6 +67,7 @@ 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]); +test_ptx!(vector_extract, [1u8, 2u8, 3u8, 4u8], [3u8, 4u8, 1u8, 2u8]); struct DisplayError { err: T, diff --git a/ptx/src/test/spirv_run/vector_extract.ptx b/ptx/src/test/spirv_run/vector_extract.ptx new file mode 100644 index 0000000..8624f8a --- /dev/null +++ b/ptx/src/test/spirv_run/vector_extract.ptx @@ -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; +} diff --git a/ptx/src/test/spirv_run/vector_extract.spvtxt b/ptx/src/test/spirv_run/vector_extract.spvtxt new file mode 100644 index 0000000..ff0ee97 --- /dev/null +++ b/ptx/src/test/spirv_run/vector_extract.spvtxt @@ -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 -- cgit v1.2.3