SKIP: FAILED

<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: Tables = struct @align(4) {
  edges:array<u32, 256> @offset(0)
  tris:array<i32, 4096> @offset(1024)
}

IsosurfaceVolume = struct @align(16) {
  min:vec3<f32> @offset(0)
  max:vec3<f32> @offset(16)
  step:vec3<f32> @offset(32)
  size:vec3<u32> @offset(48)
  threshold:f32 @offset(60)
  values:array<f32> @offset(64)
}

PositionBuffer = struct @align(4) {
  values:array<f32> @offset(0)
}

NormalBuffer = struct @align(4) {
  values:array<f32> @offset(0)
}

IndexBuffer = struct @align(4) {
  tris:array<u32> @offset(0)
}

DrawIndirectArgs = struct @align(4) {
  vc:u32 @offset(0)
  vertexCount:atomic<u32> @offset(4)
  firstVertex:u32 @offset(8)
  firstInstance:u32 @offset(12)
  indexCount:atomic<u32> @offset(16)
  indexedInstanceCount:u32 @offset(20)
  indexedFirstIndex:u32 @offset(24)
  indexedBaseVertex:u32 @offset(28)
  indexedFirstInstance:u32 @offset(32)
}

%b1 = block {  # root
  %tables:ptr<storage, Tables, read> = var @binding_point(0, 0)
  %volume:ptr<storage, IsosurfaceVolume, read_write> = var @binding_point(0, 1)
  %positionsOut:ptr<storage, PositionBuffer, read_write> = var @binding_point(0, 2)
  %normalsOut:ptr<storage, NormalBuffer, read_write> = var @binding_point(0, 3)
  %indicesOut:ptr<storage, IndexBuffer, read_write> = var @binding_point(0, 4)
  %drawOut:ptr<storage, DrawIndirectArgs, read_write> = var @binding_point(0, 5)
  %positions:ptr<private, array<vec3<f32>, 12>, read_write> = var
  %normals:ptr<private, array<vec3<f32>, 12>, read_write> = var
  %indices:ptr<private, array<u32, 12>, read_write> = var
  %cubeVerts:ptr<private, u32, read_write> = var, 0u
}

