SKIP: FAILED

../../src/tint/lang/msl/writer/printer/printer.cc:500 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: {  # 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: {
    %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: {  # 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
    %30:u32 = add %22, %29
    %valueIndex:u32 = let %30
    %32:ptr<storage, f32, read_write> = access %volume, 5u, %valueIndex
    %33:f32 = load %32
    ret %33
  }
}
%positionAt = func(%index_1:vec3<u32>):vec3<f32> {  # %index_1: 'index'
  $B4: {
    %36:ptr<storage, vec3<f32>, read_write> = access %volume, 0u
    %37:vec3<f32> = load %36
    %38:vec3<f32> = let %37
    %39:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %40:vec3<f32> = load %39
    %41:vec3<f32> = let %40
    %42:vec3<u32> = swizzle %index_1, xyz
    %43:vec3<f32> = convert %42
    %44:vec3<f32> = mul %41, %43
    %45:vec3<f32> = add %38, %44
    ret %45
  }
}
%normalAt = func(%index_2:vec3<u32>):vec3<f32> {  # %index_2: 'index'
  $B5: {
    %48:vec3<u32> = sub %index_2, vec3<u32>(1u, 0u, 0u)
    %49:f32 = call %valueAt, %48
    %50:f32 = let %49
    %51:vec3<u32> = add %index_2, vec3<u32>(1u, 0u, 0u)
    %52:f32 = call %valueAt, %51
    %53:f32 = sub %50, %52
    %54:f32 = let %53
    %55:vec3<u32> = sub %index_2, vec3<u32>(0u, 1u, 0u)
    %56:f32 = call %valueAt, %55
    %57:f32 = let %56
    %58:vec3<u32> = add %index_2, vec3<u32>(0u, 1u, 0u)
    %59:f32 = call %valueAt, %58
    %60:f32 = sub %57, %59
    %61:f32 = let %60
    %62:vec3<u32> = sub %index_2, vec3<u32>(0u, 0u, 1u)
    %63:f32 = call %valueAt, %62
    %64:f32 = let %63
    %65:vec3<u32> = add %index_2, vec3<u32>(0u, 0u, 1u)
    %66:f32 = call %valueAt, %65
    %67:f32 = sub %64, %66
    %68:vec3<f32> = construct %54, %61, %67
    ret %68
  }
}
%interpX = func(%index_3:u32, %i:vec3<u32>, %va:f32, %vb:f32):void {  # %index_3: 'index'
  $B6: {
    %74:ptr<storage, f32, read_write> = access %volume, 4u
    %75:f32 = load %74
    %76:f32 = sub %75, %va
    %77:f32 = sub %vb, %va
    %78:f32 = div %76, %77
    %mu:f32 = let %78
    %80:u32 = load %cubeVerts
    %81:ptr<private, vec3<f32>, read_write> = access %positions, %80
    %82:ptr<private, vec3<f32>, read_write> = let %81
    %83:vec3<f32> = call %positionAt, %i
    %84:vec3<f32> = let %83
    %85:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %86:f32 = load_vector_element %85, 0u
    %87:f32 = mul %86, %mu
    %88:vec3<f32> = construct %87, 0.0f, 0.0f
    %89:vec3<f32> = add %84, %88
    store %82, %89
    %90:vec3<f32> = call %normalAt, %i
    %na:vec3<f32> = let %90
    %92:vec3<u32> = add %i, vec3<u32>(1u, 0u, 0u)
    %93:vec3<f32> = call %normalAt, %92
    %nb:vec3<f32> = let %93
    %95:u32 = load %cubeVerts
    %96:ptr<private, vec3<f32>, read_write> = access %normals, %95
    %97:ptr<private, vec3<f32>, read_write> = let %96
    %98:vec3<f32> = construct %mu, %mu, %mu
    %99:vec3<f32> = mix %na, %nb, %98
    store %97, %99
    %100:ptr<private, u32, read_write> = access %indices, %index_3
    %101:u32 = load %cubeVerts
    store %100, %101
    %102:u32 = load %cubeVerts
    %103:u32 = add %102, 1u
    store %cubeVerts, %103
    ret
  }
}
%interpY = func(%index_4:u32, %i_1:vec3<u32>, %va_1:f32, %vb_1:f32):void {  # %index_4: 'index', %i_1: 'i', %va_1: 'va', %vb_1: 'vb'
  $B7: {
    %109:ptr<storage, f32, read_write> = access %volume, 4u
    %110:f32 = load %109
    %111:f32 = sub %110, %va_1
    %112:f32 = sub %vb_1, %va_1
    %113:f32 = div %111, %112
    %mu_1:f32 = let %113  # %mu_1: 'mu'
    %115:u32 = load %cubeVerts
    %116:ptr<private, vec3<f32>, read_write> = access %positions, %115
    %117:ptr<private, vec3<f32>, read_write> = let %116
    %118:vec3<f32> = call %positionAt, %i_1
    %119:vec3<f32> = let %118
    %120:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %121:f32 = load_vector_element %120, 1u
    %122:f32 = mul %121, %mu_1
    %123:vec3<f32> = construct 0.0f, %122, 0.0f
    %124:vec3<f32> = add %119, %123
    store %117, %124
    %125:vec3<f32> = call %normalAt, %i_1
    %na_1:vec3<f32> = let %125  # %na_1: 'na'
    %127:vec3<u32> = add %i_1, vec3<u32>(0u, 1u, 0u)
    %128:vec3<f32> = call %normalAt, %127
    %nb_1:vec3<f32> = let %128  # %nb_1: 'nb'
    %130:u32 = load %cubeVerts
    %131:ptr<private, vec3<f32>, read_write> = access %normals, %130
    %132:ptr<private, vec3<f32>, read_write> = let %131
    %133:vec3<f32> = construct %mu_1, %mu_1, %mu_1
    %134:vec3<f32> = mix %na_1, %nb_1, %133
    store %132, %134
    %135:ptr<private, u32, read_write> = access %indices, %index_4
    %136:u32 = load %cubeVerts
    store %135, %136
    %137:u32 = load %cubeVerts
    %138:u32 = add %137, 1u
    store %cubeVerts, %138
    ret
  }
}
%interpZ = func(%index_5:u32, %i_2:vec3<u32>, %va_2:f32, %vb_2:f32):void {  # %index_5: 'index', %i_2: 'i', %va_2: 'va', %vb_2: 'vb'
  $B8: {
    %144:ptr<storage, f32, read_write> = access %volume, 4u
    %145:f32 = load %144
    %146:f32 = sub %145, %va_2
    %147:f32 = sub %vb_2, %va_2
    %148:f32 = div %146, %147
    %mu_2:f32 = let %148  # %mu_2: 'mu'
    %150:u32 = load %cubeVerts
    %151:ptr<private, vec3<f32>, read_write> = access %positions, %150
    %152:ptr<private, vec3<f32>, read_write> = let %151
    %153:vec3<f32> = call %positionAt, %i_2
    %154:vec3<f32> = let %153
    %155:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
    %156:f32 = load_vector_element %155, 2u
    %157:f32 = mul %156, %mu_2
    %158:vec3<f32> = construct 0.0f, 0.0f, %157
    %159:vec3<f32> = add %154, %158
    store %152, %159
    %160:vec3<f32> = call %normalAt, %i_2
    %na_2:vec3<f32> = let %160  # %na_2: 'na'
    %162:vec3<u32> = add %i_2, vec3<u32>(0u, 0u, 1u)
    %163:vec3<f32> = call %normalAt, %162
    %nb_2:vec3<f32> = let %163  # %nb_2: 'nb'
    %165:u32 = load %cubeVerts
    %166:ptr<private, vec3<f32>, read_write> = access %normals, %165
    %167:ptr<private, vec3<f32>, read_write> = let %166
    %168:vec3<f32> = construct %mu_2, %mu_2, %mu_2
    %169:vec3<f32> = mix %na_2, %nb_2, %168
    store %167, %169
    %170:ptr<private, u32, read_write> = access %indices, %index_5
    %171:u32 = load %cubeVerts
    store %170, %171
    %172:u32 = load %cubeVerts
    %173:u32 = add %172, 1u
    store %cubeVerts, %173
    ret
  }
}
%computeMain = @compute @workgroup_size(4, 4, 4) func(%global_id:vec3<u32> [@global_invocation_id]):void {
  $B9: {
    %i0:vec3<u32> = let %global_id
    %177:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 0u)
    %i1:vec3<u32> = let %177
    %179:vec3<u32> = add %global_id, vec3<u32>(1u, 1u, 0u)
    %i2:vec3<u32> = let %179
    %181:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 0u)
    %i3:vec3<u32> = let %181
    %183:vec3<u32> = add %global_id, vec3<u32>(0u, 0u, 1u)
    %i4:vec3<u32> = let %183
    %185:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 1u)
    %i5:vec3<u32> = let %185
    %187:vec3<u32> = add %global_id, vec3<u32>(1u)
    %i6:vec3<u32> = let %187
    %189:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 1u)
    %i7:vec3<u32> = let %189
    %191:f32 = call %valueAt, %i0
    %v0:f32 = let %191
    %193:f32 = call %valueAt, %i1
    %v1:f32 = let %193
    %195:f32 = call %valueAt, %i2
    %v2:f32 = let %195
    %197:f32 = call %valueAt, %i3
    %v3:f32 = let %197
    %199:f32 = call %valueAt, %i4
    %v4:f32 = let %199
    %201:f32 = call %valueAt, %i5
    %v5:f32 = let %201
    %203:f32 = call %valueAt, %i6
    %v6:f32 = let %203
    %205:f32 = call %valueAt, %i7
    %v7:f32 = let %205
    %cubeIndex:ptr<function, u32, read_write> = var, 0u
    %208:ptr<storage, f32, read_write> = access %volume, 4u
    %209:f32 = load %208
    %210:bool = lt %v0, %209
    if %210 [t: $B10] {  # if_2
      $B10: {  # true
        %211:u32 = load %cubeIndex
        %212:u32 = or %211, 1u
        store %cubeIndex, %212
        exit_if  # if_2
      }
    }
    %213:ptr<storage, f32, read_write> = access %volume, 4u
    %214:f32 = load %213
    %215:bool = lt %v1, %214
    if %215 [t: $B11] {  # if_3
      $B11: {  # true
        %216:u32 = load %cubeIndex
        %217:u32 = or %216, 2u
        store %cubeIndex, %217
        exit_if  # if_3
      }
    }
    %218:ptr<storage, f32, read_write> = access %volume, 4u
    %219:f32 = load %218
    %220:bool = lt %v2, %219
    if %220 [t: $B12] {  # if_4
      $B12: {  # true
        %221:u32 = load %cubeIndex
        %222:u32 = or %221, 4u
        store %cubeIndex, %222
        exit_if  # if_4
      }
    }
    %223:ptr<storage, f32, read_write> = access %volume, 4u
    %224:f32 = load %223
    %225:bool = lt %v3, %224
    if %225 [t: $B13] {  # if_5
      $B13: {  # true
        %226:u32 = load %cubeIndex
        %227:u32 = or %226, 8u
        store %cubeIndex, %227
        exit_if  # if_5
      }
    }
    %228:ptr<storage, f32, read_write> = access %volume, 4u
    %229:f32 = load %228
    %230:bool = lt %v4, %229
    if %230 [t: $B14] {  # if_6
      $B14: {  # true
        %231:u32 = load %cubeIndex
        %232:u32 = or %231, 16u
        store %cubeIndex, %232
        exit_if  # if_6
      }
    }
    %233:ptr<storage, f32, read_write> = access %volume, 4u
    %234:f32 = load %233
    %235:bool = lt %v5, %234
    if %235 [t: $B15] {  # if_7
      $B15: {  # true
        %236:u32 = load %cubeIndex
        %237:u32 = or %236, 32u
        store %cubeIndex, %237
        exit_if  # if_7
      }
    }
    %238:ptr<storage, f32, read_write> = access %volume, 4u
    %239:f32 = load %238
    %240:bool = lt %v6, %239
    if %240 [t: $B16] {  # if_8
      $B16: {  # true
        %241:u32 = load %cubeIndex
        %242:u32 = or %241, 64u
        store %cubeIndex, %242
        exit_if  # if_8
      }
    }
    %243:ptr<storage, f32, read_write> = access %volume, 4u
    %244:f32 = load %243
    %245:bool = lt %v7, %244
    if %245 [t: $B17] {  # if_9
      $B17: {  # true
        %246:u32 = load %cubeIndex
        %247:u32 = or %246, 128u
        store %cubeIndex, %247
        exit_if  # if_9
      }
    }
    %248:u32 = load %cubeIndex
    %249:ptr<storage, u32, read> = access %tables, 0u, %248
    %250:u32 = load %249
    %edges:u32 = let %250
    %252:u32 = and %edges, 1u
    %253:bool = neq %252, 0u
    if %253 [t: $B18] {  # if_10
      $B18: {  # true
        %254:void = call %interpX, 0u, %i0, %v0, %v1
        exit_if  # if_10
      }
    }
    %255:u32 = and %edges, 2u
    %256:bool = neq %255, 0u
    if %256 [t: $B19] {  # if_11
      $B19: {  # true
        %257:void = call %interpY, 1u, %i1, %v1, %v2
        exit_if  # if_11
      }
    }
    %258:u32 = and %edges, 4u
    %259:bool = neq %258, 0u
    if %259 [t: $B20] {  # if_12
      $B20: {  # true
        %260:void = call %interpX, 2u, %i3, %v3, %v2
        exit_if  # if_12
      }
    }
    %261:u32 = and %edges, 8u
    %262:bool = neq %261, 0u
    if %262 [t: $B21] {  # if_13
      $B21: {  # true
        %263:void = call %interpY, 3u, %i0, %v0, %v3
        exit_if  # if_13
      }
    }
    %264:u32 = and %edges, 16u
    %265:bool = neq %264, 0u
    if %265 [t: $B22] {  # if_14
      $B22: {  # true
        %266:void = call %interpX, 4u, %i4, %v4, %v5
        exit_if  # if_14
      }
    }
    %267:u32 = and %edges, 32u
    %268:bool = neq %267, 0u
    if %268 [t: $B23] {  # if_15
      $B23: {  # true
        %269:void = call %interpY, 5u, %i5, %v5, %v6
        exit_if  # if_15
      }
    }
    %270:u32 = and %edges, 64u
    %271:bool = neq %270, 0u
    if %271 [t: $B24] {  # if_16
      $B24: {  # true
        %272:void = call %interpX, 6u, %i7, %v7, %v6
        exit_if  # if_16
      }
    }
    %273:u32 = and %edges, 128u
    %274:bool = neq %273, 0u
    if %274 [t: $B25] {  # if_17
      $B25: {  # true
        %275:void = call %interpY, 7u, %i4, %v4, %v7
        exit_if  # if_17
      }
    }
    %276:u32 = and %edges, 256u
    %277:bool = neq %276, 0u
    if %277 [t: $B26] {  # if_18
      $B26: {  # true
        %278:void = call %interpZ, 8u, %i0, %v0, %v4
        exit_if  # if_18
      }
    }
    %279:u32 = and %edges, 512u
    %280:bool = neq %279, 0u
    if %280 [t: $B27] {  # if_19
      $B27: {  # true
        %281:void = call %interpZ, 9u, %i1, %v1, %v5
        exit_if  # if_19
      }
    }
    %282:u32 = and %edges, 1024u
    %283:bool = neq %282, 0u
    if %283 [t: $B28] {  # if_20
      $B28: {  # true
        %284:void = call %interpZ, 10u, %i2, %v2, %v6
        exit_if  # if_20
      }
    }
    %285:u32 = and %edges, 2048u
    %286:bool = neq %285, 0u
    if %286 [t: $B29] {  # if_21
      $B29: {  # true
        %287:void = call %interpZ, 11u, %i3, %v3, %v7
        exit_if  # if_21
      }
    }
    %288:u32 = load %cubeIndex
    %289:u32 = and 4u, 31u
    %290:u32 = shl %288, %289
    %291:u32 = add %290, 1u
    %triTableOffset:u32 = let %291
    %293:u32 = sub %triTableOffset, 1u
    %294:ptr<storage, i32, read> = access %tables, 1u, %293
    %295:i32 = load %294
    %296:u32 = convert %295
    %indexCount:u32 = let %296
    %298:ptr<storage, atomic<u32>, read_write> = access %drawOut, 1u
    %299:u32 = load %cubeVerts
    %300:u32 = atomicAdd %298, %299
    %firstVertex:ptr<function, u32, read_write> = var, %300
    %302:u32 = access %global_id, 0u
    %303:u32 = access %global_id, 1u
    %304:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %305:u32 = load_vector_element %304, 0u
    %306:u32 = mul %303, %305
    %307:u32 = add %302, %306
    %308:u32 = access %global_id, 2u
    %309:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %310:u32 = load_vector_element %309, 0u
    %311:u32 = mul %308, %310
    %312:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
    %313:u32 = load_vector_element %312, 1u
    %314:u32 = mul %311, %313
    %315:u32 = add %307, %314
    %bufferOffset:u32 = let %315
    %317:u32 = mul %bufferOffset, 15u
    %firstIndex:u32 = let %317
    loop [i: $B30, b: $B31, c: $B32] {  # loop_1
      $B30: {  # initializer
        %i_3:ptr<function, u32, read_write> = var, 0u  # %i_3: 'i'
        next_iteration  # -> $B31
      }
      $B31: {  # body
        %320:u32 = load %i_3
        %321:u32 = load %cubeVerts
        %322:bool = lt %320, %321
        if %322 [t: $B33, f: $B34] {  # if_22
          $B33: {  # true
            exit_if  # if_22
          }
          $B34: {  # false
            exit_loop  # loop_1
          }
        }
        %323:u32 = load %firstVertex
        %324:u32 = mul %323, 3u
        %325:u32 = load %i_3
        %326:u32 = mul %325, 3u
        %327:u32 = add %324, %326
        %328:ptr<storage, f32, read_write> = access %positionsOut, 0u, %327
        %329:u32 = load %i_3
        %330:ptr<private, vec3<f32>, read_write> = access %positions, %329
        %331:f32 = load_vector_element %330, 0u
        store %328, %331
        %332:u32 = load %firstVertex
        %333:u32 = mul %332, 3u
        %334:u32 = load %i_3
        %335:u32 = mul %334, 3u
        %336:u32 = add %333, %335
        %337:u32 = add %336, 1u
        %338:ptr<storage, f32, read_write> = access %positionsOut, 0u, %337
        %339:u32 = load %i_3
        %340:ptr<private, vec3<f32>, read_write> = access %positions, %339
        %341:f32 = load_vector_element %340, 1u
        store %338, %341
        %342:u32 = load %firstVertex
        %343:u32 = mul %342, 3u
        %344:u32 = load %i_3
        %345:u32 = mul %344, 3u
        %346:u32 = add %343, %345
        %347:u32 = add %346, 2u
        %348:ptr<storage, f32, read_write> = access %positionsOut, 0u, %347
        %349:u32 = load %i_3
        %350:ptr<private, vec3<f32>, read_write> = access %positions, %349
        %351:f32 = load_vector_element %350, 2u
        store %348, %351
        %352:u32 = load %firstVertex
        %353:u32 = mul %352, 3u
        %354:u32 = load %i_3
        %355:u32 = mul %354, 3u
        %356:u32 = add %353, %355
        %357:ptr<storage, f32, read_write> = access %normalsOut, 0u, %356
        %358:u32 = load %i_3
        %359:ptr<private, vec3<f32>, read_write> = access %normals, %358
        %360:f32 = load_vector_element %359, 0u
        store %357, %360
        %361:u32 = load %firstVertex
        %362:u32 = mul %361, 3u
        %363:u32 = load %i_3
        %364:u32 = mul %363, 3u
        %365:u32 = add %362, %364
        %366:u32 = add %365, 1u
        %367:ptr<storage, f32, read_write> = access %normalsOut, 0u, %366
        %368:u32 = load %i_3
        %369:ptr<private, vec3<f32>, read_write> = access %normals, %368
        %370:f32 = load_vector_element %369, 1u
        store %367, %370
        %371:u32 = load %firstVertex
        %372:u32 = mul %371, 3u
        %373:u32 = load %i_3
        %374:u32 = mul %373, 3u
        %375:u32 = add %372, %374
        %376:u32 = add %375, 2u
        %377:ptr<storage, f32, read_write> = access %normalsOut, 0u, %376
        %378:u32 = load %i_3
        %379:ptr<private, vec3<f32>, read_write> = access %normals, %378
        %380:f32 = load_vector_element %379, 2u
        store %377, %380
        continue  # -> $B32
      }
      $B32: {  # continuing
        %381:u32 = load %i_3
        %382:u32 = add %381, 1u
        store %i_3, %382
        next_iteration  # -> $B31
      }
    }
    loop [i: $B35, b: $B36, c: $B37] {  # loop_2
      $B35: {  # initializer
        %i_4:ptr<function, u32, read_write> = var, 0u  # %i_4: 'i'
        next_iteration  # -> $B36
      }
      $B36: {  # body
        %384:u32 = load %i_4
        %385:bool = lt %384, %indexCount
        if %385 [t: $B38, f: $B39] {  # if_23
          $B38: {  # true
            exit_if  # if_23
          }
          $B39: {  # false
            exit_loop  # loop_2
          }
        }
        %386:u32 = load %i_4
        %387:u32 = add %triTableOffset, %386
        %388:ptr<storage, i32, read> = access %tables, 1u, %387
        %389:i32 = load %388
        %index_6:i32 = let %389  # %index_6: 'index'
        %391:u32 = load %i_4
        %392:u32 = add %firstIndex, %391
        %393:ptr<storage, u32, read_write> = access %indicesOut, 0u, %392
        %394:u32 = load %firstVertex
        %395:ptr<private, u32, read_write> = access %indices, %index_6
        %396:u32 = load %395
        %397:u32 = add %394, %396
        store %393, %397
        continue  # -> $B37
      }
      $B37: {  # continuing
        %398:u32 = load %i_4
        %399:u32 = add %398, 1u
        store %i_4, %399
        next_iteration  # -> $B36
      }
    }
    loop [i: $B40, b: $B41, c: $B42] {  # loop_3
      $B40: {  # initializer
        %i_5:ptr<function, u32, read_write> = var, %indexCount  # %i_5: 'i'
        next_iteration  # -> $B41
      }
      $B41: {  # body
        %401:u32 = load %i_5
        %402:bool = lt %401, 15u
        if %402 [t: $B43, f: $B44] {  # if_24
          $B43: {  # true
            exit_if  # if_24
          }
          $B44: {  # false
            exit_loop  # loop_3
          }
        }
        %403:u32 = load %i_5
        %404:u32 = add %firstIndex, %403
        %405:ptr<storage, u32, read_write> = access %indicesOut, 0u, %404
        %406:u32 = load %firstVertex
        store %405, %406
        continue  # -> $B42
      }
      $B42: {  # continuing
        %407:u32 = load %i_5
        %408:u32 = add %407, 1u
        store %i_5, %408
        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.  *
********************************************************************