blob: 2ae4ba72a2e66fe80ba902f6b93783c1bdaed2f7 [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>;
fn foo(m : subgroup_matrix_left<f32, 8, 8>, m_array : array<subgroup_matrix_left<f32, 8, 8>, 4>, m_nested_array : array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>, m_struct : S, 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);
}
@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);
}