blob: 1b80aa4bab4affffbf02be4fde5523b6e15cbd38 [file] [log] [blame]
SKIP: FAILED
enable chromium_experimental_subgroup_matrix;
enable f16;
@group(0) @binding(0) var<storage, read_write> prevent_dce : array<f16, 1024>;
fn subgroupMatrixMultiply_009084() -> subgroup_matrix_result<f16, 8, 8> {
var res : subgroup_matrix_result<f16, 8, 8> = subgroupMatrixMultiply<f16>(subgroup_matrix_left<i8, 8, 8>(), subgroup_matrix_right<i8, 8, 8>());
return res;
}
@compute @workgroup_size(1)
fn compute_main() {
subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixMultiply_009084(), false, 64);
}
Failed to generate SPIR-V: :14:44 error: spirv.cooperative_matrix_mul_add: no matching call to 'spirv.cooperative_matrix_mul_add(subgroup_matrix_left<i8, 8, 8>, subgroup_matrix_right<i8, 8, 8>, subgroup_matrix_result<f16, 8, 8>, u32)'
1 candidate function:
'spirv.cooperative_matrix_mul_add(subgroup_matrix<left, T, K, R> ✗ , subgroup_matrix<right, T, C, K> ✗ , subgroup_matrix<result, TR, C, R> ✓ , u32 ✓ ) -> subgroup_matrix<result, TR, C, R>' where:
'T' is 'f32', 'f16', 'u32' or 'i32'
'TR' is 'f32', 'f16', 'u32' or 'i32'
%6:subgroup_matrix_result<f16, 8, 8> = spirv.cooperative_matrix_mul_add %3, %4, %5, 3u
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
:10:3 note: in block
$B2: {
^^^
note: # Disassembly
prevent_dce_block = struct @align(2), @block {
inner:array<f16, 1024> @offset(0)
}
$B1: { # root
%1:ptr<storage, prevent_dce_block, read_write> = var undef @binding_point(0, 0)
}
%subgroupMatrixMultiply_009084 = func():subgroup_matrix_result<f16, 8, 8> {
$B2: {
%3:subgroup_matrix_left<i8, 8, 8> = construct
%4:subgroup_matrix_right<i8, 8, 8> = construct
%5:subgroup_matrix_result<f16, 8, 8> = construct
%6:subgroup_matrix_result<f16, 8, 8> = spirv.cooperative_matrix_mul_add %3, %4, %5, 3u
%res:ptr<function, subgroup_matrix_result<f16, 8, 8>, read_write> = var %6
%8:subgroup_matrix_result<f16, 8, 8> = load %res
ret %8
}
}
%compute_main = @compute @workgroup_size(1i, 1i, 1i) func():void {
$B3: {
%10:subgroup_matrix_result<f16, 8, 8> = call %subgroupMatrixMultiply_009084
%11:ptr<storage, array<f16, 1024>, read_write> = access %1, 0u
%12:ptr<storage, f16, read_write> = access %11, 0u
%13:void = spirv.cooperative_matrix_store %12, %10, 0u, 64u, 32u
ret
}
}
tint executable returned error: exit status 1