Add test for indirect shared mem use

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(
+.func incr_shared_2_global()
-    .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()
 {
     .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;

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__ptr_Workgroup_uint = OpTypePointer Function %_ptr_Workgroup_uint
-%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_Workgroup_ulong = OpTypePointer Workgroup %ulong
-%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
+    %ulong_2 = OpConstant %ulong 2
-          %1 = OpFunction %void None %32
+         %62 = OpTypeFunction %void %ulong %ulong %_ptr_Workgroup_uchar
-          %7 = OpFunctionParameter %ulong
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
-          %8 = OpFunctionParameter %ulong
+          %2 = OpFunction %void None %55
-         %27 = OpLabel
+         %40 = OpFunctionParameter %_ptr_Workgroup_uchar
-          %2 = OpVariable %_ptr_Function_ulong Function
+         %56 = OpLabel
+         %41 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
           %3 = OpVariable %_ptr_Function_ulong Function
-          %4 = OpVariable %_ptr_Function_ulong Function
+               OpStore %41 %40
-          %5 = OpVariable %_ptr_Function_ulong Function
+               OpBranch %13
-          %6 = OpVariable %_ptr_Function_float Function
+         %13 = OpLabel
-               OpStore %2 %7
+         %42 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
-               OpStore %3 %8
+          %5 = OpLoad %_ptr_Workgroup_uint %42
-         %10 = OpLoad %ulong %2
+         %11 = OpBitcast %_ptr_Workgroup_ulong %5
-          %9 = OpCopyObject %ulong %10
+          %4 = OpLoad %ulong %11
-               OpStore %4 %9
+               OpStore %3 %4
-         %12 = OpLoad %ulong %3
+          %7 = OpLoad %ulong %3
-         %11 = OpCopyObject %ulong %12
+          %6 = OpIAdd %ulong %7 %ulong_2
-               OpStore %5 %11
+               OpStore %3 %6
-         %14 = OpLoad %ulong %4
+         %43 = OpBitcast %_ptr_Function__ptr_Workgroup_uint %41
-         %22 = OpCopyObject %ulong %14
+          %8 = OpLoad %_ptr_Workgroup_uint %43
-         %21 = OpCopyObject %ulong %22
+          %9 = OpLoad %ulong %3
-         %13 = OpCopyObject %ulong %21
+         %12 = OpBitcast %_ptr_Workgroup_ulong %8
-               OpStore %4 %13
+               OpStore %12 %9
-         %16 = OpLoad %ulong %5
+               OpReturn
-         %24 = OpCopyObject %ulong %16
+               OpFunctionEnd
-         %23 = OpCopyObject %ulong %24
+         %14 = OpFunction %void None %62
-         %15 = OpCopyObject %ulong %23
+         %20 = OpFunctionParameter %ulong
-               OpStore %5 %15
+         %21 = OpFunctionParameter %ulong
-         %18 = OpLoad %ulong %4
+         %44 = OpFunctionParameter %_ptr_Workgroup_uchar
-         %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %18
+         %63 = OpLabel
-         %17 = OpLoad %float %25
+         %45 = OpVariable %_ptr_Function__ptr_Workgroup_uchar Function
-               OpStore %6 %17
+         %15 = OpVariable %_ptr_Function_ulong Function
-         %19 = OpLoad %ulong %5
+         %16 = OpVariable %_ptr_Function_ulong Function
-         %20 = OpLoad %float %6
+         %17 = OpVariable %_ptr_Function_ulong Function
-         %26 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %19
+         %18 = OpVariable %_ptr_Function_ulong Function
-               OpStore %26 %20
+         %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

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,