blob: 4aa038da18d31733de468c804f8395b50465d13a [file] [log] [blame]
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();
var 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);
}