test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl

SKIP: FAILED

../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(16) {
  arr:array<vec4<i32>, 4> @offset(0)
}

$B1: {  # root
  %src_private:ptr<private, array<vec4<i32>, 4>, read_write> = var
  %src_workgroup:ptr<workgroup, array<vec4<i32>, 4>, read_write> = var
  %src_uniform:ptr<uniform, S, read> = var @binding_point(0, 0)
  %src_storage:ptr<storage, S, read_write> = var @binding_point(0, 1)
  %dst:ptr<workgroup, array<vec4<i32>, 4>, read_write> = var
  %dst_nested:ptr<workgroup, array<array<array<i32, 2>, 3>, 4>, read_write> = var
}

%ret_arr = func():array<vec4<i32>, 4> {
  $B2: {
    ret array<vec4<i32>, 4>(vec4<i32>(0i))
  }
}
%ret_struct_arr = func():S {
  $B3: {
    ret S(array<vec4<i32>, 4>(vec4<i32>(0i)))
  }
}
%foo = func(%src_param:array<vec4<i32>, 4>):void {
  $B4: {
    %src_function:ptr<function, array<vec4<i32>, 4>, read_write> = var
    store %dst, array<vec4<i32>, 4>(vec4<i32>(1i), vec4<i32>(2i), vec4<i32>(3i), vec4<i32>(3i))
    store %dst, %src_param
    %12:array<vec4<i32>, 4> = call %ret_arr
    store %dst, %12
    %src_let:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
    store %dst, %src_let
    %14:array<vec4<i32>, 4> = load %src_function
    store %dst, %14
    %15:array<vec4<i32>, 4> = load %src_private
    store %dst, %15
    %16:array<vec4<i32>, 4> = load %src_workgroup
    store %dst, %16
    %17:S = call %ret_struct_arr
    %18:array<vec4<i32>, 4> = access %17, 0u
    store %dst, %18
    %19:ptr<uniform, array<vec4<i32>, 4>, read> = access %src_uniform, 0u
    %20:array<vec4<i32>, 4> = load %19
    store %dst, %20
    %21:ptr<storage, array<vec4<i32>, 4>, read_write> = access %src_storage, 0u
    %22:array<vec4<i32>, 4> = load %21
    store %dst, %22
    %src_nested:ptr<function, array<array<array<i32, 2>, 3>, 4>, read_write> = var
    %24:array<array<array<i32, 2>, 3>, 4> = load %src_nested
    store %dst_nested, %24
    ret
  }
}
%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
  $B5: {
    loop [i: $B6, b: $B7, c: $B8] {  # loop_1
      $B6: {  # initializer
        next_iteration %tint_local_index  # -> $B7
      }
      $B7 (%idx:u32): {  # body
        %28:bool = gte %idx, 4u
        if %28 [t: $B9] {  # if_1
          $B9: {  # true
            exit_loop  # loop_1
          }
        }
        %29:ptr<workgroup, vec4<i32>, read_write> = access %src_workgroup, %idx
        store %29, vec4<i32>(0i)
        %30:ptr<workgroup, vec4<i32>, read_write> = access %dst, %idx
        store %30, vec4<i32>(0i)
        continue  # -> $B8
      }
      $B8: {  # continuing
        %31:u32 = add %idx, 1u
        next_iteration %31  # -> $B7
      }
    }
    loop [i: $B10, b: $B11, c: $B12] {  # loop_2
      $B10: {  # initializer
        next_iteration %tint_local_index  # -> $B11
      }
      $B11 (%idx_1:u32): {  # body
        %33:bool = gte %idx_1, 24u
        if %33 [t: $B13] {  # if_2
          $B13: {  # true
            exit_loop  # loop_2
          }
        }
        %34:u32 = mod %idx_1, 2u
        %35:u32 = div %idx_1, 2u
        %36:u32 = mod %35, 3u
        %37:u32 = div %idx_1, 6u
        %38:ptr<workgroup, i32, read_write> = access %dst_nested, %37, %36, %34
        store %38, 0i
        continue  # -> $B12
      }
      $B12: {  # continuing
        %39:u32 = add %idx_1, 1u
        next_iteration %39  # -> $B11
      }
    }
    %40:void = msl.threadgroup_barrier 4u
    %val:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
    %42:void = call %foo, %val
    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.  *
********************************************************************