test/tint/builtins/gen/var/subgroupMatrixStore/8a88da.wgsl.expected.spvasm

SKIP: FAILED


enable chromium_experimental_subgroup_matrix;

var<workgroup> arg_0 : array<u32, 1024>;

fn subgroupMatrixStore_8a88da() {
  var arg_1 = 1u;
  var arg_2 = subgroup_matrix_result<u8, 8, 8>();
  const arg_3 = true;
  var arg_4 = 8u;
  subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
}

@compute @workgroup_size(1)
fn compute_main() {
  subgroupMatrixStore_8a88da();
}

Failed to generate SPIR-V: :22:20 error: spirv.cooperative_matrix_store: no matching call to 'spirv.cooperative_matrix_store(ptr<workgroup, u32, read_write>, subgroup_matrix_result<u8, 8, 8>, u32, u32, u32)'

1 candidate function:
 • 'spirv.cooperative_matrix_store(ptr<workgroup' or 'storage, S, write' or 'read_write>  ✓ , subgroup_matrix<K, S, C, R>  ✗ , u32  ✓ , u32  ✓ , u32  ✓ )' where:
      ✓  'S' is 'f32', 'i32', 'u32' or 'f16'

        %16:void = spirv.cooperative_matrix_store %15, %8, 1u, %10, 32u
                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

:20:7 note: in block
      $B3: {  # true
      ^^^

note: # Disassembly
$B1: {  # root
  %arg_0:ptr<workgroup, array<u32, 1024>, read_write> = var undef
}

%subgroupMatrixStore_8a88da = func():void {
  $B2: {
    %arg_1:ptr<function, u32, read_write> = var 1u
    %4:subgroup_matrix_result<u8, 8, 8> = construct
    %arg_2:ptr<function, subgroup_matrix_result<u8, 8, 8>, read_write> = var %4
    %arg_4:ptr<function, u32, read_write> = var 8u
    %7:u32 = load %arg_1
    %8:subgroup_matrix_result<u8, 8, 8> = load %arg_2
    %9:u32 = load %arg_4
    %10:u32 = max %9, 8u
    %11:u32 = mul %10, 7u
    %12:u32 = add %7, %11
    %13:u32 = add %12, 8u
    %14:bool = lte %13, 4096u
    if %14 [t: $B3] {  # if_1
      $B3: {  # true
        %15:ptr<workgroup, u32, read_write> = access %arg_0, %7
        %16:void = spirv.cooperative_matrix_store %15, %8, 1u, %10, 32u
        exit_if  # if_1
      }
    }
    ret
  }
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func(%tint_local_index:u32 [@local_invocation_index]):void {
  $B4: {
    loop [i: $B5, b: $B6, c: $B7] {  # loop_1
      $B5: {  # initializer
        next_iteration %tint_local_index  # -> $B6
      }
      $B6 (%idx:u32): {  # body
        %20:bool = gte %idx, 1024u
        if %20 [t: $B8] {  # if_2
          $B8: {  # true
            exit_loop  # loop_1
          }
        }
        %21:ptr<workgroup, u32, read_write> = access %arg_0, %idx
        store %21, 0u
        continue  # -> $B7
      }
      $B7: {  # continuing
        %22:u32 = add %idx, 1u
        next_iteration %22  # -> $B6
      }
    }
    %23:void = workgroupBarrier
    %24:void = call %subgroupMatrixStore_8a88da
    ret
  }
}


tint executable returned error: exit status 1