%valueAt = func(%index:vec3<u32>):f32 -> %b2 {
  %b2 = block {
    %13:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %14:vec3<u32> = load %13
    %15:vec3<bool> = gte %index, %14
    %16:bool = any %15
    if %16 [t: %b3] {  # if_1
      %b3 = block {  # true
        ret 0.0f
      }
    }
    %17:u32 = access %index, 0u
    %18:u32 = access %index, 1u
    %19:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %20:u32 = load_vector_element %19, 0u
    %21:u32 = mul %18, %20
    %22:u32 = add %17, %21
    %23:u32 = access %index, 2u
    %24:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %25:u32 = load_vector_element %24, 0u
    %26:u32 = mul %23, %25
    %27:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %28:u32 = load_vector_element %27, 1u
    %29:u32 = mul %26, %28
    %valueIndex:u32 = add %22, %29
    %31:ptr<storage, f32, read_write> = access %volume, 5u, %valueIndex
    %32:f32 = load %31
    ret %32
  }
}
%positionAt = func(%index_1:vec3<u32>):vec3<f32> -> %b4 {  # %index_1: 'index'
  %b4 = block {
    %35:ptr<storage, vec3<f32>, read_write> = access %volume, 0u
    %36:vec3<f32> = load %35
    %37:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %38:vec3<f32> = load %37
    %39:vec3<u32> = swizzle %index_1, xyz
    %40:vec3<f32> = convert %39
    %41:vec3<f32> = mul %38, %40
    %42:vec3<f32> = add %36, %41
    ret %42
  }
}
%normalAt = func(%index_2:vec3<u32>):vec3<f32> -> %b5 {  # %index_2: 'index'
  %b5 = block {
    %45:vec3<u32> = sub %index_2, vec3<u32>(1u, 0u, 0u)
    %46:f32 = call %valueAt, %45
    %47:vec3<u32> = add %index_2, vec3<u32>(1u, 0u, 0u)
    %48:f32 = call %valueAt, %47
    %49:f32 = sub %46, %48
    %50:vec3<u32> = sub %index_2, vec3<u32>(0u, 1u, 0u)
    %51:f32 = call %valueAt, %50
    %52:vec3<u32> = add %index_2, vec3<u32>(0u, 1u, 0u)
    %53:f32 = call %valueAt, %52
    %54:f32 = sub %51, %53
    %55:vec3<u32> = sub %index_2, vec3<u32>(0u, 0u, 1u)
    %56:f32 = call %valueAt, %55
    %57:vec3<u32> = add %index_2, vec3<u32>(0u, 0u, 1u)
    %58:f32 = call %valueAt, %57
    %59:f32 = sub %56, %58
    %60:vec3<f32> = construct %49, %54, %59
    ret %60
  }
}
%interpX = func(%index_3:u32, %i:vec3<u32>, %va:f32, %vb:f32):void -> %b6 {  # %index_3: 'index'
  %b6 = block {
    %66:ptr<storage, f32, read_write> = access %volume, 4u
    %67:f32 = load %66
    %68:f32 = sub %67, %va
    %69:f32 = sub %vb, %va
    %mu:f32 = div %68, %69
    %71:u32 = load %cubeVerts
    %72:ptr<private, vec3<f32>, read_write> = access %positions, %71
    %73:vec3<f32> = call %positionAt, %i
    %74:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %75:f32 = load_vector_element %74, 0u
    %76:f32 = mul %75, %mu
    %77:vec3<f32> = construct %76, 0.0f, 0.0f
    %78:vec3<f32> = add %73, %77
    store %72, %78
    %na:vec3<f32> = call %normalAt, %i
    %80:vec3<u32> = add %i, vec3<u32>(1u, 0u, 0u)
    %nb:vec3<f32> = call %normalAt, %80
    %82:u32 = load %cubeVerts
    %83:ptr<private, vec3<f32>, read_write> = access %normals, %82
    %84:vec3<f32> = construct %mu, %mu, %mu
    %85:vec3<f32> = mix %na, %nb, %84
    store %83, %85
    %86:ptr<private, u32, read_write> = access %indices, %index_3
    %87:u32 = load %cubeVerts
    store %86, %87
    %88:u32 = load %cubeVerts
    %89:u32 = add %88, 1u
    store %cubeVerts, %89
    ret
  }
}
%interpY = func(%index_4:u32, %i_1:vec3<u32>, %va_1:f32, %vb_1:f32):void -> %b7 {  # %index_4: 'index', %i_1: 'i', %va_1: 'va', %vb_1: 'vb'
  %b7 = block {
    %95:ptr<storage, f32, read_write> = access %volume, 4u
    %96:f32 = load %95
    %97:f32 = sub %96, %va_1
    %98:f32 = sub %vb_1, %va_1
    %mu_1:f32 = div %97, %98  # %mu_1: 'mu'
    %100:u32 = load %cubeVerts
    %101:ptr<private, vec3<f32>, read_write> = access %positions, %100
    %102:vec3<f32> = call %positionAt, %i_1
    %103:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %104:f32 = load_vector_element %103, 1u
    %105:f32 = mul %104, %mu_1
    %106:vec3<f32> = construct 0.0f, %105, 0.0f
    %107:vec3<f32> = add %102, %106
    store %101, %107
    %na_1:vec3<f32> = call %normalAt, %i_1  # %na_1: 'na'
    %109:vec3<u32> = add %i_1, vec3<u32>(0u, 1u, 0u)
    %nb_1:vec3<f32> = call %normalAt, %109  # %nb_1: 'nb'
    %111:u32 = load %cubeVerts
    %112:ptr<private, vec3<f32>, read_write> = access %normals, %111
    %113:vec3<f32> = construct %mu_1, %mu_1, %mu_1
    %114:vec3<f32> = mix %na_1, %nb_1, %113
    store %112, %114
    %115:ptr<private, u32, read_write> = access %indices, %index_4
    %116:u32 = load %cubeVerts
    store %115, %116
    %117:u32 = load %cubeVerts
    %118:u32 = add %117, 1u
    store %cubeVerts, %118
    ret
  }
}
%interpZ = func(%index_5:u32, %i_2:vec3<u32>, %va_2:f32, %vb_2:f32):void -> %b8 {  # %index_5: 'index', %i_2: 'i', %va_2: 'va', %vb_2: 'vb'
  %b8 = block {
    %124:ptr<storage, f32, read_write> = access %volume, 4u
    %125:f32 = load %124
    %126:f32 = sub %125, %va_2
    %127:f32 = sub %vb_2, %va_2
    %mu_2:f32 = div %126, %127  # %mu_2: 'mu'
    %129:u32 = load %cubeVerts
    %130:ptr<private, vec3<f32>, read_write> = access %positions, %129
    %131:vec3<f32> = call %positionAt, %i_2
    %132:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %133:f32 = load_vector_element %132, 2u
    %134:f32 = mul %133, %mu_2
    %135:vec3<f32> = construct 0.0f, 0.0f, %134
    %136:vec3<f32> = add %131, %135
    store %130, %136
    %na_2:vec3<f32> = call %normalAt, %i_2  # %na_2: 'na'
    %138:vec3<u32> = add %i_2, vec3<u32>(0u, 0u, 1u)
    %nb_2:vec3<f32> = call %normalAt, %138  # %nb_2: 'nb'
    %140:u32 = load %cubeVerts
    %141:ptr<private, vec3<f32>, read_write> = access %normals, %140
    %142:vec3<f32> = construct %mu_2, %mu_2, %mu_2
    %143:vec3<f32> = mix %na_2, %nb_2, %142
    store %141, %143
    %144:ptr<private, u32, read_write> = access %indices, %index_5
    %145:u32 = load %cubeVerts
    store %144, %145
    %146:u32 = load %cubeVerts
    %147:u32 = add %146, 1u
    store %cubeVerts, %147
    ret
  }
}
%computeMain = @compute @workgroup_size(4, 4, 4) func(%global_id:vec3<u32> [@global_invocation_id]):void -> %b9 {
  %b9 = block {
    %i0:vec3<u32> = let %global_id
    %i1:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 0u)
    %i2:vec3<u32> = add %global_id, vec3<u32>(1u, 1u, 0u)
    %i3:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 0u)
    %i4:vec3<u32> = add %global_id, vec3<u32>(0u, 0u, 1u)
    %i5:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 1u)
    %i6:vec3<u32> = add %global_id, vec3<u32>(1u)
    %i7:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 1u)
    %v0:f32 = call %valueAt, %i0
    %v1:f32 = call %valueAt, %i1
    %v2:f32 = call %valueAt, %i2
    %v3:f32 = call %valueAt, %i3
    %v4:f32 = call %valueAt, %i4
    %v5:f32 = call %valueAt, %i5
    %v6:f32 = call %valueAt, %i6
    %v7:f32 = call %valueAt, %i7
    %cubeIndex:ptr<function, u32, read_write> = var, 0u
    %167:ptr<storage, f32, read_write> = access %volume, 4u
    %168:f32 = load %167
    %169:bool = lt %v0, %168
    if %169 [t: %b10] {  # if_2
      %b10 = block {  # true
        %170:u32 = load %cubeIndex
        %171:u32 = or %170, 1u
        store %cubeIndex, %171
        exit_if  # if_2
      }
    }
    %172:ptr<storage, f32, read_write> = access %volume, 4u
    %173:f32 = load %172
    %174:bool = lt %v1, %173
    if %174 [t: %b11] {  # if_3
      %b11 = block {  # true
        %175:u32 = load %cubeIndex
        %176:u32 = or %175, 2u
        store %cubeIndex, %176
        exit_if  # if_3
      }
    }
    %177:ptr<storage, f32, read_write> = access %volume, 4u
    %178:f32 = load %177
    %179:bool = lt %v2, %178
    if %179 [t: %b12] {  # if_4
      %b12 = block {  # true
        %180:u32 = load %cubeIndex
        %181:u32 = or %180, 4u
        store %cubeIndex, %181
        exit_if  # if_4
      }
    }
    %182:ptr<storage, f32, read_write> = access %volume, 4u
    %183:f32 = load %182
    %184:bool = lt %v3, %183
    if %184 [t: %b13] {  # if_5
      %b13 = block {  # true
        %185:u32 = load %cubeIndex
        %186:u32 = or %185, 8u
        store %cubeIndex, %186
        exit_if  # if_5
      }
    }
    %187:ptr<storage, f32, read_write> = access %volume, 4u
    %188:f32 = load %187
    %189:bool = lt %v4, %188
    if %189 [t: %b14] {  # if_6
      %b14 = block {  # true
        %190:u32 = load %cubeIndex
        %191:u32 = or %190, 16u
        store %cubeIndex, %191
        exit_if  # if_6
      }
    }
    %192:ptr<storage, f32, read_write> = access %volume, 4u
    %193:f32 = load %192
    %194:bool = lt %v5, %193
    if %194 [t: %b15] {  # if_7
      %b15 = block {  # true
        %195:u32 = load %cubeIndex
        %196:u32 = or %195, 32u
        store %cubeIndex, %196
        exit_if  # if_7
      }
    }
    %197:ptr<storage, f32, read_write> = access %volume, 4u
    %198:f32 = load %197
    %199:bool = lt %v6, %198
    if %199 [t: %b16] {  # if_8
      %b16 = block {  # true
        %200:u32 = load %cubeIndex
        %201:u32 = or %200, 64u
        store %cubeIndex, %201
        exit_if  # if_8
      }
    }
    %202:ptr<storage, f32, read_write> = access %volume, 4u
    %203:f32 = load %202
    %204:bool = lt %v7, %203
    if %204 [t: %b17] {  # if_9
      %b17 = block {  # true
        %205:u32 = load %cubeIndex
        %206:u32 = or %205, 128u
        store %cubeIndex, %206
        exit_if  # if_9
      }
    }
    %207:u32 = load %cubeIndex
    %208:ptr<storage, u32, read> = access %tables, 0u, %207
    %edges:u32 = load %208
    %210:u32 = and %edges, 1u
    %211:bool = neq %210, 0u
    if %211 [t: %b18] {  # if_10
      %b18 = block {  # true
        %212:void = call %interpX, 0u, %i0, %v0, %v1
        exit_if  # if_10
      }
    }
    %213:u32 = and %edges, 2u
    %214:bool = neq %213, 0u
    if %214 [t: %b19] {  # if_11
      %b19 = block {  # true
        %215:void = call %interpY, 1u, %i1, %v1, %v2
        exit_if  # if_11
      }
    }
    %216:u32 = and %edges, 4u
    %217:bool = neq %216, 0u
    if %217 [t: %b20] {  # if_12
      %b20 = block {  # true
        %218:void = call %interpX, 2u, %i3, %v3, %v2
        exit_if  # if_12
      }
    }
    %219:u32 = and %edges, 8u
    %220:bool = neq %219, 0u
    if %220 [t: %b21] {  # if_13
      %b21 = block {  # true
        %221:void = call %interpY, 3u, %i0, %v0, %v3
        exit_if  # if_13
      }
    }
    %222:u32 = and %edges, 16u
    %223:bool = neq %222, 0u
    if %223 [t: %b22] {  # if_14
      %b22 = block {  # true
        %224:void = call %interpX, 4u, %i4, %v4, %v5
        exit_if  # if_14
      }
    }
    %225:u32 = and %edges, 32u
    %226:bool = neq %225, 0u
    if %226 [t: %b23] {  # if_15
      %b23 = block {  # true
        %227:void = call %interpY, 5u, %i5, %v5, %v6
        exit_if  # if_15
      }
    }
    %228:u32 = and %edges, 64u
    %229:bool = neq %228, 0u
    if %229 [t: %b24] {  # if_16
      %b24 = block {  # true
        %230:void = call %interpX, 6u, %i7, %v7, %v6
        exit_if  # if_16
      }
    }
    %231:u32 = and %edges, 128u
    %232:bool = neq %231, 0u
    if %232 [t: %b25] {  # if_17
      %b25 = block {  # true
        %233:void = call %interpY, 7u, %i4, %v4, %v7
        exit_if  # if_17
      }
    }
    %234:u32 = and %edges, 256u
    %235:bool = neq %234, 0u
    if %235 [t: %b26] {  # if_18
      %b26 = block {  # true
        %236:void = call %interpZ, 8u, %i0, %v0, %v4
        exit_if  # if_18
      }
    }
    %237:u32 = and %edges, 512u
    %238:bool = neq %237, 0u
    if %238 [t: %b27] {  # if_19
      %b27 = block {  # true
        %239:void = call %interpZ, 9u, %i1, %v1, %v5
        exit_if  # if_19
      }
    }
    %240:u32 = and %edges, 1024u
    %241:bool = neq %240, 0u
    if %241 [t: %b28] {  # if_20
      %b28 = block {  # true
        %242:void = call %interpZ, 10u, %i2, %v2, %v6
        exit_if  # if_20
      }
    }
    %243:u32 = and %edges, 2048u
    %244:bool = neq %243, 0u
    if %244 [t: %b29] {  # if_21
      %b29 = block {  # true
        %245:void = call %interpZ, 11u, %i3, %v3, %v7
        exit_if  # if_21
      }
    }
    %246:u32 = load %cubeIndex
    %247:u32 = shiftl %246, 4u
    %triTableOffset:u32 = add %247, 1u
    %249:u32 = sub %triTableOffset, 1u
    %250:ptr<storage, i32, read> = access %tables, 1u, %249
    %251:i32 = load %250
    %indexCount:u32 = convert %251
    %253:ptr<storage, atomic<u32>, read_write> = access %drawOut, 1u
    %254:u32 = load %cubeVerts
    %255:u32 = atomicAdd %253, %254
    %firstVertex:ptr<function, u32, read_write> = var, %255
    %257:u32 = access %global_id, 0u
    %258:u32 = access %global_id, 1u
    %259:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %260:u32 = load_vector_element %259, 0u
    %261:u32 = mul %258, %260
    %262:u32 = add %257, %261
    %263:u32 = access %global_id, 2u
    %264:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %265:u32 = load_vector_element %264, 0u
    %266:u32 = mul %263, %265
    %267:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %268:u32 = load_vector_element %267, 1u
    %269:u32 = mul %266, %268
    %bufferOffset:u32 = add %262, %269
    %firstIndex:u32 = mul %bufferOffset, 15u
    loop [i: %b30, b: %b31, c: %b32] {  # loop_1
      %b30 = block {  # initializer
        %i_3:ptr<function, u32, read_write> = var, 0u  # %i_3: 'i'
        next_iteration %b31
      }
      %b31 = block {  # body
        %273:u32 = load %i_3
        %274:u32 = load %cubeVerts
        %275:bool = lt %273, %274
        if %275 [t: %b33, f: %b34] {  # if_22
          %b33 = block {  # true
            exit_if  # if_22
          }
          %b34 = block {  # false
            exit_loop  # loop_1
          }
        }
        %276:u32 = load %firstVertex
        %277:u32 = mul %276, 3u
        %278:u32 = load %i_3
        %279:u32 = mul %278, 3u
        %280:u32 = add %277, %279
        %281:ptr<storage, f32, read_write> = access %positionsOut, 0u, %280
        %282:u32 = load %i_3
        %283:ptr<private, vec3<f32>, read_write> = access %positions, %282
        %284:f32 = load_vector_element %283, 0u
        store %281, %284
        %285:u32 = load %firstVertex
        %286:u32 = mul %285, 3u
        %287:u32 = load %i_3
        %288:u32 = mul %287, 3u
        %289:u32 = add %286, %288
        %290:u32 = add %289, 1u
        %291:ptr<storage, f32, read_write> = access %positionsOut, 0u, %290
        %292:u32 = load %i_3
        %293:ptr<private, vec3<f32>, read_write> = access %positions, %292
        %294:f32 = load_vector_element %293, 1u
        store %291, %294
        %295:u32 = load %firstVertex
        %296:u32 = mul %295, 3u
        %297:u32 = load %i_3
        %298:u32 = mul %297, 3u
        %299:u32 = add %296, %298
        %300:u32 = add %299, 2u
        %301:ptr<storage, f32, read_write> = access %positionsOut, 0u, %300
        %302:u32 = load %i_3
        %303:ptr<private, vec3<f32>, read_write> = access %positions, %302
        %304:f32 = load_vector_element %303, 2u
        store %301, %304
        %305:u32 = load %firstVertex
        %306:u32 = mul %305, 3u
        %307:u32 = load %i_3
        %308:u32 = mul %307, 3u
        %309:u32 = add %306, %308
        %310:ptr<storage, f32, read_write> = access %normalsOut, 0u, %309
        %311:u32 = load %i_3
        %312:ptr<private, vec3<f32>, read_write> = access %normals, %311
        %313:f32 = load_vector_element %312, 0u
        store %310, %313
        %314:u32 = load %firstVertex
        %315:u32 = mul %314, 3u
        %316:u32 = load %i_3
        %317:u32 = mul %316, 3u
        %318:u32 = add %315, %317
        %319:u32 = add %318, 1u
        %320:ptr<storage, f32, read_write> = access %normalsOut, 0u, %319
        %321:u32 = load %i_3
        %322:ptr<private, vec3<f32>, read_write> = access %normals, %321
        %323:f32 = load_vector_element %322, 1u
        store %320, %323
        %324:u32 = load %firstVertex
        %325:u32 = mul %324, 3u
        %326:u32 = load %i_3
        %327:u32 = mul %326, 3u
        %328:u32 = add %325, %327
        %329:u32 = add %328, 2u
        %330:ptr<storage, f32, read_write> = access %normalsOut, 0u, %329
        %331:u32 = load %i_3
        %332:ptr<private, vec3<f32>, read_write> = access %normals, %331
        %333:f32 = load_vector_element %332, 2u
        store %330, %333
        continue %b32
      }
      %b32 = block {  # continuing
        %334:u32 = load %i_3
        %335:u32 = add %334, 1u
        store %i_3, %335
        next_iteration %b31
      }
    }
    loop [i: %b35, b: %b36, c: %b37] {  # loop_2
      %b35 = block {  # initializer
        %i_4:ptr<function, u32, read_write> = var, 0u  # %i_4: 'i'
        next_iteration %b36
      }
      %b36 = block {  # body
        %337:u32 = load %i_4
        %338:bool = lt %337, %indexCount
        if %338 [t: %b38, f: %b39] {  # if_23
          %b38 = block {  # true
            exit_if  # if_23
          }
          %b39 = block {  # false
            exit_loop  # loop_2
          }
        }
        %339:u32 = load %i_4
        %340:u32 = add %triTableOffset, %339
        %341:ptr<storage, i32, read> = access %tables, 1u, %340
        %index_6:i32 = load %341  # %index_6: 'index'
        %343:u32 = load %i_4
        %344:u32 = add %firstIndex, %343
        %345:ptr<storage, u32, read_write> = access %indicesOut, 0u, %344
        %346:u32 = load %firstVertex
        %347:ptr<private, u32, read_write> = access %indices, %index_6
        %348:u32 = load %347
        %349:u32 = add %346, %348
        store %345, %349
        continue %b37
      }
      %b37 = block {  # continuing
        %350:u32 = load %i_4
        %351:u32 = add %350, 1u
        store %i_4, %351
        next_iteration %b36
      }
    }
    loop [i: %b40, b: %b41, c: %b42] {  # loop_3
      %b40 = block {  # initializer
        %i_5:ptr<function, u32, read_write> = var, %indexCount  # %i_5: 'i'
        next_iteration %b41
      }
      %b41 = block {  # body
        %353:u32 = load %i_5
        %354:bool = lt %353, 15u
        if %354 [t: %b43, f: %b44] {  # if_24
          %b43 = block {  # true
            exit_if  # if_24
          }
          %b44 = block {  # false
            exit_loop  # loop_3
          }
        }
        %355:u32 = load %i_5
        %356:u32 = add %firstIndex, %355
        %357:ptr<storage, u32, read_write> = access %indicesOut, 0u, %356
        %358:u32 = load %firstVertex
        store %357, %358
        continue %b42
      }
      %b42 = block {  # continuing
        %359:u32 = load %i_5
        %360:u32 = add %359, 1u
        store %i_5, %360
        next_iteration %b41
      }
    }
    ret
  }
}

unhandled variable address space
********************************************************************
*  The tint shader compiler has encountered an unexpected error.   *
*                                                                  *
*  Please help us fix this issue by submitting a bug report at     *
*  crbug.com/tint with the source program that triggered the bug.  *
********************************************************************