SKIP: FAILED

<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: Params = struct @align(4) {
  filterDim:u32 @offset(0)
  blockDim:u32 @offset(4)
}

Flip = struct @align(4) {
  value:u32 @offset(0)
}

%b1 = block {  # root
  %samp:ptr<handle, sampler, read> = var @binding_point(0, 0)
  %params:ptr<uniform, Params, read> = var @binding_point(0, 1)
  %inputTex:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 1)
  %outputTex:ptr<handle, texture_storage_2d<rgba8unorm, write>, read> = var @binding_point(1, 2)
  %flip:ptr<uniform, Flip, read> = var @binding_point(1, 3)
  %tile:ptr<workgroup, array<array<vec3<f32>, 256>, 4>, read_write> = var
}

%tint_symbol = @compute @workgroup_size(64, 1, 1) func(%WorkGroupID:vec3<u32> [@workgroup_id], %LocalInvocationID:vec3<u32> [@local_invocation_id]):void -> %b2 {
  %b2 = block {
    %10:ptr<uniform, u32, read> = access %params, 0u
    %11:u32 = load %10
    %12:u32 = sub %11, 1u
    %filterOffset:u32 = div %12, 2u
    %14:texture_2d<f32> = load %inputTex
    %dims:vec2<u32> = textureDimensions %14, 0i
    %16:vec2<u32> = swizzle %WorkGroupID, xy
    %17:ptr<uniform, u32, read> = access %params, 1u
    %18:u32 = load %17
    %19:vec2<u32> = construct %18, 4u
    %20:vec2<u32> = mul %16, %19
    %21:vec2<u32> = swizzle %LocalInvocationID, xy
    %22:vec2<u32> = mul %21, vec2<u32>(4u, 1u)
    %23:vec2<u32> = add %20, %22
    %24:vec2<u32> = construct %filterOffset, 0u
    %baseIndex:vec2<u32> = sub %23, %24
    loop [i: %b3, b: %b4, c: %b5] {  # loop_1
      %b3 = block {  # initializer
        %r:ptr<function, u32, read_write> = var, 0u
        next_iteration %b4
      }
      %b4 = block {  # body
        %27:u32 = load %r
        %28:bool = lt %27, 4u
        if %28 [t: %b6, f: %b7] {  # if_1
          %b6 = block {  # true
            exit_if  # if_1
          }
          %b7 = block {  # false
            exit_loop  # loop_1
          }
        }
        loop [i: %b8, b: %b9, c: %b10] {  # loop_2
          %b8 = block {  # initializer
            %c:ptr<function, u32, read_write> = var, 0u
            next_iteration %b9
          }
          %b9 = block {  # body
            %30:u32 = load %c
            %31:bool = lt %30, 4u
            if %31 [t: %b11, f: %b12] {  # if_2
              %b11 = block {  # true
                exit_if  # if_2
              }
              %b12 = block {  # false
                exit_loop  # loop_2
              }
            }
            %32:u32 = load %c
            %33:u32 = load %r
            %34:vec2<u32> = construct %32, %33
            %35:vec2<u32> = add %baseIndex, %34
            %loadIndex:ptr<function, vec2<u32>, read_write> = var, %35
            %37:ptr<uniform, u32, read> = access %flip, 0u
            %38:u32 = load %37
            %39:bool = neq %38, 0u
            if %39 [t: %b13] {  # if_3
              %b13 = block {  # true
                %40:vec2<u32> = load %loadIndex
                %41:vec2<u32> = swizzle %40, yx
                store %loadIndex, %41
                exit_if  # if_3
              }
            }
            %42:u32 = load %r
            %43:u32 = access %LocalInvocationID, 0u
            %44:u32 = mul 4u, %43
            %45:u32 = load %c
            %46:u32 = add %44, %45
            %47:ptr<workgroup, vec3<f32>, read_write> = access %tile, %42, %46
            %48:texture_2d<f32> = load %inputTex
            %49:sampler = load %samp
            %50:vec2<u32> = load %loadIndex
            %51:vec2<f32> = convert %50
            %52:vec2<f32> = add %51, vec2<f32>(0.25f)
            %53:vec2<f32> = convert %dims
            %54:vec2<f32> = div %52, %53
            %55:vec4<f32> = textureSampleLevel %48, %49, %54, 0.0f
            %56:vec3<f32> = swizzle %55, xyz
            store %47, %56
            continue %b10
          }
          %b10 = block {  # continuing
            %57:u32 = load %c
            %58:u32 = add %57, 1u
            store %c, %58
            next_iteration %b9
          }
        }
        continue %b5
      }
      %b5 = block {  # continuing
        %59:u32 = load %r
        %60:u32 = add %59, 1u
        store %r, %60
        next_iteration %b4
      }
    }
    %61:void = workgroupBarrier
    loop [i: %b14, b: %b15, c: %b16] {  # loop_3
      %b14 = block {  # initializer
        %r_1:ptr<function, u32, read_write> = var, 0u  # %r_1: 'r'
        next_iteration %b15
      }
      %b15 = block {  # body
        %63:u32 = load %r_1
        %64:bool = lt %63, 4u
        if %64 [t: %b17, f: %b18] {  # if_4
          %b17 = block {  # true
            exit_if  # if_4
          }
          %b18 = block {  # false
            exit_loop  # loop_3
          }
        }
        loop [i: %b19, b: %b20, c: %b21] {  # loop_4
          %b19 = block {  # initializer
            %c_1:ptr<function, u32, read_write> = var, 0u  # %c_1: 'c'
            next_iteration %b20
          }
          %b20 = block {  # body
            %66:u32 = load %c_1
            %67:bool = lt %66, 4u
            if %67 [t: %b22, f: %b23] {  # if_5
              %b22 = block {  # true
                exit_if  # if_5
              }
              %b23 = block {  # false
                exit_loop  # loop_4
              }
            }
            %68:u32 = load %c_1
            %69:u32 = load %r_1
            %70:vec2<u32> = construct %68, %69
            %71:vec2<u32> = add %baseIndex, %70
            %writeIndex:ptr<function, vec2<u32>, read_write> = var, %71
            %73:ptr<uniform, u32, read> = access %flip, 0u
            %74:u32 = load %73
            %75:bool = neq %74, 0u
            if %75 [t: %b24] {  # if_6
              %b24 = block {  # true
                %76:vec2<u32> = load %writeIndex
                %77:vec2<u32> = swizzle %76, yx
                store %writeIndex, %77
                exit_if  # if_6
              }
            }
            %78:u32 = access %LocalInvocationID, 0u
            %79:u32 = mul 4u, %78
            %80:u32 = load %c_1
            %center:u32 = add %79, %80
            %82:bool = gte %center, %filterOffset
            %83:bool = if %82 [t: %b25, f: %b26] {  # if_7
              %b25 = block {  # true
                %84:u32 = sub 256u, %filterOffset
                %85:bool = lt %center, %84
                exit_if %85  # if_7
              }
              %b26 = block {  # false
                exit_if false  # if_7
              }
            }
            %86:bool = if %83 [t: %b27, f: %b28] {  # if_8
              %b27 = block {  # true
                %87:vec2<u32> = load %writeIndex
                %88:vec2<bool> = lt %87, %dims
                %89:bool = all %88
                exit_if %89  # if_8
              }
              %b28 = block {  # false
                exit_if false  # if_8
              }
            }
            if %86 [t: %b29] {  # if_9
              %b29 = block {  # true
                %acc:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
                loop [i: %b30, b: %b31, c: %b32] {  # loop_5
                  %b30 = block {  # initializer
                    %f:ptr<function, u32, read_write> = var, 0u
                    next_iteration %b31
                  }
                  %b31 = block {  # body
                    %92:u32 = load %f
                    %93:ptr<uniform, u32, read> = access %params, 0u
                    %94:u32 = load %93
                    %95:bool = lt %92, %94
                    if %95 [t: %b33, f: %b34] {  # if_10
                      %b33 = block {  # true
                        exit_if  # if_10
                      }
                      %b34 = block {  # false
                        exit_loop  # loop_5
                      }
                    }
                    %96:u32 = load %f
                    %97:u32 = add %center, %96
                    %98:u32 = sub %97, %filterOffset
                    %i:ptr<function, u32, read_write> = var, %98
                    %100:vec3<f32> = load %acc
                    %101:ptr<uniform, u32, read> = access %params, 0u
                    %102:u32 = load %101
                    %103:f32 = convert %102
                    %104:f32 = div 1.0f, %103
                    %105:u32 = load %r_1
                    %106:u32 = load %i
                    %107:ptr<workgroup, vec3<f32>, read_write> = access %tile, %105, %106
                    %108:vec3<f32> = load %107
                    %109:vec3<f32> = mul %104, %108
                    %110:vec3<f32> = add %100, %109
                    store %acc, %110
                    continue %b32
                  }
                  %b32 = block {  # continuing
                    %111:u32 = load %f
                    %112:u32 = add %111, 1u
                    store %f, %112
                    next_iteration %b31
                  }
                }
                %113:texture_storage_2d<rgba8unorm, write> = load %outputTex
                %114:vec2<u32> = load %writeIndex
                %115:vec3<f32> = load %acc
                %116:vec4<f32> = construct %115, 1.0f
                %117:void = textureStore %113, %114, %116
                exit_if  # if_9
              }
            }
            continue %b21
          }
          %b21 = block {  # continuing
            %118:u32 = load %c_1
            %119:u32 = add %118, 1u
            store %c_1, %119
            next_iteration %b20
          }
        }
        continue %b16
      }
      %b16 = block {  # continuing
        %120:u32 = load %r_1
        %121:u32 = add %120, 1u
        store %r_1, %121
        next_iteration %b15
      }
    }
    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.  *
********************************************************************