blob: eb41b4d2263fd023dc6b4cb287eefec07470c004 [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>;
@compute @workgroup_size(64)
fn main() {
// Zero init.
subgroupMatrixStore<row_major>(&buffer, 0, subgroup_matrix_left<f32, 8, 8>(), 64);
subgroupMatrixStore<row_major>(&buffer, 0, array<subgroup_matrix_left<f32, 8, 8>, 4>()[1], 64);
subgroupMatrixStore<row_major>(&buffer, 0, array<array<subgroup_matrix_left<f32, 8, 8>, 4>, 4>()[2][3], 64);
subgroupMatrixStore<row_major>(&buffer, 0, S().l, 64);
subgroupMatrixStore<row_major>(&buffer, 0, S_Nested().s.r, 64);
// Non-zero init.
subgroupMatrixStore<row_major>(&buffer, 0, subgroup_matrix_left<f32, 8, 8>(42), 64);
subgroupMatrixStore<row_major>(&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], 64);
subgroupMatrixStore<row_major>(&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], 64);
subgroupMatrixStore<row_major>(&buffer, 0, S(
subgroup_matrix_left<f32, 8, 8>(42),
subgroup_matrix_right<f32, 8, 8>(100),
).l, 64);
subgroupMatrixStore<row_major>(&buffer, 0, S_Nested(
S(
subgroup_matrix_left<f32, 8, 8>(42),
subgroup_matrix_right<f32, 8, 8>(100),
),
).s.r, 64);
}