xref: /aosp_15_r20/external/skia/tests/sksl/compute/AtomicOperationsOverArrayAndStruct.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  globalCounts: GlobalCounts,
8};
9@group(0) @binding(0) var<storage, read_write> _storage0 : ssbo;
10struct GlobalCounts {
11  firstHalfCount: atomic<u32>,
12  secondHalfCount: atomic<u32>,
13};
14var<workgroup> localCounts: array<atomic<u32>, 2>;
15fn _skslMain(_stageIn: CSIn) {
16  {
17    if _stageIn.sk_LocalInvocationID.x == 0u {
18      {
19        atomicStore(&localCounts[0], 0u);
20        atomicStore(&localCounts[1], 0u);
21      }
22    }
23    workgroupBarrier();
24    let idx: u32 = u32(select(1, 0, _stageIn.sk_LocalInvocationID.x < 128u));
25    let _skTemp1 = atomicAdd(&localCounts[idx], 1u);
26    workgroupBarrier();
27    if _stageIn.sk_LocalInvocationID.x == 0u {
28      {
29        let _skTemp2 = atomicLoad(&localCounts[0]);
30        let _skTemp3 = atomicAdd(&_storage0.globalCounts.firstHalfCount, _skTemp2);
31        let _skTemp4 = atomicLoad(&localCounts[1]);
32        let _skTemp5 = atomicAdd(&_storage0.globalCounts.secondHalfCount, _skTemp4);
33      }
34    }
35  }
36}
37@compute @workgroup_size(256, 1, 1) fn main(_stageIn: CSIn) {
38  _skslMain(_stageIn);
39}
40