SKIP: FAILED

../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: LightData = struct @align(16) {
  position:vec4<f32> @offset(0)
  color:vec3<f32> @offset(16)
  radius:f32 @offset(28)
}

LightsBuffer = struct @align(16) {
  lights:array<LightData> @offset(0)
}

TileLightIdData = struct @align(4) {
  count:atomic<u32> @offset(0)
  lightId:array<u32, 64> @offset(4)
}

Tiles = struct @align(4) {
  data:array<TileLightIdData, 4> @offset(0)
}

Config = struct @align(4) {
  numLights:u32 @offset(0)
  numTiles:u32 @offset(4)
  tileCountX:u32 @offset(8)
  tileCountY:u32 @offset(12)
  numTileLightSlot:u32 @offset(16)
  tileSize:u32 @offset(20)
}

Uniforms = struct @align(16) {
  min:vec4<f32> @offset(0)
  max:vec4<f32> @offset(16)
  viewMatrix:mat4x4<f32> @offset(32)
  projectionMatrix:mat4x4<f32> @offset(96)
  fullScreenSize:vec4<f32> @offset(160)
}

$B1: {  # root
  %lightsBuffer:ptr<storage, LightsBuffer, read_write> = var @binding_point(0, 0)
  %tileLightId:ptr<storage, Tiles, read_write> = var @binding_point(1, 0)
  %config:ptr<uniform, Config, read> = var @binding_point(2, 0)
  %uniforms:ptr<uniform, Uniforms, read> = var @binding_point(3, 0)
}

