Add test for indirect shared mem use

3 files changed, 88 insertions, 55 deletions

diff --git a/ptx/src/test/spirv_run/extern_shared_call.ptx b/ptx/src/test/spirv_run/extern_shared_call.ptx
index 6626783..76b2ae7 100644
--- a/ptx/src/test/spirv_run/extern_shared_call.ptx
+++ b/ptx/src/test/spirv_run/extern_shared_call.ptx
@@ -4,28 +4,17 @@
 

 .extern .shared .align 4 .b32 shared_mem[];

 

-.func (.param .u64 output) incr_shared_2_param(

-    .param .u64 .ptr .shared shared_mem_addr

-)

-{

-    .reg .u64 	    temp;

-    ld.shared.u64   temp, [shared_mem_addr];

-    add.u64         temp, temp, 2;

-    st.param.u64    [output], temp;

-    ret;

-}

-

-.func (.param .u64 output) incr_shared_2_global()

+.func incr_shared_2_global()

 {

     .reg .u64 	    temp;

     ld.shared.u64   temp, [shared_mem];

     add.u64         temp, temp, 2;

-    st.param.u64    [output], temp;

+    st.shared.u64    [shared_mem], temp;

     ret;

 }

 

 

-.visible .entry extern_shared(

+.visible .entry extern_shared_call(

 	.param .u64 input,

 	.param .u64 output

 )

@@ -39,6 +28,7 @@
 

     ld.global.u64       temp, [in_addr];

     st.shared.u64       [shared_mem], temp;

+    call incr_shared_2_global;

     ld.shared.u64       temp, [shared_mem];

     st.global.u64       [out_addr], temp;

 	ret;

diff --git a/ptx/src/test/spirv_run/extern_shared_call.spvtxt b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
index 84e7eac..d979193 100644
--- a/ptx/src/test/spirv_run/extern_shared_call.spvtxt
+++ b/ptx/src/test/spirv_run/extern_shared_call.spvtxt
@@ -2,52 +2,94 @@
                OpCapability Linkage
                OpCapability Addresses
                OpCapability Kernel
-               OpCapability Int64
                OpCapability Int8
-         %29 = OpExtInstImport "OpenCL.std"
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %48 = OpExtInstImport "OpenCL.std"
                OpMemoryModel Physical64 OpenCL
-               OpEntryPoint Kernel %1 "cvta"
+               OpEntryPoint Kernel %14 "extern_shared_call" %1
+               OpDecorate %1 Alignment 4
        %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup__ptr_Workgroup_uint = OpTypePointer Workgroup %_ptr_Workgroup_uint
+          %1 = OpVariable %_ptr_Workgroup__ptr_Workgroup_uint Workgroup
+      %uchar = OpTypeInt 8 0
+%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
+         %55 = OpTypeFunction %void %_ptr_Workgroup_uchar
+%_ptr_Function__ptr_Workgroup_uchar = OpTypePointer Function %_ptr_Workgroup_uchar
       %ulong = OpTypeInt 64 0
-         %32 = OpTypeFunction %void %ulong %ulong
 %_ptr_Function_ulong = OpTypePointer Function %ulong
-      %float = OpTypeFloat 32
-%_ptr_Function_float = OpTypePointer Function %float
-%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
-          %1 = OpFunction %void None %32
-          %7 = OpFunctionParameter %ulong
-          %8 = OpFunctionParameter %ulong
-         %27 = OpLabel
-          %2 = OpVariable %_ptr_Function_ulong Function
+%_ptr_Function__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint
+%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
+    %ulong_2 = OpConstant %ulong 2
+         %62 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
+          %2 = OpFunction %void None %55
+         %40 = OpFunctionParameter %_ptr_Workgroup_uchar
+         %56 = OpLabel
+         %41 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
           %3 = OpVariable %_ptr_Function_ulong Function
-          %4 = OpVariable %_ptr_Function_ulong Function
-          %5 = OpVariable %_ptr_Function_ulong Function
-          %6 = OpVariable %_ptr_Function_float Function
-               OpStore %2 %7
-               OpStore %3 %8
-         %10 = OpLoad %ulong %2
-          %9 = OpCopyObject %ulong %10
-               OpStore %4 %9
-         %12 = OpLoad %ulong %3
-         %11 = OpCopyObject %ulong %12
-               OpStore %5 %11
-         %14 = OpLoad %ulong %4
-         %22 = OpCopyObject %ulong %14
-         %21 = OpCopyObject %ulong %22
-         %13 = OpCopyObject %ulong %21
-               OpStore %4 %13
-         %16 = OpLoad %ulong %5
-         %24 = OpCopyObject %ulong %16
-         %23 = OpCopyObject %ulong %24
-         %15 = OpCopyObject %ulong %23
-               OpStore %5 %15
-         %18 = OpLoad %ulong %4
-         %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18
-         %17 = OpLoad %float %25
-               OpStore %6 %17
-         %19 = OpLoad %ulong %5
-         %20 = OpLoad %float %6
-         %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19
-               OpStore %26 %20
+               OpStore %41 %40
+               OpBranch %13
+         %13 = OpLabel
+         %42 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
+          %5 = OpLoad %_ptr_Workgroup_uint %42
+         %11 = OpBitcast %_ptr_Workgroup_ulong %5
+          %4 = OpLoad %ulong %11
+               OpStore %3 %4
+          %7 = OpLoad %ulong %3
+          %6 = OpIAdd %ulong %7 %ulong_2
+               OpStore %3 %6
+         %43 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
+          %8 = OpLoad %_ptr_Workgroup_uint %43
+          %9 = OpLoad %ulong %3
+         %12 = OpBitcast %_ptr_Workgroup_ulong %8
+               OpStore %12 %9
+               OpReturn
+               OpFunctionEnd
+         %14 = OpFunction %void None %62
+         %20 = OpFunctionParameter %ulong
+         %21 = OpFunctionParameter %ulong
+         %44 = OpFunctionParameter %_ptr_Workgroup_uchar
+         %63 = OpLabel
+         %45 = OpVariable %_ptr_Function__ptr_Workgroup_uchar 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
+         %19 = OpVariable %_ptr_Function_ulong Function
+               OpStore %45 %44
+               OpBranch %38
+         %38 = OpLabel
+               OpStore %15 %20
+               OpStore %16 %21
+         %23 = OpLoad %ulong %15
+         %22 = OpCopyObject %ulong %23
+               OpStore %17 %22
+         %25 = OpLoad %ulong %16
+         %24 = OpCopyObject %ulong %25
+               OpStore %18 %24
+         %27 = OpLoad %ulong %17
+         %34 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %27
+         %26 = OpLoad %ulong %34
+               OpStore %19 %26
+         %46 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %45
+         %28 = OpLoad %_ptr_Workgroup_uint %46
+         %29 = OpLoad %ulong %19
+         %35 = OpBitcast %_ptr_Workgroup_ulong %28
+               OpStore %35 %29
+         %65 = OpFunctionCall %void %2 %44
+         %47 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %45
+         %31 = OpLoad %_ptr_Workgroup_uint %47
+         %36 = OpBitcast %_ptr_Workgroup_ulong %31
+         %30 = OpLoad %ulong %36
+               OpStore %19 %30
+         %32 = OpLoad %ulong %18
+         %33 = OpLoad %ulong %19
+         %37 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %32
+               OpStore %37 %33
                OpReturn
                OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 4c5f9b3..3a8acb1 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -79,6 +79,7 @@ test_ptx!(min, [555i32, 444i32], [444i32]);
 test_ptx!(max, [555i32, 444i32], [555i32]);

 test_ptx!(global_array, [0xDEADu32], [1u32]);

 test_ptx!(extern_shared, [127u64], [127u64]);

+test_ptx!(extern_shared_call, [121u64], [123u64]);

 

 struct DisplayError<T: Debug> {

     err: T,