mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-10-25 19:20:30 +00:00
268d7b8357
The following operations are supported: OpAtomicLoad OpAtomicStore OpAtomicExchange OpAtomicCompareExchange OpAtomicCompareExchangeWeak OpAtomicIIncrement OpAtomicIDecrement OpAtomicIAdd OpAtomicISub OpAtomicSMin OpAtomicUMin OpAtomicSMax OpAtomicUMax OpAtomicAnd OpAtomicOr OpAtomicXor These are not, but may be supported in the future: OpAtomicFlagTestAndSet OpAtomicFlagClear OpAtomicFMinEXT OpAtomicFMaxEXT OpAtomicFAddEXT Bug: tint:1441 Change-Id: Ifd53643b38d43664905a0dddfca609add4914670 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94121 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Ben Clayton <bclayton@google.com> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
45 lines
1007 B
WebGPU Shading Language
45 lines
1007 B
WebGPU Shading Language
type Arr = array<u32, 1u>;
|
|
|
|
type Arr_1 = array<Arr, 2u>;
|
|
|
|
type Arr_2 = array<Arr_1, 3u>;
|
|
|
|
var<private> local_invocation_index_1 : u32;
|
|
|
|
var<workgroup> wg : array<array<array<atomic<u32>, 1u>, 2u>, 3u>;
|
|
|
|
fn compute_main_inner(local_invocation_index : u32) {
|
|
var idx : u32 = 0u;
|
|
idx = local_invocation_index;
|
|
loop {
|
|
let x_25 : u32 = idx;
|
|
if (!((x_25 < 6u))) {
|
|
break;
|
|
}
|
|
let x_31 : u32 = idx;
|
|
let x_33 : u32 = idx;
|
|
let x_35 : u32 = idx;
|
|
atomicStore(&(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u);
|
|
|
|
continuing {
|
|
let x_42 : u32 = idx;
|
|
idx = (x_42 + 1u);
|
|
}
|
|
}
|
|
workgroupBarrier();
|
|
atomicStore(&(wg[2i][1i][0i]), 1u);
|
|
return;
|
|
}
|
|
|
|
fn compute_main_1() {
|
|
let x_57 : u32 = local_invocation_index_1;
|
|
compute_main_inner(x_57);
|
|
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();
|
|
}
|