%tint_symbol = @compute @workgroup_size(64, 1, 1) func(%GlobalInvocationID:vec3<u32> [@global_invocation_id]):void {
  $B2: {
    %7:u32 = access %GlobalInvocationID, 0u
    %index:ptr<function, u32, read_write> = var, %7
    %9:u32 = load %index
    %10:ptr<uniform, u32, read> = access %config, 0u
    %11:u32 = load %10
    %12:bool = gte %9, %11
    if %12 [t: $B3] {  # if_1
      $B3: {  # true
        ret
      }
    }
    %13:u32 = load %index
    %14:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %13, 0u
    %15:ptr<storage, vec4<f32>, read_write> = let %14
    %16:u32 = load %index
    %17:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %16, 0u
    %18:f32 = load_vector_element %17, 1u
    %19:f32 = sub %18, 0.10000000149011611938f
    %20:f32 = let %19
    %21:u32 = load %index
    %22:f32 = convert %21
    %23:f32 = let %22
    %24:u32 = load %index
    %25:f32 = convert %24
    %26:f32 = div %25, 64.0f
    %27:f32 = floor %26
    %28:f32 = mul 64.0f, %27
    %29:f32 = sub %23, %28
    %30:f32 = mul 0.00100000004749745131f, %29
    %31:f32 = add %20, %30
    store_vector_element %15, 1u, %31
    %32:u32 = load %index
    %33:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %32, 0u
    %34:f32 = load_vector_element %33, 1u
    %35:ptr<uniform, vec4<f32>, read> = access %uniforms, 0u
    %36:f32 = load_vector_element %35, 1u
    %37:bool = lt %34, %36
    if %37 [t: $B4] {  # if_2
      $B4: {  # true
        %38:u32 = load %index
        %39:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %38, 0u
        %40:ptr<uniform, vec4<f32>, read> = access %uniforms, 1u
        %41:f32 = load_vector_element %40, 1u
        store_vector_element %39, 1u, %41
        exit_if  # if_2
      }
    }
    %42:ptr<uniform, mat4x4<f32>, read> = access %uniforms, 3u
    %43:mat4x4<f32> = load %42
    %M:ptr<function, mat4x4<f32>, read_write> = var, %43
    %45:ptr<function, vec4<f32>, read_write> = access %M, 3i
    %46:f32 = load_vector_element %45, 2i
    %47:f32 = negation %46
    %48:ptr<function, vec4<f32>, read_write> = access %M, 2i
    %49:f32 = load_vector_element %48, 2i
    %50:f32 = add -1.0f, %49
    %51:f32 = div %47, %50
    %viewNear:ptr<function, f32, read_write> = var, %51
    %53:ptr<function, vec4<f32>, read_write> = access %M, 3i
    %54:f32 = load_vector_element %53, 2i
    %55:f32 = negation %54
    %56:ptr<function, vec4<f32>, read_write> = access %M, 2i
    %57:f32 = load_vector_element %56, 2i
    %58:f32 = add 1.0f, %57
    %59:f32 = div %55, %58
    %viewFar:ptr<function, f32, read_write> = var, %59
    %61:u32 = load %index
    %62:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %61, 0u
    %63:vec4<f32> = load %62
    %lightPos:ptr<function, vec4<f32>, read_write> = var, %63
    %65:ptr<uniform, mat4x4<f32>, read> = access %uniforms, 2u
    %66:mat4x4<f32> = load %65
    %67:vec4<f32> = load %lightPos
    %68:vec4<f32> = mul %66, %67
    store %lightPos, %68
    %69:vec4<f32> = load %lightPos
    %70:f32 = load_vector_element %lightPos, 3u
    %71:vec4<f32> = div %69, %70
    store %lightPos, %71
    %72:u32 = load %index
    %73:ptr<storage, f32, read_write> = access %lightsBuffer, 0u, %72, 2u
    %74:f32 = load %73
    %lightRadius:ptr<function, f32, read_write> = var, %74
    %76:vec4<f32> = load %lightPos
    %77:vec4<f32> = let %76
    %78:f32 = load %lightRadius
    %79:vec3<f32> = construct %78
    %80:vec4<f32> = construct %79, 0.0f
    %81:vec4<f32> = sub %77, %80
    %boxMin:ptr<function, vec4<f32>, read_write> = var, %81
    %83:vec4<f32> = load %lightPos
    %84:vec4<f32> = let %83
    %85:f32 = load %lightRadius
    %86:vec3<f32> = construct %85
    %87:vec4<f32> = construct %86, 0.0f
    %88:vec4<f32> = add %84, %87
    %boxMax:ptr<function, vec4<f32>, read_write> = var, %88
    %frustumPlanes:ptr<function, array<vec4<f32>, 6>, read_write> = var
    %91:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 4i
    %92:f32 = load %viewNear
    %93:vec4<f32> = construct 0.0f, 0.0f, -1.0f, %92
    store %91, %93
    %94:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 5i
    %95:f32 = load %viewFar
    %96:f32 = negation %95
    %97:vec4<f32> = construct 0.0f, 0.0f, 1.0f, %96
    store %94, %97
    %TILE_SIZE:i32 = let 16i
    %TILE_COUNT_X:i32 = let 2i
    %TILE_COUNT_Y:i32 = let 2i
    loop [i: $B5, b: $B6, c: $B7] {  # loop_1
      $B5: {  # initializer
        %y:ptr<function, i32, read_write> = var, 0i
        next_iteration  # -> $B6
      }
      $B6: {  # body
        %102:i32 = load %y
        %103:bool = lt %102, %TILE_COUNT_Y
        if %103 [t: $B8, f: $B9] {  # if_3
          $B8: {  # true
            exit_if  # if_3
          }
          $B9: {  # false
            exit_loop  # loop_1
          }
        }
        loop [i: $B10, b: $B11, c: $B12] {  # loop_2
          $B10: {  # initializer
            %x:ptr<function, i32, read_write> = var, 0i
            next_iteration  # -> $B11
          }
          $B11: {  # body
            %105:i32 = load %x
            %106:bool = lt %105, %TILE_COUNT_X
            if %106 [t: $B13, f: $B14] {  # if_4
              $B13: {  # true
                exit_if  # if_4
              }
              $B14: {  # false
                exit_loop  # loop_2
              }
            }
            %107:i32 = load %x
            %108:i32 = mul %107, %TILE_SIZE
            %109:i32 = load %y
            %110:i32 = mul %109, %TILE_SIZE
            %111:vec2<i32> = construct %108, %110
            %tilePixel0Idx:ptr<function, vec2<i32>, read_write> = var, %111
            %113:vec2<i32> = load %tilePixel0Idx
            %114:vec2<f32> = convert %113
            %115:vec2<f32> = mul 2.0f, %114
            %116:ptr<uniform, vec4<f32>, read> = access %uniforms, 4u
            %117:vec4<f32> = load %116
            %118:vec2<f32> = swizzle %117, xy
            %119:vec2<f32> = div %115, %118
            %120:vec2<f32> = sub %119, vec2<f32>(1.0f)
            %floorCoord:ptr<function, vec2<f32>, read_write> = var, %120
            %122:vec2<i32> = load %tilePixel0Idx
            %123:vec2<i32> = let %122
            %124:vec2<i32> = construct %TILE_SIZE
            %125:vec2<i32> = add %123, %124
            %126:vec2<f32> = convert %125
            %127:vec2<f32> = mul 2.0f, %126
            %128:ptr<uniform, vec4<f32>, read> = access %uniforms, 4u
            %129:vec4<f32> = load %128
            %130:vec2<f32> = swizzle %129, xy
            %131:vec2<f32> = div %127, %130
            %132:vec2<f32> = sub %131, vec2<f32>(1.0f)
            %ceilCoord:ptr<function, vec2<f32>, read_write> = var, %132
            %134:f32 = load %viewNear
            %135:f32 = negation %134
            %136:f32 = load_vector_element %floorCoord, 0u
            %137:f32 = mul %135, %136
            %138:ptr<function, vec4<f32>, read_write> = access %M, 2i
            %139:f32 = load_vector_element %138, 0i
            %140:f32 = load %viewNear
            %141:f32 = mul %139, %140
            %142:f32 = sub %137, %141
            %143:ptr<function, vec4<f32>, read_write> = access %M, 0i
            %144:f32 = load_vector_element %143, 0i
            %145:f32 = div %142, %144
            %146:f32 = load %viewNear
            %147:f32 = negation %146
            %148:f32 = load_vector_element %floorCoord, 1u
            %149:f32 = mul %147, %148
            %150:ptr<function, vec4<f32>, read_write> = access %M, 2i
            %151:f32 = load_vector_element %150, 1i
            %152:f32 = load %viewNear
            %153:f32 = mul %151, %152
            %154:f32 = sub %149, %153
            %155:ptr<function, vec4<f32>, read_write> = access %M, 1i
            %156:f32 = load_vector_element %155, 1i
            %157:f32 = div %154, %156
            %158:vec2<f32> = construct %145, %157
            %viewFloorCoord:ptr<function, vec2<f32>, read_write> = var, %158
            %160:f32 = load %viewNear
            %161:f32 = negation %160
            %162:f32 = load_vector_element %ceilCoord, 0u
            %163:f32 = mul %161, %162
            %164:ptr<function, vec4<f32>, read_write> = access %M, 2i
            %165:f32 = load_vector_element %164, 0i
            %166:f32 = load %viewNear
            %167:f32 = mul %165, %166
            %168:f32 = sub %163, %167
            %169:ptr<function, vec4<f32>, read_write> = access %M, 0i
            %170:f32 = load_vector_element %169, 0i
            %171:f32 = div %168, %170
            %172:f32 = load %viewNear
            %173:f32 = negation %172
            %174:f32 = load_vector_element %ceilCoord, 1u
            %175:f32 = mul %173, %174
            %176:ptr<function, vec4<f32>, read_write> = access %M, 2i
            %177:f32 = load_vector_element %176, 1i
            %178:f32 = load %viewNear
            %179:f32 = mul %177, %178
            %180:f32 = sub %175, %179
            %181:ptr<function, vec4<f32>, read_write> = access %M, 1i
            %182:f32 = load_vector_element %181, 1i
            %183:f32 = div %180, %182
            %184:vec2<f32> = construct %171, %183
            %viewCeilCoord:ptr<function, vec2<f32>, read_write> = var, %184
            %186:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 0i
            %187:f32 = load_vector_element %viewFloorCoord, 0u
            %188:f32 = negation %187
            %189:f32 = load %viewNear
            %190:f32 = div %188, %189
            %191:vec4<f32> = construct 1.0f, 0.0f, %190, 0.0f
            store %186, %191
            %192:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 1i
            %193:f32 = load_vector_element %viewCeilCoord, 0u
            %194:f32 = load %viewNear
            %195:f32 = div %193, %194
            %196:vec4<f32> = construct -1.0f, 0.0f, %195, 0.0f
            store %192, %196
            %197:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 2i
            %198:f32 = load_vector_element %viewFloorCoord, 1u
            %199:f32 = negation %198
            %200:f32 = load %viewNear
            %201:f32 = div %199, %200
            %202:vec4<f32> = construct 0.0f, 1.0f, %201, 0.0f
            store %197, %202
            %203:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 3i
            %204:f32 = load_vector_element %viewCeilCoord, 1u
            %205:f32 = load %viewNear
            %206:f32 = div %204, %205
            %207:vec4<f32> = construct 0.0f, -1.0f, %206, 0.0f
            store %203, %207
            %dp:ptr<function, f32, read_write> = var, 0.0f
            loop [i: $B15, b: $B16, c: $B17] {  # loop_3
              $B15: {  # initializer
                %i:ptr<function, u32, read_write> = var, 0u
                next_iteration  # -> $B16
              }
              $B16: {  # body
                %210:u32 = load %i
                %211:bool = lt %210, 6u
                if %211 [t: $B18, f: $B19] {  # if_5
                  $B18: {  # true
                    exit_if  # if_5
                  }
                  $B19: {  # false
                    exit_loop  # loop_3
                  }
                }
                %p:ptr<function, vec4<f32>, read_write> = var
                %213:u32 = load %i
                %214:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %213
                %215:f32 = load_vector_element %214, 0u
                %216:bool = gt %215, 0.0f
                if %216 [t: $B20, f: $B21] {  # if_6
                  $B20: {  # true
                    %217:f32 = load_vector_element %boxMax, 0u
                    store_vector_element %p, 0u, %217
                    exit_if  # if_6
                  }
                  $B21: {  # false
                    %218:f32 = load_vector_element %boxMin, 0u
                    store_vector_element %p, 0u, %218
                    exit_if  # if_6
                  }
                }
                %219:u32 = load %i
                %220:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %219
                %221:f32 = load_vector_element %220, 1u
                %222:bool = gt %221, 0.0f
                if %222 [t: $B22, f: $B23] {  # if_7
                  $B22: {  # true
                    %223:f32 = load_vector_element %boxMax, 1u
                    store_vector_element %p, 1u, %223
                    exit_if  # if_7
                  }
                  $B23: {  # false
                    %224:f32 = load_vector_element %boxMin, 1u
                    store_vector_element %p, 1u, %224
                    exit_if  # if_7
                  }
                }
                %225:u32 = load %i
                %226:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %225
                %227:f32 = load_vector_element %226, 2u
                %228:bool = gt %227, 0.0f
                if %228 [t: $B24, f: $B25] {  # if_8
                  $B24: {  # true
                    %229:f32 = load_vector_element %boxMax, 2u
                    store_vector_element %p, 2u, %229
                    exit_if  # if_8
                  }
                  $B25: {  # false
                    %230:f32 = load_vector_element %boxMin, 2u
                    store_vector_element %p, 2u, %230
                    exit_if  # if_8
                  }
                }
                store_vector_element %p, 3u, 1.0f
                %231:f32 = load %dp
                %232:f32 = let %231
                %233:vec4<f32> = load %p
                %234:u32 = load %i
                %235:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %234
                %236:vec4<f32> = load %235
                %237:f32 = dot %233, %236
                %238:f32 = min 0.0f, %237
                %239:f32 = add %232, %238
                store %dp, %239
                continue  # -> $B17
              }
              $B17: {  # continuing
                %240:u32 = load %i
                %241:u32 = add %240, 1u
                store %i, %241
                next_iteration  # -> $B16
              }
            }
            %242:f32 = load %dp
            %243:bool = gte %242, 0.0f
            if %243 [t: $B26] {  # if_9
              $B26: {  # true
                %244:i32 = load %x
                %245:i32 = load %y
                %246:i32 = mul %245, %TILE_COUNT_X
                %247:i32 = add %244, %246
                %248:u32 = convert %247
                %tileId:ptr<function, u32, read_write> = var, %248
                %250:u32 = load %tileId
                %251:bool = lt %250, 0u
                %252:bool = if %251 [t: $B27, f: $B28] {  # if_10
                  $B27: {  # true
                    exit_if true  # if_10
                  }
                  $B28: {  # false
                    %253:u32 = load %tileId
                    %254:ptr<uniform, u32, read> = access %config, 1u
                    %255:u32 = load %254
                    %256:bool = gte %253, %255
                    exit_if %256  # if_10
                  }
                }
                if %252 [t: $B29] {  # if_11
                  $B29: {  # true
                    continue  # -> $B12
                  }
                }
                %257:u32 = load %tileId
                %258:ptr<storage, atomic<u32>, read_write> = access %tileLightId, 0u, %257, 0u
                %259:u32 = atomicAdd %258, 1u
                %offset:ptr<function, u32, read_write> = var, %259
                %261:u32 = load %offset
                %262:ptr<uniform, u32, read> = access %config, 4u
                %263:u32 = load %262
                %264:bool = gte %261, %263
                if %264 [t: $B30] {  # if_12
                  $B30: {  # true
                    continue  # -> $B12
                  }
                }
                %265:u32 = load %tileId
                %266:u32 = load %offset
                %267:ptr<storage, u32, read_write> = access %tileLightId, 0u, %265, 1u, %266
                %268:u32 = access %GlobalInvocationID, 0u
                store %267, %268
                exit_if  # if_9
              }
            }
            continue  # -> $B12
          }
          $B12: {  # continuing
            %269:i32 = load %x
            %270:i32 = add %269, 1i
            store %x, %270
            next_iteration  # -> $B11
          }
        }
        continue  # -> $B7
      }
      $B7: {  # continuing
        %271:i32 = load %y
        %272:i32 = add %271, 1i
        store %y, %272
        next_iteration  # -> $B6
      }
    }
    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.  *
********************************************************************