engineComputeShaders top-level constant

Implementation

const Map<String, WebGpuComputeStage> engineComputeShaders =
    <String, WebGpuComputeStage>{
      'PrefixSum': WebGpuComputeStage(
        wgsl: r'''
struct Values {
    values: array<u32>,
}

var<private> gl_LocalInvocationID_1: vec3<u32>;
@group(0) @binding(0)
var<storage, read_write> values_buffer: Values;
var<workgroup> partial: array<u32, 256>;

fn main_1() {
    var i: u32;
    var base: u32;
    var s0_: u32;
    var s1_: u32;
    var s2_: u32;
    var s3_: u32;
    var offset: u32;
    var add: u32;
    var local: u32;
    var before: u32;
    var local_1: u32;

    let _e24 = gl_LocalInvocationID_1[0u];
    i = _e24;
    let _e25 = i;
    base = (_e25 * 4u);
    let _e27 = base;
    let _e30 = values_buffer.values[_e27];
    s0_ = _e30;
    let _e31 = s0_;
    let _e32 = base;
    let _e36 = values_buffer.values[(_e32 + 1u)];
    s1_ = (_e31 + _e36);
    let _e38 = s1_;
    let _e39 = base;
    let _e43 = values_buffer.values[(_e39 + 2u)];
    s2_ = (_e38 + _e43);
    let _e45 = s2_;
    let _e46 = base;
    let _e50 = values_buffer.values[(_e46 + 3u)];
    s3_ = (_e45 + _e50);
    let _e52 = i;
    let _e53 = s3_;
    partial[_e52] = _e53;
    workgroupBarrier();
    offset = 1u;
    loop {
        let _e55 = offset;
        if (_e55 < 256u) {
            let _e57 = i;
            let _e58 = offset;
            if (_e57 >= _e58) {
                let _e60 = i;
                let _e61 = offset;
                let _e64 = partial[(_e60 - _e61)];
                local = _e64;
            } else {
                local = 0u;
            }
            let _e65 = local;
            add = _e65;
            workgroupBarrier();
            let _e66 = i;
            let _e67 = add;
            let _e69 = partial[_e66];
            partial[_e66] = (_e69 + _e67);
            workgroupBarrier();
            continue;
        } else {
            break;
        }
        continuing {
            let _e72 = offset;
            offset = (_e72 << bitcast<u32>(1u));
        }
    }
    let _e75 = i;
    if (_e75 == 0u) {
        local_1 = 0u;
    } else {
        let _e77 = i;
        let _e80 = partial[(_e77 - 1u)];
        local_1 = _e80;
    }
    let _e81 = local_1;
    before = _e81;
    let _e82 = base;
    let _e83 = before;
    let _e84 = s0_;
    values_buffer.values[_e82] = (_e83 + _e84);
    let _e88 = base;
    let _e90 = before;
    let _e91 = s1_;
    values_buffer.values[(_e88 + 1u)] = (_e90 + _e91);
    let _e95 = base;
    let _e97 = before;
    let _e98 = s2_;
    values_buffer.values[(_e95 + 2u)] = (_e97 + _e98);
    let _e102 = base;
    let _e104 = before;
    let _e105 = s3_;
    values_buffer.values[(_e102 + 3u)] = (_e104 + _e105);
    return;
}

@compute @workgroup_size(256, 1, 1)
fn main(@builtin(local_invocation_id) gl_LocalInvocationID: vec3<u32>) {
    gl_LocalInvocationID_1 = gl_LocalInvocationID;
    main_1();
}
''',
        workgroupSize: (256, 1, 1),
        storage: <WebGpuStorageBinding>[
          WebGpuStorageBinding(
            name: 'Values',
            group: 0,
            binding: 0,
            readOnly: false,
          ),
        ],
        blocks: <WebGpuBlock>[],
      ),
    };