enable chromium_experimental_subgroup_matrix, subgroups;

@group(0) @binding(0) var<storage, read_write> s_var : array<f32>;

override o : u32 = 1024;

var<workgroup> wg_var : array<f32, o>;

@compute @workgroup_size(32)
fn main() {
  let m = subgroupMatrixLoad<subgroup_matrix_left<f32, 8, 8>, row_major>(&(s_var), 0, 8);
  subgroupMatrixStore<row_major>(&(wg_var), 0, m, 8);
  workgroupBarrier();
  let m2 = subgroupMatrixLoad<subgroup_matrix_right<f32, 8, 8>, col_major>(&(wg_var), 0, 8);
  subgroupMatrixStore<col_major>(&(s_var), 0, m2, 8);
}