test/tint/extensions/subgroup_matrix/pointers.wgsl

// [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: 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<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);
}