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: ptr<function, subgroup_matrix_left<f32, 8, 8>>,
  m_array: ptr<function, array<subgroup_matrix_left<f32, 8, 8>, 4>>,
  m_nested_array: ptr<function, array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>>,
  m_struct: ptr<function, S>,
  m_nested_struct: ptr<function, 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);
}

@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);
}