blob: 4b04255c0811205d926a57d23d8b3dd67f2a0464 [file]
// [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: 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<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);
}