| // [hlsl-dxc] flags: --hlsl-shader-model 6.10 | |
| enable chromium_experimental_subgroup_matrix; | |
| struct S { | |
| l: subgroup_matrix_left<f32, 8, 8>, | |
| r: subgroup_matrix_right<f32, 8, 8>, | |
| } | |
| struct S_Nested { | |
| s: S, | |
| } | |
| @group(0) @binding(0) | |
| var<storage, read_write> buffer: array<f32>; | |
| fn foo( | |
| m: subgroup_matrix_left<f32, 8, 8>, | |
| m_array: array<subgroup_matrix_left<f32, 8, 8>, 4>, | |
| m_nested_array: array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>, | |
| m_struct: S, | |
| m_nested_struct: S_Nested, | |
| ) { | |
| subgroupMatrixStore<row_major>(&buffer, 0, m, 64); | |
| subgroupMatrixStore<row_major>(&buffer, 0, m_array[0], 64); | |
| subgroupMatrixStore<row_major>(&buffer, 0, m_nested_array[1][2], 64); | |
| subgroupMatrixStore<row_major>(&buffer, 0, m_struct.l, 64); | |
| subgroupMatrixStore<row_major>(&buffer, 0, m_nested_struct.s.r, 64); | |
| } | |
| @compute @workgroup_size(64) | |
| fn main() { | |
| var m: subgroup_matrix_left<f32, 8, 8>; | |
| var m_array: array<subgroup_matrix_left<f32, 8, 8>, 4>; | |
| var m_nested_array: array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>; | |
| var m_struct: S; | |
| var m_nested_struct: S_Nested; | |
| foo(m, m_array, m_nested_array, m_struct, m_nested_struct); | |
| } |