Implement function pointers and activemask

4 files changed, 114 insertions, 9 deletions

diff --git a/ptx/src/test/spirv_run/activemask.spvtxt b/ptx/src/test/spirv_run/activemask.spvtxt
index c4ad55d..0753c95 100644
--- a/ptx/src/test/spirv_run/activemask.spvtxt
+++ b/ptx/src/test/spirv_run/activemask.spvtxt
@@ -7,21 +7,22 @@
                OpCapability Int64
                OpCapability Float16
                OpCapability Float64
-         %16 = OpExtInstImport "OpenCL.std"
+         %18 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
                OpEntryPoint Kernel %1 "activemask"
                OpExecutionMode %1 ContractionOff
+               OpDecorate %15 LinkageAttributes "__zluda_ptx_impl__activemask" Import
        %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+         %21 = OpTypeFunction %uint
       %ulong = OpTypeInt 64 0
-         %19 = OpTypeFunction %void %ulong %ulong
+         %23 = OpTypeFunction %void %ulong %ulong
 %_ptr_Function_ulong = OpTypePointer Function %ulong
-       %uint = OpTypeInt 32 0
 %_ptr_Function_uint = OpTypePointer Function %uint
-     %v4uint = OpTypeVector %uint 4
-       %bool = OpTypeBool
-       %true = OpConstantTrue %bool
 %_ptr_Generic_uint = OpTypePointer Generic %uint
-          %1 = OpFunction %void None %19
+         %15 = OpFunction %uint None %21
+               OpFunctionEnd
+          %1 = OpFunction %void None %23
           %6 = OpFunctionParameter %ulong
           %7 = OpFunctionParameter %ulong
          %14 = OpLabel
@@ -33,8 +34,7 @@
                OpStore %3 %7
           %8 = OpLoad %ulong %3 Aligned 8
                OpStore %4 %8
-         %26 = OpSubgroupBallotKHR %v4uint %true
-          %9 = OpCompositeExtract %uint %26 0
+          %9 = OpFunctionCall %uint %15
                OpStore %5 %9
          %10 = OpLoad %ulong %4
          %11 = OpLoad %uint %5
diff --git a/ptx/src/test/spirv_run/func_ptr.ptx b/ptx/src/test/spirv_run/func_ptr.ptx
new file mode 100644
index 0000000..aa94f2b
--- /dev/null
+++ b/ptx/src/test/spirv_run/func_ptr.ptx
@@ -0,0 +1,31 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.func (.reg .f32 out) foobar(.reg .f32 x, .reg .f32 y)
+{
+    add.f32     out, x, y;
+    ret;
+}
+
+.visible .entry func_ptr(
+	.param .u64 input,
+	.param .u64 output
+)
+{
+	.reg .u64 	    in_addr;
+    .reg .u64 	    out_addr;
+    .reg .u64 	    temp;
+    .reg .u64 	    temp2;
+    .reg .u64 	    f_addr;
+
+	ld.param.u64 	in_addr, [input];
+    ld.param.u64 	out_addr, [output];
+
+    ld.u64          temp, [in_addr];
+	add.u64		    temp2, temp, 1;
+    mov.u64         f_addr, foobar;
+    add.u64		    temp2, temp2, f_addr;
+    st.u64          [out_addr], temp2;
+	ret;
+}
diff --git a/ptx/src/test/spirv_run/func_ptr.spvtxt b/ptx/src/test/spirv_run/func_ptr.spvtxt
new file mode 100644
index 0000000..adc71eb
--- /dev/null
+++ b/ptx/src/test/spirv_run/func_ptr.spvtxt
@@ -0,0 +1,73 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %38 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %11 "func_ptr"
+               OpExecutionMode %11 ContractionOff
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+         %41 = OpTypeFunction %float %float %float
+%_ptr_Function_float = OpTypePointer Function %float
+      %ulong = OpTypeInt 64 0
+         %44 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+    %ulong_1 = OpConstant %ulong 1
+    %ulong_0 = OpConstant %ulong 0
+          %1 = OpFunction %float None %41
+          %5 = OpFunctionParameter %float
+          %6 = OpFunctionParameter %float
+         %10 = OpLabel
+          %3 = OpVariable %_ptr_Function_float Function
+          %4 = OpVariable %_ptr_Function_float Function
+          %2 = OpVariable %_ptr_Function_float Function
+               OpStore %3 %5
+               OpStore %4 %6
+          %8 = OpLoad %float %3
+          %9 = OpLoad %float %4
+          %7 = OpFAdd %float %8 %9
+               OpStore %2 %7
+               OpFunctionEnd
+         %11 = OpFunction %void None %44
+         %19 = OpFunctionParameter %ulong
+         %20 = OpFunctionParameter %ulong
+         %36 = OpLabel
+         %12 = OpVariable %_ptr_Function_ulong Function
+         %13 = OpVariable %_ptr_Function_ulong Function
+         %14 = OpVariable %_ptr_Function_ulong Function
+         %15 = OpVariable %_ptr_Function_ulong Function
+         %16 = OpVariable %_ptr_Function_ulong Function
+         %17 = OpVariable %_ptr_Function_ulong Function
+         %18 = OpVariable %_ptr_Function_ulong Function
+               OpStore %12 %19
+               OpStore %13 %20
+         %21 = OpLoad %ulong %12 Aligned 8
+               OpStore %14 %21
+         %22 = OpLoad %ulong %13 Aligned 8
+               OpStore %15 %22
+         %24 = OpLoad %ulong %14
+         %34 = OpConvertUToPtr %_ptr_Generic_ulong %24
+         %23 = OpLoad %ulong %34 Aligned 8
+               OpStore %16 %23
+         %26 = OpLoad %ulong %16
+         %25 = OpIAdd %ulong %26 %ulong_1
+               OpStore %17 %25
+         %27 = OpCopyObject %ulong %ulong_0
+               OpStore %18 %27
+         %29 = OpLoad %ulong %17
+         %30 = OpLoad %ulong %18
+         %28 = OpIAdd %ulong %29 %30
+               OpStore %17 %28
+         %31 = OpLoad %ulong %15
+         %32 = OpLoad %ulong %17
+         %35 = OpConvertUToPtr %_ptr_Generic_ulong %31
+               OpStore %35 %32 Aligned 8
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index f6b556e..0dcd0bb 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -209,6 +209,7 @@ test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]);
 test_ptx!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32]);

 test_ptx!(activemask, [0u32], [1u32]);

 test_ptx!(membar, [152731u32], [152731u32]);

+test_ptx!(func_ptr, [152731u64], [152732u64]);

 

 struct DisplayError<T: Debug> {

     err: T,