test/tint/builtins/gen/literal/subgroupMatrixStore/c9f1a5.wgsl.expected.spvasm

SKIP: FAILED


enable chromium_experimental_subgroup_matrix;

struct SB_RW {
  arg_0 : array<u32, 1024>,
}

@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;

fn subgroupMatrixStore_c9f1a5() {
  subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_result<u8, 8, 8>(), true, 8u);
}

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

Failed to generate SPIR-V: :18:15 error: spirv.cooperative_matrix_store: no matching call to 'spirv.cooperative_matrix_store(ptr<storage, 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'

    %6:void = spirv.cooperative_matrix_store %5, %4, 1u, 8u, 32u
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

:14:3 note: in block
  $B2: {
  ^^^

note: # Disassembly
SB_RW = struct @align(4) {
  arg_0:array<u32, 1024> @offset(0)
}

sb_rw_block = struct @align(4), @block {
  inner:SB_RW @offset(0)
}

$B1: {  # root
  %1:ptr<storage, sb_rw_block, read_write> = var undef @binding_point(0, 0)
}

%subgroupMatrixStore_c9f1a5 = func():void {
  $B2: {
    %3:ptr<storage, array<u32, 1024>, read_write> = access %1, 0u, 0u
    %4:subgroup_matrix_result<u8, 8, 8> = construct
    %5:ptr<storage, u32, read_write> = access %3, 1u
    %6:void = spirv.cooperative_matrix_store %5, %4, 1u, 8u, 32u
    ret
  }
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func():void {
  $B3: {
    %8:void = call %subgroupMatrixStore_c9f1a5
    ret
  }
}


tint executable returned error: exit status 1