blob: 191517f1383a8487e8ceb676ebf8992975ec63c6 [file] [log] [blame]
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() {
subgroupMatrixStore(&(buffer), 0, subgroup_matrix_left<f32, 8, 8>(), false, 64);
subgroupMatrixStore(&(buffer), 0, array<subgroup_matrix_left<f32, 8, 8>, 4>()[1], false, 64);
subgroupMatrixStore(&(buffer), 0, array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>()[2][3], false, 64);
subgroupMatrixStore(&(buffer), 0, S().l, false, 64);
subgroupMatrixStore(&(buffer), 0, S_Nested().s.r, false, 64);
subgroupMatrixStore(&(buffer), 0, subgroup_matrix_left<f32, 8, 8>(42), false, 64);
subgroupMatrixStore(&(buffer), 0, array<subgroup_matrix_left<f32, 8, 8>, 2>(subgroup_matrix_left<f32, 8, 8>(42), subgroup_matrix_left<f32, 8, 8>(100))[1], false, 64);
subgroupMatrixStore(&(buffer), 0, array<array<subgroup_matrix_left<f32, 8, 8>, 2>, 2>(array<subgroup_matrix_left<f32, 8, 8>, 2>(subgroup_matrix_left<f32, 8, 8>(42), subgroup_matrix_left<f32, 8, 8>(100)), array<subgroup_matrix_left<f32, 8, 8>, 2>(subgroup_matrix_left<f32, 8, 8>(-(7)), subgroup_matrix_left<f32, 8, 8>(-(42))))[1][0], false, 64);
subgroupMatrixStore(&(buffer), 0, S(subgroup_matrix_left<f32, 8, 8>(42), subgroup_matrix_right<f32, 8, 8>(100)).l, false, 64);
subgroupMatrixStore(&(buffer), 0, S_Nested(S(subgroup_matrix_left<f32, 8, 8>(42), subgroup_matrix_right<f32, 8, 8>(100))).s.r, false, 64);
}