| 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 |