xref: /aosp_15_r20/external/skia/tests/sksl/compute/AtomicOperations.wgsl (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)
1diagnostic(off, derivative_uniformity);
2diagnostic(off, chromium.unreachable_code);
3struct CSIn {
4  @builtin(local_invocation_id) sk_LocalInvocationID: vec3<u32>,
5};
6struct ssbo {
7  globalCounter: atomic<u32>,
8};
9@group(0) @binding(0) var<storage, read_write> _storage0 : ssbo;
10var<workgroup> localCounter: atomic<u32>;
11fn _skslMain(_stageIn: CSIn) {
12  {
13    if _stageIn.sk_LocalInvocationID.x == 0u {
14      {
15        atomicStore(&localCounter, 0u);
16      }
17    }
18    workgroupBarrier();
19    let _skTemp1 = atomicAdd(&localCounter, 1u);
20    workgroupBarrier();
21    if _stageIn.sk_LocalInvocationID.x == 0u {
22      {
23        let _skTemp2 = atomicLoad(&localCounter);
24        let _skTemp3 = atomicAdd(&_storage0.globalCounter, _skTemp2);
25      }
26    }
27  }
28}
29@compute @workgroup_size(64, 1, 1) fn main(_stageIn: CSIn) {
30  _skslMain(_stageIn);
31}
32