gistfile1.txt

__kernel void test_fn(__local uint2 *sSharedStorage, __global uint2 *srcValues, __global uint *offsets, __global uint2 *destBuffer, uint alignmentOffset )    
{
    int tid = get_global_id( 0 );
 sSharedStorage[ offsets[tid] ] = (uint2)(uint)0;
 sSharedStorage[ offsets[tid] +1 ] =  sSharedStorage[ offsets[tid] ];
   barrier( CLK_LOCAL_MEM_FENCE );

    vstore2( srcValues[ tid ], offsets[ tid ], ( (__local uint *)sSharedStorage ) + alignmentOffset );

   barrier( CLK_LOCAL_MEM_FENCE );

  int i;
  __local uint *sp = (__local uint*) (sSharedStorage + offsets[tid]) + alignmentOffset;
  __global uint *dp = (__global uint*) (destBuffer + offsets[tid]) + alignmentOffset;
  for( i = 0; (size_t)i < sizeof( sSharedStorage[0]) / sizeof( *sp ); i++ )
       dp[i] = sp[i];
}

; SPIR-V
; Version: 1.0
; Generator: Khronos; 17
; Bound: 60
; Schema: 0
               OpCapability Addresses
               OpCapability Kernel
               OpCapability Int64
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %2 "test_fn" %__spirv_BuiltInGlobalInvocationId
          %4 = OpString "kernel_arg_type.test_fn.uint2*,uint2*,uint*,uint2*,uint,"
               OpSource OpenCL_C 102000
               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
               OpName %sSharedStorage "sSharedStorage"
               OpName %srcValues "srcValues"
               OpName %offsets "offsets"
               OpName %destBuffer "destBuffer"
               OpName %alignmentOffset "alignmentOffset"
               OpName %entry "entry"
               OpName %call "call"
               OpName %idxprom "idxprom"
               OpName %arrayidx "arrayidx"
               OpName %idxprom1 "idxprom1"
               OpName %arrayidx2 "arrayidx2"
               OpName %idxprom5 "idxprom5"
               OpName %arrayidx6 "arrayidx6"
               OpName %add "add"
               OpName %idxprom9 "idxprom9"
               OpName %arrayidx10 "arrayidx10"
               OpName %arrayidx12 "arrayidx12"
               OpName %conv15 "conv15"
               OpName %idx_ext "idx.ext"
               OpName %add_ptr "add.ptr"
               OpName %idx_ext18 "idx.ext18"
               OpName %add_ptr21 "add.ptr21"
               OpName %add_ptr27 "add.ptr27"
               OpName %arrayidx31_1 "arrayidx31.1"
               OpName %arrayidx33_1 "arrayidx33.1"
               OpDecorate %30 FuncParamAttr NoCapture
         %30 = OpDecorationGroup
               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
               OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
               OpGroupDecorate %30 %srcValues %offsets %destBuffer
      %ulong = OpTypeInt 64 0
       %uint = OpTypeInt 32 0
   %ulong_32 = OpConstant %ulong 32
     %uint_1 = OpConstant %uint 1
     %uint_2 = OpConstant %uint 2
   %uint_272 = OpConstant %uint 272
    %ulong_0 = OpConstant %ulong 0
    %ulong_1 = OpConstant %ulong 1
    %v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
       %void = OpTypeVoid
     %v2uint = OpTypeVector %uint 2
%_ptr_Workgroup_v2uint = OpTypePointer Workgroup %v2uint
%_ptr_CrossWorkgroup_v2uint = OpTypePointer CrossWorkgroup %v2uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
         %46 = OpTypeFunction %void %_ptr_Workgroup_v2uint %_ptr_CrossWorkgroup_v2uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_v2uint %uint
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
         %48 = OpConstantNull %v2uint
          %2 = OpFunction %void None %46
%sSharedStorage = OpFunctionParameter %_ptr_Workgroup_v2uint
  %srcValues = OpFunctionParameter %_ptr_CrossWorkgroup_v2uint
    %offsets = OpFunctionParameter %_ptr_CrossWorkgroup_uint
 %destBuffer = OpFunctionParameter %_ptr_CrossWorkgroup_v2uint
%alignmentOffset = OpFunctionParameter %uint
      %entry = OpLabel
         %49 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
       %call = OpCompositeExtract %ulong %49 0
         %50 = OpShiftLeftLogical %ulong %call %ulong_32
    %idxprom = OpShiftRightArithmetic %ulong %50 %ulong_32
   %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %offsets %idxprom
         %51 = OpLoad %uint %arrayidx Aligned 4
   %idxprom1 = OpUConvert %ulong %51
  %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_Workgroup_v2uint %sSharedStorage %idxprom1
               OpStore %arrayidx2 %48 Aligned 8
         %52 = OpLoad %uint %arrayidx Aligned 4
   %idxprom5 = OpUConvert %ulong %52
  %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_Workgroup_v2uint %sSharedStorage %idxprom5
         %53 = OpLoad %v2uint %arrayidx6 Aligned 8
        %add = OpIAdd %uint %52 %uint_1
   %idxprom9 = OpUConvert %ulong %add
 %arrayidx10 = OpInBoundsPtrAccessChain %_ptr_Workgroup_v2uint %sSharedStorage %idxprom9
               OpStore %arrayidx10 %53 Aligned 8
               OpControlBarrier %uint_2 %uint_2 %uint_272
 %arrayidx12 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v2uint %srcValues %idxprom
         %54 = OpLoad %v2uint %arrayidx12 Aligned 8
         %55 = OpLoad %uint %arrayidx Aligned 4
     %conv15 = OpUConvert %ulong %55
    %idx_ext = OpUConvert %ulong %alignmentOffset
    %add_ptr = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %sSharedStorage %ulong_0 %idx_ext
         %56 = OpExtInst %void %1 vstoren %54 %conv15 %add_ptr
               OpControlBarrier %uint_2 %uint_2 %uint_272
         %57 = OpLoad %uint %arrayidx Aligned 4
  %idx_ext18 = OpUConvert %ulong %57
  %add_ptr21 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %sSharedStorage %idx_ext18 %idx_ext
  %add_ptr27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %destBuffer %idx_ext18 %idx_ext
         %58 = OpLoad %uint %add_ptr21 Aligned 4
               OpStore %add_ptr27 %58 Aligned 4
%arrayidx31_1 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %add_ptr21 %ulong_1
         %59 = OpLoad %uint %arrayidx31_1 Aligned 4
%arrayidx33_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %add_ptr27 %ulong_1
               OpStore %arrayidx33_1 %59 Aligned 4
               OpReturn
               OpFunctionEnd