test/tint/builtins/gen/var/subgroupMatrixMultiply/5677fc.wgsl.expected.msl

SKIP: FAILED

tint_module_vars_struct = struct @align(1) {
  prevent_dce:ptr<storage, array<f32, 1024>, read_write> @offset(0)
}

%subgroupMatrixMultiply_5677fc = func():subgroup_matrix_result<f32, 8, 8> {
  $B1: {
    %2:subgroup_matrix_left<f16, 8, 8> = msl.make_filled_simdgroup_matrix<subgroup_matrix_left<f16, 8, 8>> 0.0h
    %arg_0:ptr<function, subgroup_matrix_left<f16, 8, 8>, read_write> = var %2
    %4:subgroup_matrix_right<f16, 8, 8> = msl.make_filled_simdgroup_matrix<subgroup_matrix_right<f16, 8, 8>> 0.0h
    %arg_1:ptr<function, subgroup_matrix_right<f16, 8, 8>, read_write> = var %4
    %6:subgroup_matrix_left<f16, 8, 8> = load %arg_0
    %7:subgroup_matrix_right<f16, 8, 8> = load %arg_1
    %8:ptr<function, subgroup_matrix_result<f32, 8, 8>, read_write> = var undef
    %9:subgroup_matrix_result<f32, 8, 8> = load %8
    %10:void = msl.simdgroup_multiply %9, %6, %7
    %11:subgroup_matrix_result<f32, 8, 8> = load %8
    %res:ptr<function, subgroup_matrix_result<f32, 8, 8>, read_write> = var %11
    %13:subgroup_matrix_result<f32, 8, 8> = load %res
    ret %13
  }
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func(%prevent_dce:ptr<storage, array<f32, 1024>, read_write> [@binding_point(0, 0)]):void {
  $B2: {
    %16:tint_module_vars_struct = construct %prevent_dce
    %tint_module_vars:tint_module_vars_struct = let %16
    %18:subgroup_matrix_result<f32, 8, 8> = call %subgroupMatrixMultiply_5677fc
    %19:ptr<storage, array<f32, 1024>, read_write> = access %tint_module_vars, 0u
    %20:ptr<storage, f32, read_write> = access %19, 0u
    %21:u64 = msl.convert 64u
    %22:void = msl.simdgroup_store %18, %20, %21, vec2<u64>(0u64), false
    ret
  }
}
Failed to generate: :15:16 error: msl.simdgroup_multiply: no matching call to 'msl.simdgroup_multiply(subgroup_matrix_result<f32, 8, 8>, subgroup_matrix_left<f16, 8, 8>, subgroup_matrix_right<f16, 8, 8>)'

1 candidate function:
 • 'msl.simdgroup_multiply(subgroup_matrix<result, S, C, R>  ✓ , subgroup_matrix<left, S, K, R>  ✗ , subgroup_matrix<right, S, C, K>  ✗ )' where:
      ✓  'S' is 'f32' or 'f16'

    %10:void = msl.simdgroup_multiply %9, %6, %7
               ^^^^^^^^^^^^^^^^^^^^^^

:6:3 note: in block
  $B1: {
  ^^^

note: # Disassembly
tint_module_vars_struct = struct @align(1) {
  prevent_dce:ptr<storage, array<f32, 1024>, read_write> @offset(0)
}

%subgroupMatrixMultiply_5677fc = func():subgroup_matrix_result<f32, 8, 8> {
  $B1: {
    %2:subgroup_matrix_left<f16, 8, 8> = msl.make_filled_simdgroup_matrix<subgroup_matrix_left<f16, 8, 8>> 0.0h
    %arg_0:ptr<function, subgroup_matrix_left<f16, 8, 8>, read_write> = var %2
    %4:subgroup_matrix_right<f16, 8, 8> = msl.make_filled_simdgroup_matrix<subgroup_matrix_right<f16, 8, 8>> 0.0h
    %arg_1:ptr<function, subgroup_matrix_right<f16, 8, 8>, read_write> = var %4
    %6:subgroup_matrix_left<f16, 8, 8> = load %arg_0
    %7:subgroup_matrix_right<f16, 8, 8> = load %arg_1
    %8:ptr<function, subgroup_matrix_result<f32, 8, 8>, read_write> = var undef
    %9:subgroup_matrix_result<f32, 8, 8> = load %8
    %10:void = msl.simdgroup_multiply %9, %6, %7
    %11:subgroup_matrix_result<f32, 8, 8> = load %8
    %res:ptr<function, subgroup_matrix_result<f32, 8, 8>, read_write> = var %11
    %13:subgroup_matrix_result<f32, 8, 8> = load %res
    ret %13
  }
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func(%prevent_dce:ptr<storage, array<f32, 1024>, read_write> [@binding_point(0, 0)]):void {
  $B2: {
    %16:tint_module_vars_struct = construct %prevent_dce
    %tint_module_vars:tint_module_vars_struct = let %16
    %18:subgroup_matrix_result<f32, 8, 8> = call %subgroupMatrixMultiply_5677fc
    %19:ptr<storage, array<f32, 1024>, read_write> = access %tint_module_vars, 0u
    %20:ptr<storage, f32, read_write> = access %19, 0u
    %21:u64 = msl.convert 64u
    %22:void = msl.simdgroup_store %18, %20, %21, vec2<u64>(0u64), false
    ret
  }
}


tint executable returned error: exit status 1