Implement bfi instruction

3 files changed, 111 insertions, 0 deletions

diff --git a/ptx/src/test/spirv_run/bfi.ptx b/ptx/src/test/spirv_run/bfi.ptx
new file mode 100644
index 0000000..f2bca91
--- /dev/null
+++ b/ptx/src/test/spirv_run/bfi.ptx
@@ -0,0 +1,24 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry bfi(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	    in_addr;

+    .reg .u64 	    out_addr;

+    .reg .u32       temp<4>;

+

+	ld.param.u64 	in_addr, [input];

+    ld.param.u64 	out_addr, [output];

+

+    ld.u32          temp0, [in_addr];

+	ld.u32          temp1, [in_addr+4];

+	ld.u32          temp2, [in_addr+8];

+	ld.u32          temp3, [in_addr+12];

+	bfi.b32  		temp0, temp0, temp1, temp2, temp3;

+    st.u32          [out_addr], temp0;

+	ret;

+}

diff --git a/ptx/src/test/spirv_run/bfi.spvtxt b/ptx/src/test/spirv_run/bfi.spvtxt
new file mode 100644
index 0000000..a226f78
--- /dev/null
+++ b/ptx/src/test/spirv_run/bfi.spvtxt
@@ -0,0 +1,82 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %51 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "bfi"
+               OpDecorate %44 LinkageAttributes "__zluda_ptx_impl__bfi_b32" Import
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+         %54 = OpTypeFunction %uint %uint %uint %uint %uint
+      %ulong = OpTypeInt 64 0
+         %56 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Function_uint = OpTypePointer Function %uint
+%_ptr_Generic_uint = OpTypePointer Generic %uint
+    %ulong_4 = OpConstant %ulong 4
+    %ulong_8 = OpConstant %ulong 8
+   %ulong_12 = OpConstant %ulong 12
+         %44 = OpFunction %uint None %54
+         %46 = OpFunctionParameter %uint
+         %47 = OpFunctionParameter %uint
+         %48 = OpFunctionParameter %uint
+         %49 = OpFunctionParameter %uint
+               OpFunctionEnd
+          %1 = OpFunction %void None %56
+         %10 = OpFunctionParameter %ulong
+         %11 = OpFunctionParameter %ulong
+         %43 = 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_uint Function
+          %7 = OpVariable %_ptr_Function_uint Function
+          %8 = OpVariable %_ptr_Function_uint Function
+          %9 = OpVariable %_ptr_Function_uint Function
+               OpStore %2 %10
+               OpStore %3 %11
+         %12 = OpLoad %ulong %2 Aligned 8
+               OpStore %4 %12
+         %13 = OpLoad %ulong %3 Aligned 8
+               OpStore %5 %13
+         %15 = OpLoad %ulong %4
+         %35 = OpConvertUToPtr %_ptr_Generic_uint %15
+         %14 = OpLoad %uint %35 Aligned 4
+               OpStore %6 %14
+         %17 = OpLoad %ulong %4
+         %30 = OpIAdd %ulong %17 %ulong_4
+         %36 = OpConvertUToPtr %_ptr_Generic_uint %30
+         %16 = OpLoad %uint %36 Aligned 4
+               OpStore %7 %16
+         %19 = OpLoad %ulong %4
+         %32 = OpIAdd %ulong %19 %ulong_8
+         %37 = OpConvertUToPtr %_ptr_Generic_uint %32
+         %18 = OpLoad %uint %37 Aligned 4
+               OpStore %8 %18
+         %21 = OpLoad %ulong %4
+         %34 = OpIAdd %ulong %21 %ulong_12
+         %38 = OpConvertUToPtr %_ptr_Generic_uint %34
+         %20 = OpLoad %uint %38 Aligned 4
+               OpStore %9 %20
+         %23 = OpLoad %uint %6
+         %24 = OpLoad %uint %7
+         %25 = OpLoad %uint %8
+         %26 = OpLoad %uint %9
+         %40 = OpCopyObject %uint %23
+         %41 = OpCopyObject %uint %24
+         %39 = OpFunctionCall %uint %44 %41 %40 %25 %26
+         %22 = OpCopyObject %uint %39
+               OpStore %6 %22
+         %27 = OpLoad %ulong %5
+         %28 = OpLoad %uint %6
+         %42 = OpConvertUToPtr %_ptr_Generic_uint %27
+               OpStore %42 %28 Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 7c790eb..91e6113 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -139,6 +139,11 @@ test_ptx!(
     [0b11111000_11000001_00100010_10100000u32, 16u32, 8u32],

     [0b11000001u32]

 );

+test_ptx!(

+    bfi,

+    [0b10u32, 0b101u32, 0u32, 2u32],

+    [0b110u32]

+);

 test_ptx!(stateful_ld_st_simple, [121u64], [121u64]);

 test_ptx!(stateful_ld_st_ntid, [123u64], [123u64]);

 test_ptx!(stateful_ld_st_ntid_chain, [12651u64], [12651u64]);