prefix_sum_sub.wgsl.h

Go to the documentation of this file.

#pragma once
 
#include <string>
 
shader_source = R"(
// Define the binding group layout
@group(0) @binding(0) var<storage, read_write> data: array<u32>;
@group(0) @binding(1) var<storage, read_write> result: array<u32>;
@group(0) @binding(2) var<storage, read_write> aux: array<u32>;
 
// Shared memory for the workgroup
const WORKGROUP_SIZE : u32 = 256u;
var<workgroup> shared_data: array<u32, WORKGROUP_SIZE>;
 
@compute @workgroup_size(WORKGROUP_SIZE)
fn PrefixSumSub(
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(local_invocation_id) local_id: vec3<u32>,
    @builtin(workgroup_id) group_id: vec3<u32>
) {
    // ids
    let tid = local_id.x;
    let gid = global_id.x;
 
    // Load data into shared memory
    if gid < arrayLength(&data) {
        shared_data[tid] = data[gid];
    }
    workgroupBarrier();
 
    // Up-sweep (reduce) phase
    var offset = 1u;
    while offset < WORKGROUP_SIZE {
        let idx = (tid + 1u) * offset * 2u - 1u;
        if idx < WORKGROUP_SIZE {
            shared_data[idx] += shared_data[idx - offset];
        }
        offset *= 2u;
        workgroupBarrier();
    }
 
    // Clear the last element
    if tid == 0u {
        shared_data[WORKGROUP_SIZE - 1u] = 0u;
    }
    workgroupBarrier();
 
    // Down-sweep phase
    offset = WORKGROUP_SIZE / 2u;
    while offset > 0u {
        let idx = (tid + 1u) * offset * 2u - 1u;
        if idx < WORKGROUP_SIZE {
            let t = shared_data[idx - offset];
            shared_data[idx - offset] = shared_data[idx];
            shared_data[idx] += t;
        }
        offset /= 2u;
        workgroupBarrier();
    }
 
    // Write the result back to the global memory
    if gid < arrayLength(&data) {
        result[gid] = shared_data[tid];
        if tid == (WORKGROUP_SIZE - 1u) {
            aux[group_id.x] = data[gid] + result[gid];
        }
    }
}
)";