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