xref: /aosp_15_r20/external/skia/tests/sksl/compute/AtomicOperationsOverArrayAndStruct.wgsl (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)

diagnostic(off, derivative_uniformity);
diagnostic(off, chromium.unreachable_code);
struct CSIn {
  @builtin(local_invocation_id) sk_LocalInvocationID: vec3<u32>,
};
struct ssbo {
  globalCounts: GlobalCounts,
};
@group(0) @binding(0) var<storage, read_write> _storage0 : ssbo;
struct GlobalCounts {
  firstHalfCount: atomic<u32>,
  secondHalfCount: atomic<u32>,
};
var<workgroup> localCounts: array<atomic<u32>, 2>;
fn _skslMain(_stageIn: CSIn) {
  {
    if _stageIn.sk_LocalInvocationID.x == 0u {
      {
        atomicStore(&localCounts[0], 0u);
        atomicStore(&localCounts[1], 0u);
      }
    }
    workgroupBarrier();
    let idx: u32 = u32(select(1, 0, _stageIn.sk_LocalInvocationID.x < 128u));
    let _skTemp1 = atomicAdd(&localCounts[idx], 1u);
    workgroupBarrier();
    if _stageIn.sk_LocalInvocationID.x == 0u {
      {
        let _skTemp2 = atomicLoad(&localCounts[0]);
        let _skTemp3 = atomicAdd(&_storage0.globalCounts.firstHalfCount, _skTemp2);
        let _skTemp4 = atomicLoad(&localCounts[1]);
        let _skTemp5 = atomicAdd(&_storage0.globalCounts.secondHalfCount, _skTemp4);
      }
    }
  }
}
@compute @workgroup_size(256, 1, 1) fn main(_stageIn: CSIn) {
  _skslMain(_stageIn);
}