2022-06-24 22:28:23 +00:00
|
|
|
struct S0_atomic {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : atomic<u32>,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(8) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(12) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
|
|
|
}
|
|
|
|
|
|
|
|
struct S0 {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : u32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(8) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(12) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
|
|
|
}
|
|
|
|
|
|
|
|
struct S1_atomic {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : S0_atomic,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(20) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(24) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
|
|
|
}
|
|
|
|
|
|
|
|
struct S1 {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : S0,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(20) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(24) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
|
|
|
}
|
|
|
|
|
|
|
|
struct S2_atomic {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(8) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(12) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : S1_atomic,
|
|
|
|
}
|
|
|
|
|
|
|
|
struct S2 {
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(0) */
|
2022-06-24 22:28:23 +00:00
|
|
|
x : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(4) */
|
2022-06-24 22:28:23 +00:00
|
|
|
y : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(8) */
|
2022-06-24 22:28:23 +00:00
|
|
|
z : i32,
|
2022-12-01 17:37:56 +00:00
|
|
|
/* @offset(12) */
|
2022-06-24 22:28:23 +00:00
|
|
|
a : S1,
|
|
|
|
}
|
|
|
|
|
|
|
|
var<private> local_invocation_index_1 : u32;
|
|
|
|
|
|
|
|
var<workgroup> wg : S2_atomic;
|
|
|
|
|
2023-02-21 17:33:54 +00:00
|
|
|
fn compute_main_inner(local_invocation_index_2 : u32) {
|
2022-06-24 22:28:23 +00:00
|
|
|
wg.x = 0i;
|
|
|
|
wg.y = 0i;
|
|
|
|
wg.z = 0i;
|
|
|
|
wg.a.x = 0i;
|
|
|
|
wg.a.a.x = 0i;
|
|
|
|
atomicStore(&(wg.a.a.a), 0u);
|
|
|
|
wg.a.a.y = 0i;
|
|
|
|
wg.a.a.z = 0i;
|
|
|
|
wg.a.y = 0i;
|
|
|
|
wg.a.z = 0i;
|
|
|
|
workgroupBarrier();
|
|
|
|
atomicStore(&(wg.a.a.a), 1u);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
fn compute_main_1() {
|
|
|
|
let x_44 : u32 = local_invocation_index_1;
|
|
|
|
compute_main_inner(x_44);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
@compute @workgroup_size(1i, 1i, 1i)
|
|
|
|
fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
|
|
|
|
local_invocation_index_1 = local_invocation_index_1_param;
|
|
|
|
compute_main_1();
|
|
|
|
}
|