external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber

   1 # Test OpPtrAccessChain applied to workgroup memory.
   2 #
   3 # Derived from the following OpenCL C:
   4 #
   5 #int get_data(local int *d);
   6 #
   7 #int get_data(local int *d) {
   8 #   return d[1];
   9 #}
  10 #
  11 #kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
  12 #   local int data[17];
  13 #   uint i = get_local_id(0);
  14 #   data[i] = A[i] * B[i];
  15 #
  16 #   if (i == 0) data[16] = 0;
  17 #
  18 #   C[i] = get_data(&data[i]);
  19 #}
  20 #
  21 # Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>` with an incorrect ArrayStride decoration
  22 # added. This decoration should be ignored, so it should give the same results as ArrayStride == 4.
  23
  24 [compute shader spirv]
  25                OpCapability Shader
  26                OpCapability VariablePointers
  27                OpExtension "SPV_KHR_storage_buffer_storage_class"
  28                OpExtension "SPV_KHR_variable_pointers"
  29                OpMemoryModel Logical GLSL450
  30                OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
  31                OpExecutionMode %30 LocalSize 16 1 1
  32                OpSource OpenCL_C 120
  33                OpDecorate %_runtimearr_uint ArrayStride 4
  34                OpMemberDecorate %_struct_3 0 Offset 0
  35                OpDecorate %_struct_3 Block
  36                OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
  37                OpDecorate %22 DescriptorSet 0
  38                OpDecorate %22 Binding 0
  39                OpDecorate %23 DescriptorSet 0
  40                OpDecorate %23 Binding 1
  41                OpDecorate %24 DescriptorSet 0
  42                OpDecorate %24 Binding 2
  43                OpDecorate %_arr_uint_uint_17 ArrayStride 4
  44                OpDecorate %_ptr_Workgroup_uint ArrayStride 8
  45        %uint = OpTypeInt 32 0
  46 %_runtimearr_uint = OpTypeRuntimeArray %uint
  47   %_struct_3 = OpTypeStruct %_runtimearr_uint
  48 %_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
  49 %_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
  50           %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
  51        %void = OpTypeVoid
  52           %8 = OpTypeFunction %void
  53      %v3uint = OpTypeVector %uint 3
  54 %_ptr_Input_v3uint = OpTypePointer Input %v3uint
  55 %_ptr_Input_uint = OpTypePointer Input %uint
  56 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
  57     %uint_17 = OpConstant %uint 17
  58 %_arr_uint_uint_17 = OpTypeArray %uint %uint_17
  59 %_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
  60        %bool = OpTypeBool
  61      %uint_1 = OpConstant %uint 1
  62      %uint_0 = OpConstant %uint 0
  63     %uint_16 = OpConstant %uint 16
  64          %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
  65 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
  66          %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
  67          %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
  68          %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
  69          %25 = OpFunction %uint Pure %6
  70          %26 = OpFunctionParameter %_ptr_Workgroup_uint
  71          %27 = OpLabel
  72          %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
  73          %29 = OpLoad %uint %28
  74                OpReturnValue %29
  75                OpFunctionEnd
  76          %30 = OpFunction %void None %8
  77          %31 = OpLabel
  78          %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
  79          %33 = OpLoad %uint %32
  80          %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
  81          %35 = OpLoad %uint %34
  82          %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
  83          %37 = OpLoad %uint %36
  84          %38 = OpIMul %uint %37 %35
  85          %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
  86                OpStore %39 %38
  87          %40 = OpIEqual %bool %33 %uint_0
  88                OpSelectionMerge %43 None
  89                OpBranchConditional %40 %41 %43
  90          %41 = OpLabel
  91          %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
  92                OpStore %42 %uint_0
  93                OpBranch %43
  94          %43 = OpLabel
  95          %44 = OpFunctionCall %uint %25 %39
  96          %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
  97                OpStore %45 %44
  98                OpReturn
  99                OpFunctionEnd
 100
 101 [test]
 102 # A[]
 103 ssbo 0:0 subdata int 0  0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
 104 # B[]
 105 ssbo 0:1 subdata int 0  1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
 106 # The answer array C[]
 107 ssbo 0:2 subdata int 0  -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
 108
 109 compute 1 1 1
 110
 111 probe ssbo int 0:2 0 ==  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0