Updated version of piet-gpu#199 barrier test with explicit zeroing

barrier_test.wgsl

@group(0) @binding(0)
var<storage, read_write> bump: atomic<u32>;

@group(0) @binding(1)
var<storage, read_write> paths: array<u32>;

let WG_SIZE = 256u;

var<workgroup> sh_tile_offset: u32;

@compute @workgroup_size(256)
fn main(
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(local_invocation_id) local_id: vec3<u32>,
) {
    sh_tile_offset = 0u;
    workgroupBarrier();
    let drawobj_ix = global_id.x;
    let tile_count_in = local_id.x + 1u;
    if local_id.x == WG_SIZE - 1u {
        sh_tile_offset = 1u + atomicAdd(&bump, tile_count_in);
    }
    workgroupBarrier();
    let tile_offset = sh_tile_offset;
    if drawobj_ix < 3u {
        paths[drawobj_ix] = tile_offset;
    }

}