| SKIP: FAILED |
| |
| |
| enable chromium_experimental_subgroup_matrix; |
| |
| struct SB_RW { |
| arg_0 : array<u32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW; |
| |
| fn subgroupMatrixStore_152780() { |
| var arg_1 = 1u; |
| var arg_2 = subgroup_matrix_result<u8, 8, 8>(); |
| const arg_3 = true; |
| var arg_4 = 8u; |
| subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4); |
| } |
| |
| @compute @workgroup_size(1) |
| fn compute_main() { |
| subgroupMatrixStore_152780(); |
| } |
| |
| Failed to generate SPIR-V: :29:20 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' |
| |
| %19:void = spirv.cooperative_matrix_store %18, %9, 1u, %11, 32u |
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| :27:7 note: in block |
| $B3: { # true |
| ^^^ |
| |
| note: # Disassembly |
| SB_RW = struct @align(4), @block { |
| arg_0:array<u32> @offset(0) |
| } |
| |
| $B1: { # root |
| %sb_rw:ptr<storage, SB_RW, read_write> = var undef @binding_point(0, 0) |
| } |
| |
| %subgroupMatrixStore_152780 = func():void { |
| $B2: { |
| %arg_1:ptr<function, u32, read_write> = var 1u |
| %4:subgroup_matrix_result<u8, 8, 8> = construct |
| %arg_2:ptr<function, subgroup_matrix_result<u8, 8, 8>, read_write> = var %4 |
| %arg_4:ptr<function, u32, read_write> = var 8u |
| %7:ptr<storage, array<u32>, read_write> = access %sb_rw, 0u |
| %8:u32 = load %arg_1 |
| %9:subgroup_matrix_result<u8, 8, 8> = load %arg_2 |
| %10:u32 = load %arg_4 |
| %11:u32 = max %10, 8u |
| %12:u32 = spirv.array_length %sb_rw, 0u |
| %13:u32 = mul %12, 4u |
| %14:u32 = mul %11, 7u |
| %15:u32 = add %8, %14 |
| %16:u32 = add %15, 8u |
| %17:bool = lte %16, %13 |
| if %17 [t: $B3] { # if_1 |
| $B3: { # true |
| %18:ptr<storage, u32, read_write> = access %7, %8 |
| %19:void = spirv.cooperative_matrix_store %18, %9, 1u, %11, 32u |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| %compute_main = @compute @workgroup_size(1i, 1i, 1i) func():void { |
| $B4: { |
| %21:void = call %subgroupMatrixStore_152780 |
| ret |
| } |
| } |
| |
| |
| tint executable returned error: exit status 1 |