Add test for vector extract

3 files changed, 122 insertions, 0 deletions

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<T: Debug> {

     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