|  | 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>; | 
|  |  | 
|  | @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; | 
|  |  | 
|  | subgroupMatrixStore(&buffer, 0, m, false, 64); | 
|  | subgroupMatrixStore(&buffer, 0, m_array[0], false, 64); | 
|  | subgroupMatrixStore(&buffer, 0, m_nested_array[1][2], false, 64); | 
|  | subgroupMatrixStore(&buffer, 0, m_struct.l, false, 64); | 
|  | subgroupMatrixStore(&buffer, 0, m_nested_struct.s.r, false, 64); | 
|  | } |