test/tint/builtins/gen/literal/subgroupMatrixLoad/1b3162.wgsl.expected.spvasm

SKIP: FAILED


enable chromium_experimental_subgroup_matrix;

@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;

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

fn subgroupMatrixLoad_1b3162() -> subgroup_matrix_right<i8, 8, 8> {
  var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(arg_0), 1u, true, 8u);
  return res;
}

@compute @workgroup_size(1)
fn compute_main() {
  subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_1b3162(), false, 64);
}

Failed to generate SPIR-V: :13:42 error: spirv.cooperative_matrix_load: no matching call to 'spirv.cooperative_matrix_load<subgroup_matrix_right<i8, 8, 8>>(ptr<workgroup, i32, read_write>, u32, u32, u32)'

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

    %5:subgroup_matrix_right<i8, 8, 8> = spirv.cooperative_matrix_load<subgroup_matrix_right<i8, 8, 8>> %4, 1u, 8u, 32u
                                         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

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

:45:16 error: spirv.cooperative_matrix_store: no matching call to 'spirv.cooperative_matrix_store(ptr<storage, i32, read_write>, subgroup_matrix_right<i8, 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'

    %18:void = spirv.cooperative_matrix_store %17, %15, 0u, 64u, 32u
               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

:20:3 note: in block
  $B3: {
  ^^^

note: # Disassembly
prevent_dce_block = struct @align(4), @block {
  inner:array<i32, 1024> @offset(0)
}

$B1: {  # root
  %1:ptr<storage, prevent_dce_block, read_write> = var undef @binding_point(0, 0)
  %arg_0:ptr<workgroup, array<i32, 1024>, read_write> = var undef
}

%subgroupMatrixLoad_1b3162 = func():subgroup_matrix_right<i8, 8, 8> {
  $B2: {
    %4:ptr<workgroup, i32, read_write> = access %arg_0, 1u
    %5:subgroup_matrix_right<i8, 8, 8> = spirv.cooperative_matrix_load<subgroup_matrix_right<i8, 8, 8>> %4, 1u, 8u, 32u
    %res:ptr<function, subgroup_matrix_right<i8, 8, 8>, read_write> = var %5
    %7:subgroup_matrix_right<i8, 8, 8> = load %res
    ret %7
  }
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func(%tint_local_index:u32 [@local_invocation_index]):void {
  $B3: {
    loop [i: $B4, b: $B5, c: $B6] {  # loop_1
      $B4: {  # initializer
        next_iteration %tint_local_index  # -> $B5
      }
      $B5 (%idx:u32): {  # body
        %11:bool = gte %idx, 1024u
        if %11 [t: $B7] {  # if_1
          $B7: {  # true
            exit_loop  # loop_1
          }
        }
        %12:ptr<workgroup, i32, read_write> = access %arg_0, %idx
        store %12, 0i
        continue  # -> $B6
      }
      $B6: {  # continuing
        %13:u32 = add %idx, 1u
        next_iteration %13  # -> $B5
      }
    }
    %14:void = workgroupBarrier
    %15:subgroup_matrix_right<i8, 8, 8> = call %subgroupMatrixLoad_1b3162
    %16:ptr<storage, array<i32, 1024>, read_write> = access %1, 0u
    %17:ptr<storage, i32, read_write> = access %16, 0u
    %18:void = spirv.cooperative_matrix_store %17, %15, 0u, 64u, 32u
    ret
  }
}


tint executable returned error: exit status 1