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(1) @binding(0) var<storage, read_write> data_0: array<u32>;
@group(1) @binding(1) var<storage, read_write> result_0: array<u32>;
@group(1) @binding(2) var<storage, read_write> aux_0: array<u32>;
 
@group(1) @binding(3) var<storage, read_write> data_1: array<u32>;
@group(1) @binding(4) var<storage, read_write> result_1: array<u32>;
@group(1) @binding(5) var<storage, read_write> aux_1: array<u32>;
 
// Shared memory for the workgroup
const WORKGROUP_SIZE : u32 = 256u;
var<workgroup> shared_data_0: array<u32, WORKGROUP_SIZE>;
var<workgroup> shared_data_1: 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_0) {
        shared_data_0[tid] = data_0[gid];
        shared_data_1[tid] = data_1[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_0[idx] += shared_data_0[idx - offset];
            shared_data_1[idx] += shared_data_1[idx - offset];
        }
        offset *= 2u;
        workgroupBarrier();
    }
 
    // Clear the last element
    if tid == 0u {
        shared_data_0[WORKGROUP_SIZE - 1u] = 0u;
        shared_data_1[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_0 = shared_data_0[idx - offset];
            shared_data_0[idx - offset] = shared_data_0[idx];
            shared_data_0[idx] += t_0;
 
            let t_1 = shared_data_1[idx - offset];
            shared_data_1[idx - offset] = shared_data_1[idx];
            shared_data_1[idx] += t_1;
        }
        offset /= 2u;
        workgroupBarrier();
    }
 
    // Write the result back to the global memory
    if gid < arrayLength(&data_0) {
        result_0[gid] = shared_data_0[tid];
        result_1[gid] = shared_data_1[tid];
        
        if tid == (WORKGROUP_SIZE - 1u) {
            aux_0[group_id.x] = data_0[gid] + result_0[gid];
            aux_1[group_id.x] = data_1[gid] + result_1[gid];
        }
    }
}
)";