/**
 * chiral_sidon_check.wgsl — Chiral CRT Sidon Filter on GPU
 *
 * Uses the existing dna_braid.wgsl infrastructure:
 *   - braid_cross() = chiral crossing (over/under)
 *   - eigensolid_check() = Sidon convergence
 *   - Workgroup 256 = 2^8 chiral configurations
 *
 * This kernel checks ALL 256 chiral configurations simultaneously:
 *   1. Each thread = one chiral configuration (binary vector of length k)
 *   2. Each thread computes CRT embedding for its chiral config
 *   3. Each thread checks Sidon property (pairwise sums distinct)
 *   4. Results written to output buffer
 *
 * This is the GPU-accelerated Stage 6 (Sidon filter) of the six-stage pipeline.
 * With 256 threads, all configurations are checked in ONE dispatch.
 *
 * Connection to HCMR:
 *   - self_loop_prob = fraction of threads that fail Sidon (collision)
 *   - throughput = (1 - self_loop_prob) × base_rate
 *   - ring dispatch (all pass) = 0% collisions = maximum throughput
 */

const WORKGROUP_SIZE: u32 = 256u;  // 2^8 = 256 chiral configs
const N_LABELS: u32 = 8u;          // 8 Sidon labels (powers of 2)
const N_MODULI: u32 = 9u;          // L0 + 8 reflection moduli

@group(0) @binding(0) var<storage, read> labels: array<u32, 8>;        // Sidon labels
@group(0) @binding(1) var<storage, read> moduli: array<u32, 9>;        // CRT moduli
@group(0) @binding(2) var<uniform> params: ChiralParams;
@group(0) @binding(3) var<storage, read_write> results: array<u32>;    // output: 1=Sidon, 0=collision
@group(0) @binding(4) var<storage, read_write> scores: array<f32>;    // Sidon score

struct ChiralParams {
    S: u32,              // reflection point
    n_configs: u32,      // 256
    _pad: u32,
    _pad2: u32,
};

/// Compute CRT reconstruction for a single label with chiral config.
/// config_bits: each bit j determines whether axis j is flipped.
fn crt_embed_chiral(a: u32, config_bits: u32) -> u32 {
    // Identity component (always standard)
    var val: u32 = a % moduli[0];
    
    // For CRT reconstruction, we need all residues
    // Then use CRT to reconstruct the value mod M = prod(moduli)
    // For now: just compute the sum of residues as a hash
    // (full CRT reconstruction would be done on CPU or in a more complex kernel)
    
    var hash: u32 = val;
    for (var j: u32 = 1u; j < N_MODULI; j = j + 1u) {
        var residue: u32;
        if ((config_bits >> (j - 1u)) & 1u == 0u) {
            residue = (params.S - a) % moduli[j];  // standard
        } else {
            residue = (a - params.S) % moduli[j];  // flipped (chiral)
        }
        // Simple hash: multiply and add (not full CRT, but collision-detectable)
        hash = hash * moduli[j] + residue;
    }
    return hash;
}

/// Check Sidon property: all pairwise sums distinct.
/// Each thread checks one chiral configuration.
@compute @workgroup_size(WORKGROUP_SIZE)
fn chiral_sidon_check(@builtin(global_invocation_id) gid: vec3<u32>) {
    let config_id = gid.x;
    if (config_id >= params.n_configs) {
        return;
    }
    
    // Compute hash for each label under this chiral configuration
    var hashes: array<u32, 8>;
    for (var i: u32 = 0u; i < N_LABELS; i = i + 1u) {
        hashes[i] = crt_embed_chiral(labels[i], config_id);
    }
    
    // Check all pairwise sums for collisions
    var collisions: u32 = 0u;
    var total_pairs: u32 = 0u;
    
    for (var i: u32 = 0u; i < N_LABELS; i = i + 1u) {
        for (var j: u32 = i; j < N_LABELS; j = j + 1u) {
            let sum_ij = hashes[i] + hashes[j];
            total_pairs = total_pairs + 1u;
            
            // Check against all other pairs
            for (var k: u32 = 0u; k < N_LABELS; k = k + 1u) {
                for (var l: u32 = k; l < N_LABELS; l = l + 1u) {
                    if (i * 8u + j < k * 8u + l) {  // avoid double-counting
                        let sum_kl = hashes[k] + hashes[l];
                        if (sum_ij == sum_kl) {
                            collisions = collisions + 1u;
                        }
                    }
                }
            }
        }
    }
    
    // Write results
    results[config_id] = select(1u, 0u, collisions > 0u);
    scores[config_id] = 1.0 - f32(collisions) / f32(total_pairs);
}

/// COUCH gate: check if self-loop (contention) is below threshold.
/// Uses the number of under-crossings as proxy for contention.
@compute @workgroup_size(WORKGROUP_SIZE)
fn couch_stability_check(@builtin(global_invocation_id) gid: vec3<u32>) {
    let config_id = gid.x;
    if (config_id >= params.n_configs) {
        return;
    }
    
    // Count under-crossings (number of 1 bits in config_id)
    var under_count: u32 = 0u;
    var bits = config_id;
    for (var i: u32 = 0u; i < 8u; i = i + 1u) {
        under_count = under_count + (bits & 1u);
        bits = bits >> 1u;
    }
    
    // Self-loop proxy: more under-crossings = more contention
    // 0 under = ring dispatch (0.0), 8 under = AVX-512 (0.885)
    let self_loop = 57942u * under_count / 8u;  // Q16_16 raw
    
    // COUCH stable if self_loop < threshold (49152 = 0.75)
    results[config_id + params.n_configs] = select(1u, 0u, self_loop >= 49152u);
}