Constant TOPOLOGY_BOUNDARY_MATRIX 

pub const TOPOLOGY_BOUNDARY_MATRIX: &str = r#"
struct Simplex {
    vertices: array<u32, 8>,  // Max 7-simplex
    dimension: u32,
    filtration_time: f32,
    padding: array<u32, 2>,
}

struct MatrixEntry {
    row: u32,
    col: u32,
    value: i32,
    padding: u32,
}

@group(0) @binding(0)
var<storage, read> simplices: array<Simplex>;

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

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

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let simplex_idx = global_id.x;
    if (simplex_idx >= arrayLength(&simplices)) {
        return;
    }

    let s = simplices[simplex_idx];
    if (s.dimension == 0u) {
        return;  // 0-simplices have no boundary
    }

    // Generate boundary faces with alternating signs
    let dim = s.dimension;
    for (var i = 0u; i <= dim; i++) {
        let sign = select(-1, 1, i % 2u == 0u);

        // Allocate entry atomically
        let entry_idx = atomicAdd(&entry_counter, 1u);

        // Compute hash of face (for row index)
        var face_hash = 0u;
        for (var j = 0u; j <= dim; j++) {
            if (j != i) {
                face_hash = face_hash * 31u + s.vertices[j];
            }
        }

        boundary_entries[entry_idx] = MatrixEntry(face_hash, simplex_idx, sign, 0u);
    }
}
"#;

Boundary matrix construction for simplicial complex (sparse format)