facett-core 0.1.15

//! **GPU compute particles** (feature `wgpu`) — a storage-buffer boids system
//! stepped in a compute shader and drawn as soft **additive** points into an HDR
//! target. Net-new (the only prior compute shader in the tree is the 2D map cull);
//! it composes with the shared HDR offscreen colour by drawing additively over it.
//!
//! Reynolds' three flocking rules (separation / alignment / cohesion) run O(n²) in
//! `particles.wgsl`'s `cs_step` over a ping-pong pair of storage buffers (read
//! `src`, write `dst` — no read/write hazard). The sim is a **pure function** of the
//! initial state + dt + step count, so it is deterministic (FC-7) and the CPU
//! mirror ([`step_cpu`]) is the parity reference the GPU readback proof checks
//! against.
//!
//! The render pass ([`GpuParticles::render`]) expands a unit quad per particle with
//! a soft radial falloff and premultiplied output, so an additive blend accumulates
//! glow — overlapping particles get brighter, the classic light-field look.

use wgpu::util::DeviceExt;
use wgpu::TextureFormat;

/// One particle: 2D position (normalised `[0,1)` toroidal space) + velocity.
/// 16 bytes, matches `Particle` in `particles.wgsl`.
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
pub struct Particle {
    pub pos: [f32; 2],
    pub vel: [f32; 2],
}

/// The boids tuning. [`Default`] is a calm, cohesive flock.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct SimParams {
    pub max_speed: f32,
    pub sep_radius: f32,
    pub align_radius: f32,
    pub cohesion_radius: f32,
    pub sep_weight: f32,
    pub align_weight: f32,
    pub cohesion_weight: f32,
}

impl Default for SimParams {
    fn default() -> Self {
        Self {
            max_speed: 0.4,
            sep_radius: 0.04,
            align_radius: 0.10,
            cohesion_radius: 0.10,
            sep_weight: 1.5,
            align_weight: 1.0,
            cohesion_weight: 0.8,
        }
    }
}

/// How the additive points are drawn.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct DrawParams {
    /// Point sprite size in physical pixels.
    pub point_size: f32,
    /// Glow intensity multiplier (can exceed 1.0 into the HDR target).
    pub intensity: f32,
    /// Base (cool/slow) tint, linear RGB.
    pub tint: [f32; 3],
    /// How strongly speed warms the colour (0 = constant tint).
    pub speed_warm: f32,
}

impl Default for DrawParams {
    fn default() -> Self {
        Self { point_size: 6.0, intensity: 1.0, tint: [0.25, 0.55, 1.0], speed_warm: 2.0 }
    }
}

/// The CPU reference step — advance `particles` by `dt` under `p`, **mirroring**
/// `cs_step` in `particles.wgsl` (same loop order, same math). Pure + deterministic
/// (FC-7): the GPU readback proof asserts the device produces the same result.
pub fn step_cpu(particles: &[Particle], dt: f32, p: &SimParams) -> Vec<Particle> {
    let n = particles.len();
    let mut out = Vec::with_capacity(n);
    for i in 0..n {
        let me = particles[i];
        let mut sep = [0.0f32; 2];
        let mut align_sum = [0.0f32; 2];
        let mut coh_sum = [0.0f32; 2];
        let mut align_n = 0.0f32;
        let mut coh_n = 0.0f32;
        for (j, o) in particles.iter().enumerate() {
            if j == i {
                continue;
            }
            let d = [o.pos[0] - me.pos[0], o.pos[1] - me.pos[1]];
            let dist = (d[0] * d[0] + d[1] * d[1]).sqrt();
            if dist > 0.0 && dist < p.sep_radius {
                let inv = 1.0 / (dist * dist);
                sep[0] -= d[0] * inv;
                sep[1] -= d[1] * inv;
            }
            if dist < p.align_radius {
                align_sum[0] += o.vel[0];
                align_sum[1] += o.vel[1];
                align_n += 1.0;
            }
            if dist < p.cohesion_radius {
                coh_sum[0] += o.pos[0];
                coh_sum[1] += o.pos[1];
                coh_n += 1.0;
            }
        }
        let mut acc = [sep[0] * p.sep_weight, sep[1] * p.sep_weight];
        if align_n > 0.0 {
            acc[0] += (align_sum[0] / align_n - me.vel[0]) * p.align_weight;
            acc[1] += (align_sum[1] / align_n - me.vel[1]) * p.align_weight;
        }
        if coh_n > 0.0 {
            acc[0] += (coh_sum[0] / coh_n - me.pos[0]) * p.cohesion_weight;
            acc[1] += (coh_sum[1] / coh_n - me.pos[1]) * p.cohesion_weight;
        }
        let mut vel = [me.vel[0] + acc[0] * dt, me.vel[1] + acc[1] * dt];
        let sp = (vel[0] * vel[0] + vel[1] * vel[1]).sqrt();
        if sp > p.max_speed {
            let k = p.max_speed / sp;
            vel[0] *= k;
            vel[1] *= k;
        }
        let mut pos = [me.pos[0] + vel[0] * dt, me.pos[1] + vel[1] * dt];
        pos[0] -= pos[0].floor();
        pos[1] -= pos[1].floor();
        out.push(Particle { pos, vel });
    }
    out
}

/// The CPU reference **advect** — advance each particle's arc param `s` (packed in
/// `vel.x`) by `speed·dt` (`vel.y`), wrap to `[0,1)`, and resample its edge
/// poly-line for the new `pos`. Line-for-line mirror of `cs_advect` in
/// `particles.wgsl` (`fract` + [`crate::edges::point_at_frac`], the same arc-length
/// walk the GPU `sample_polyline` twins). Pure + deterministic (FC-7): the readback
/// proof asserts the device produces the same result. `polylines[part_edge[i]]` is
/// particle `i`'s edge in normalised screen space.
pub fn advect_cpu(particles: &[Particle], polylines: &[Vec<[f32; 2]>], part_edge: &[u32], dt: f32) -> Vec<Particle> {
    use egui::pos2;
    particles
        .iter()
        .enumerate()
        .map(|(i, me)| {
            let raw = me.vel[0] + me.vel[1] * dt;
            let s = raw - raw.floor(); // fract == wrap [0,1), matches WGSL
            let e = part_edge.get(i).copied().unwrap_or(0) as usize;
            let pos = match polylines.get(e) {
                Some(pl) if !pl.is_empty() => {
                    let pts: Vec<egui::Pos2> = pl.iter().map(|p| pos2(p[0], p[1])).collect();
                    let at = crate::edges::point_at_frac(&pts, s);
                    [at.x, at.y]
                }
                _ => me.pos,
            };
            Particle { pos, vel: [s, me.vel[1]] }
        })
        .collect()
}

#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct SimUniforms {
    a: [f32; 4],
    b: [f32; 4],
    c: [f32; 4],
}

#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct DrawUniforms {
    a: [f32; 4],
    tint: [f32; 4],
}

/// The `cs_advect` uniform (`AdvectU` in `particles.wgsl`), 16 bytes. `a = [dt,
/// count, _, _]`.
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct AdvectUniforms {
    a: [f32; 4],
}

/// The per-frame edge geometry for the advection pass, in CSR form: `offsets`
/// (per-edge start into `points`, len `n_edges + 1`), the concatenated normalised
/// poly-lines `points`, and the per-particle `part_edge` id. Built once per DAG
/// layout (or when the camera moves) and reused across advect steps — mirrors
/// [`GpuParticles::make_buffers`]. Screen-space is normalised `[0,1]` so the shared
/// `pt_vs` maps it straight to NDC.
pub struct FlowGeometry {
    offsets: wgpu::Buffer,
    points: wgpu::Buffer,
    part_edge: wgpu::Buffer,
}

impl FlowGeometry {
    /// Pack `polylines` (per-edge normalised screen poly-lines) + `part_edge`
    /// (per-particle edge id) into GPU storage buffers. Empty poly-lines are kept
    /// as zero-length spans so edge ids stay stable.
    pub fn new(device: &wgpu::Device, polylines: &[Vec<[f32; 2]>], part_edge: &[u32]) -> Self {
        let mut offsets: Vec<u32> = Vec::with_capacity(polylines.len() + 1);
        let mut points: Vec<[f32; 2]> = Vec::new();
        offsets.push(0);
        for pl in polylines {
            points.extend_from_slice(pl);
            offsets.push(points.len() as u32);
        }
        // create_buffer_init rejects empty contents — pad to one element.
        if points.is_empty() {
            points.push([0.0, 0.0]);
        }
        let pe: Vec<u32> = if part_edge.is_empty() { vec![0] } else { part_edge.to_vec() };
        let mk = |label: &str, bytes: &[u8]| {
            device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
                label: Some(label),
                contents: bytes,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            })
        };
        Self {
            offsets: mk("l0_flow_offsets", bytemuck::cast_slice(&offsets)),
            points: mk("l0_flow_points", bytemuck::cast_slice(&points)),
            part_edge: mk("l0_flow_part_edge", bytemuck::cast_slice(&pe)),
        }
    }
}

/// The compute shader (`cs_step`) + the additive point render (`pt_vs`/`pt_fs`).
pub const PARTICLES_WGSL: &str = include_str!("particles.wgsl");

/// A ping-pong pair of particle storage buffers + the live count. After each
/// [`GpuParticles::step`] the up-to-date particles are in `a`.
pub struct ParticleBuffers {
    a: wgpu::Buffer,
    b: wgpu::Buffer,
    count: u32,
}

impl ParticleBuffers {
    /// The live particle count.
    #[must_use]
    pub fn count(&self) -> u32 {
        self.count
    }
}

/// The compute-particles renderer: one compute pipeline (the boids step) + one
/// additive render pipeline (the point sprites), built for an HDR `target_format`.
pub struct GpuParticles {
    step_pipeline: wgpu::ComputePipeline,
    step_bgl: wgpu::BindGroupLayout,
    advect_pipeline: wgpu::ComputePipeline,
    advect_bgl: wgpu::BindGroupLayout,
    draw_pipeline: wgpu::RenderPipeline,
    draw_bgl: wgpu::BindGroupLayout,
    sim_u: wgpu::Buffer,
    advect_u: wgpu::Buffer,
    draw_u: wgpu::Buffer,
    target_format: TextureFormat,
}

impl GpuParticles {
    /// Build the compute + additive-render pipelines for `target_format` (e.g. the
    /// HDR `Rgba16Float` offscreen). Validates `particles.wgsl` on build. Requires a
    /// device created with compute + storage support.
    pub fn new(device: &wgpu::Device, target_format: TextureFormat) -> Self {
        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("l0_particles"),
            source: wgpu::ShaderSource::Wgsl(PARTICLES_WGSL.into()),
        });

        // ── compute step bind group: uniform + src(read) + dst(read_write) ──
        let step_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("l0_particles_step_bgl"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: true }, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 2,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: false }, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
            ],
        });
        let step_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
            label: Some("l0_particles_step"),
            layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("l0_particles_step_pll"),
                bind_group_layouts: &[Some(&step_bgl)],
                immediate_size: 0,
            })),
            module: &shader,
            entry_point: Some("cs_step"),
            compilation_options: Default::default(),
            cache: None,
        });

        // ── advect bind group: uniform + src(read) + dst(rw) + 3 CSR geom reads ──
        let storage = |binding: u32, read_only: bool| wgpu::BindGroupLayoutEntry {
            binding,
            visibility: wgpu::ShaderStages::COMPUTE,
            ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only }, has_dynamic_offset: false, min_binding_size: None },
            count: None,
        };
        let advect_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("l0_particles_advect_bgl"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
                storage(1, true),  // a_src
                storage(2, false), // a_dst
                storage(3, true),  // edge_off
                storage(4, true),  // edge_pts
                storage(5, true),  // part_edge
            ],
        });
        let advect_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
            label: Some("l0_particles_advect"),
            layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("l0_particles_advect_pll"),
                bind_group_layouts: &[Some(&advect_bgl)],
                immediate_size: 0,
            })),
            module: &shader,
            entry_point: Some("cs_advect"),
            compilation_options: Default::default(),
            cache: None,
        });

        // ── render bind group: uniform + points(read storage, vertex stage) ──
        let draw_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("l0_particles_draw_bgl"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::VERTEX,
                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: true }, has_dynamic_offset: false, min_binding_size: None },
                    count: None,
                },
            ],
        });
        // Additive: premultiplied source added onto the target (glow accumulates).
        let blend = Some(wgpu::BlendState {
            color: wgpu::BlendComponent { src_factor: wgpu::BlendFactor::One, dst_factor: wgpu::BlendFactor::One, operation: wgpu::BlendOperation::Add },
            alpha: wgpu::BlendComponent { src_factor: wgpu::BlendFactor::One, dst_factor: wgpu::BlendFactor::One, operation: wgpu::BlendOperation::Add },
        });
        let draw_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
            label: Some("l0_particles_draw"),
            layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("l0_particles_draw_pll"),
                bind_group_layouts: &[Some(&draw_bgl)],
                immediate_size: 0,
            })),
            vertex: wgpu::VertexState { module: &shader, entry_point: Some("pt_vs"), compilation_options: Default::default(), buffers: &[] },
            primitive: wgpu::PrimitiveState { topology: wgpu::PrimitiveTopology::TriangleList, ..Default::default() },
            depth_stencil: None,
            multisample: wgpu::MultisampleState::default(),
            fragment: Some(wgpu::FragmentState {
                module: &shader,
                entry_point: Some("pt_fs"),
                compilation_options: Default::default(),
                targets: &[Some(wgpu::ColorTargetState { format: target_format, blend, write_mask: wgpu::ColorWrites::ALL })],
            }),
            multiview_mask: None,
            cache: None,
        });

        let mkbuf = |label: &str| device.create_buffer(&wgpu::BufferDescriptor {
            label: Some(label),
            size: std::mem::size_of::<SimUniforms>().max(std::mem::size_of::<DrawUniforms>()) as u64,
            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });
        Self {
            step_pipeline,
            step_bgl,
            advect_pipeline,
            advect_bgl,
            draw_pipeline,
            draw_bgl,
            sim_u: mkbuf("l0_particles_sim_u"),
            advect_u: mkbuf("l0_particles_advect_u"),
            draw_u: mkbuf("l0_particles_draw_u"),
            target_format,
        }
    }

    /// The colour format the render pipeline targets.
    #[must_use]
    pub fn target_format(&self) -> TextureFormat {
        self.target_format
    }

    /// Allocate the ping-pong storage buffers seeded with `init`.
    ///
    /// An **empty** `init` is legal and yields `count == 0` — a 0-workgroup dispatch and
    /// a 0-instance draw, both valid wgpu. `create_buffer_init` rejects empty contents, so
    /// the allocation is padded to one element exactly as [`FlowGeometry::new`] already
    /// pads its own (the two were inconsistent: `FlowGeometry` tolerated the empty case
    /// and this panicked on it, which quietly made "no particles" a landmine every caller
    /// had to guard against separately).
    pub fn make_buffers(&self, device: &wgpu::Device, init: &[Particle]) -> ParticleBuffers {
        let usage = wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST;
        let padded = if init.is_empty() { &[Particle { pos: [0.0, 0.0], vel: [0.0, 0.0] }][..] } else { init };
        let a = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
            label: Some("l0_particles_a"),
            contents: bytemuck::cast_slice(padded),
            usage,
        });
        let b = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("l0_particles_b"),
            size: (std::mem::size_of::<Particle>() * init.len().max(1)) as u64,
            usage,
            mapped_at_creation: false,
        });
        ParticleBuffers { a, b, count: init.len() as u32 }
    }

    /// Record one boids step: dispatch `cs_step` reading `bufs.a`, writing `bufs.b`,
    /// then swap so the result lands back in `bufs.a`.
    pub fn step(
        &self,
        device: &wgpu::Device,
        queue: &wgpu::Queue,
        encoder: &mut wgpu::CommandEncoder,
        bufs: &mut ParticleBuffers,
        dt: f32,
        p: &SimParams,
    ) {
        let u = SimUniforms {
            a: [dt, bufs.count as f32, p.max_speed, p.sep_radius],
            b: [p.align_radius, p.cohesion_radius, p.sep_weight, p.align_weight],
            c: [p.cohesion_weight, 0.0, 0.0, 0.0],
        };
        queue.write_buffer(&self.sim_u, 0, bytemuck::bytes_of(&u));
        let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("l0_particles_step_bind"),
            layout: &self.step_bgl,
            entries: &[
                wgpu::BindGroupEntry { binding: 0, resource: self.sim_u.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 1, resource: bufs.a.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 2, resource: bufs.b.as_entire_binding() },
            ],
        });
        {
            let mut cp = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("l0_particles_step_pass"), timestamp_writes: None });
            cp.set_pipeline(&self.step_pipeline);
            cp.set_bind_group(0, &bind, &[]);
            cp.dispatch_workgroups(bufs.count.div_ceil(64), 1, 1);
        }
        std::mem::swap(&mut bufs.a, &mut bufs.b);
    }

    /// Record one **edge advection** step: dispatch `cs_advect` reading `bufs.a`,
    /// writing `bufs.b` (each particle advances its arc param `s` by `speed·dt` and
    /// resamples its edge poly-line), then swap so the result lands in `bufs.a`. The
    /// host then draws the advected points with the **existing** [`Self::render`]
    /// (`load = true` composes over the DAG's HDR scene). CPU mirror is
    /// [`advect_cpu`] (FC-7 parity proof `gpu_advect_matches_cpu_reference`).
    pub fn advect(
        &self,
        device: &wgpu::Device,
        queue: &wgpu::Queue,
        encoder: &mut wgpu::CommandEncoder,
        bufs: &mut ParticleBuffers,
        geo: &FlowGeometry,
        dt: f32,
    ) {
        let u = AdvectUniforms { a: [dt, bufs.count as f32, 0.0, 0.0] };
        queue.write_buffer(&self.advect_u, 0, bytemuck::bytes_of(&u));
        let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("l0_particles_advect_bind"),
            layout: &self.advect_bgl,
            entries: &[
                wgpu::BindGroupEntry { binding: 0, resource: self.advect_u.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 1, resource: bufs.a.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 2, resource: bufs.b.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 3, resource: geo.offsets.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 4, resource: geo.points.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 5, resource: geo.part_edge.as_entire_binding() },
            ],
        });
        {
            let mut cp = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("l0_particles_advect_pass"), timestamp_writes: None });
            cp.set_pipeline(&self.advect_pipeline);
            cp.set_bind_group(0, &bind, &[]);
            cp.dispatch_workgroups(bufs.count.div_ceil(64), 1, 1);
        }
        std::mem::swap(&mut bufs.a, &mut bufs.b);
    }

    /// Record the additive point render of `bufs.a` into `target` (must be
    /// `target_format`). `load` keeps the existing target contents (compose over the
    /// HDR scene); `false` clears to transparent first.
    #[allow(clippy::too_many_arguments)]
    pub fn render(
        &self,
        device: &wgpu::Device,
        queue: &wgpu::Queue,
        encoder: &mut wgpu::CommandEncoder,
        bufs: &ParticleBuffers,
        target: &wgpu::TextureView,
        d: &DrawParams,
        load: bool,
        w: u32,
        h: u32,
    ) {
        let u = DrawUniforms {
            a: [d.point_size, d.intensity, w.max(1) as f32, h.max(1) as f32],
            tint: [d.tint[0], d.tint[1], d.tint[2], d.speed_warm],
        };
        queue.write_buffer(&self.draw_u, 0, bytemuck::bytes_of(&u));
        let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("l0_particles_draw_bind"),
            layout: &self.draw_bgl,
            entries: &[
                wgpu::BindGroupEntry { binding: 0, resource: self.draw_u.as_entire_binding() },
                wgpu::BindGroupEntry { binding: 1, resource: bufs.a.as_entire_binding() },
            ],
        });
        let load_op = if load { wgpu::LoadOp::Load } else { wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT) };
        let mut rp = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
            label: Some("l0_particles_draw_pass"),
            color_attachments: &[Some(wgpu::RenderPassColorAttachment {
                view: target,
                resolve_target: None,
                depth_slice: None,
                ops: wgpu::Operations { load: load_op, store: wgpu::StoreOp::Store },
            })],
            depth_stencil_attachment: None,
            timestamp_writes: None,
            occlusion_query_set: None,
            multiview_mask: None,
        });
        rp.set_pipeline(&self.draw_pipeline);
        rp.set_bind_group(0, &bind, &[]);
        rp.draw(0..6, 0..bufs.count);
    }
}

#[cfg(test)]
mod tests {
    use super::*;

    fn ring(n: usize) -> Vec<Particle> {
        // A deterministic seeded flock: positions on a small circle, tangential vel.
        (0..n)
            .map(|i| {
                let t = i as f32 / n as f32 * std::f32::consts::TAU;
                Particle { pos: [0.5 + 0.08 * t.cos(), 0.5 + 0.08 * t.sin()], vel: [-0.1 * t.sin(), 0.1 * t.cos()] }
            })
            .collect()
    }

    /// INJECT-ASSERT (FC-7): the CPU step is a pure, deterministic function —
    /// stepping the same state twice gives **bit-identical** results.
    #[test]
    fn cpu_step_is_deterministic() {
        let p = SimParams::default();
        let init = ring(32);
        let a = step_cpu(&init, 0.016, &p);
        let b = step_cpu(&init, 0.016, &p);
        assert_eq!(a.len(), init.len());
        for (x, y) in a.iter().zip(&b) {
            assert_eq!(x.pos[0].to_bits(), y.pos[0].to_bits(), "deterministic pos.x");
            assert_eq!(x.vel[1].to_bits(), y.vel[1].to_bits(), "deterministic vel.y");
        }
    }

    /// INJECT-ASSERT: positions stay inside the `[0,1)` toroidal box (the wrap
    /// works), and the speed clamp holds — no particle exceeds `max_speed`.
    #[test]
    fn step_keeps_bounds_and_clamps_speed() {
        let p = SimParams::default();
        let mut cur = ring(48);
        for _ in 0..60 {
            cur = step_cpu(&cur, 0.02, &p);
        }
        for q in &cur {
            assert!((0.0..1.0).contains(&q.pos[0]) && (0.0..1.0).contains(&q.pos[1]), "in box: {:?}", q.pos);
            let sp = (q.vel[0] * q.vel[0] + q.vel[1] * q.vel[1]).sqrt();
            assert!(sp <= p.max_speed + 1e-4, "speed clamped: {sp}");
        }
    }

    /// INJECT-ASSERT: separation actually pushes two near-coincident particles
    /// apart — the distance grows after a step (the rule does something real).
    #[test]
    fn separation_pushes_neighbours_apart() {
        let p = SimParams { align_weight: 0.0, cohesion_weight: 0.0, ..SimParams::default() };
        let init = vec![
            Particle { pos: [0.50, 0.5], vel: [0.0, 0.0] },
            Particle { pos: [0.51, 0.5], vel: [0.0, 0.0] },
        ];
        let d0 = (init[1].pos[0] - init[0].pos[0]).abs();
        let next = step_cpu(&init, 0.05, &p);
        let d1 = (next[1].pos[0] - next[0].pos[0]).abs();
        assert!(d1 > d0, "separation increased the gap ({d0} → {d1})");
    }

    /// Headless device with the adapter's real limits (so compute + storage are
    /// available), guarded on compute support. `None` → self-skip.
    fn compute_device() -> Option<(wgpu::Device, wgpu::Queue)> {
        let instance = wgpu::Instance::default();
        let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
            power_preference: wgpu::PowerPreference::default(),
            force_fallback_adapter: false,
            compatible_surface: None,
        }))
        .ok()?;
        if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
            return None;
        }
        if adapter.limits().max_storage_buffers_per_shader_stage < 2 {
            return None;
        }
        pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
            label: Some("l0-particles-proof"),
            required_features: wgpu::Features::empty(),
            required_limits: adapter.limits(),
            memory_hints: wgpu::MemoryHints::default(),
            experimental_features: wgpu::ExperimentalFeatures::disabled(),
            trace: wgpu::Trace::Off,
        }))
        .ok()
    }

    fn read_particles(device: &wgpu::Device, queue: &wgpu::Queue, buf: &wgpu::Buffer, n: u32) -> Vec<Particle> {
        let size = (std::mem::size_of::<Particle>() as u32 * n) as u64;
        let readback = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("l0-particles-readback"),
            size,
            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
            mapped_at_creation: false,
        });
        let mut enc = device.create_command_encoder(&Default::default());
        enc.copy_buffer_to_buffer(buf, 0, &readback, 0, size);
        queue.submit(Some(enc.finish()));
        let slice = readback.slice(..);
        let (tx, rx) = std::sync::mpsc::channel();
        slice.map_async(wgpu::MapMode::Read, move |r| { let _ = tx.send(r); });
        device.poll(wgpu::PollType::wait_indefinitely()).ok();
        rx.recv().unwrap().unwrap();
        let data = slice.get_mapped_range();
        let out: Vec<Particle> = bytemuck::cast_slice(&data).to_vec();
        drop(data);
        readback.unmap();
        out
    }

    /// RENDER/COMPUTE PARITY PROOF (self-skips without a compute device): one GPU
    /// boids step matches the CPU reference within float tolerance — proving the
    /// WGSL `cs_step` and the CPU `step_cpu` are the same function (FC-7 parity).
    #[test]
    fn gpu_step_matches_cpu_reference() {
        let Some((device, queue)) = compute_device() else {
            eprintln!("[particles] no compute device — skipping GPU parity proof");
            return;
        };
        let p = SimParams::default();
        let init = ring(40);
        let dt = 0.016f32;

        let gp = GpuParticles::new(&device, TextureFormat::Rgba16Float);
        let mut bufs = gp.make_buffers(&device, &init);
        let mut enc = device.create_command_encoder(&Default::default());
        gp.step(&device, &queue, &mut enc, &mut bufs, dt, &p);
        queue.submit(Some(enc.finish()));
        device.poll(wgpu::PollType::wait_indefinitely()).ok();

        let gpu = read_particles(&device, &queue, &bufs.a, bufs.count);
        let cpu = step_cpu(&init, dt, &p);
        assert_eq!(gpu.len(), cpu.len());
        let mut moved = false;
        for (g, c) in gpu.iter().zip(&cpu) {
            assert!((g.pos[0] - c.pos[0]).abs() < 1e-3, "pos.x parity: gpu {} vs cpu {}", g.pos[0], c.pos[0]);
            assert!((g.pos[1] - c.pos[1]).abs() < 1e-3, "pos.y parity: gpu {} vs cpu {}", g.pos[1], c.pos[1]);
            assert!((g.vel[0] - c.vel[0]).abs() < 1e-3, "vel.x parity");
            if (g.pos[0] - 0.5).abs() > 1e-4 {
                moved = true;
            }
        }
        assert!(moved, "the step actually advanced particles");
    }

    /// RENDER PROOF (self-skips without a compute device): additive points
    /// accumulate — a pixel under TWO stacked particles is brighter than one under a
    /// single particle, and the frame is non-blank. Renders into `Rgba8Unorm` for an
    /// easy readback; production uses the same pipeline against `Rgba16Float`.
    #[test]
    fn additive_render_accumulates_overlap() {
        let Some((device, queue)) = compute_device() else {
            eprintln!("[particles] no compute device — skipping additive render proof");
            return;
        };
        let (w, h) = (64u32, 64u32);
        // Left: a single particle. Right: two particles at the same spot (overlap).
        let parts = vec![
            Particle { pos: [0.25, 0.5], vel: [0.0, 0.0] },
            Particle { pos: [0.75, 0.5], vel: [0.0, 0.0] },
            Particle { pos: [0.75, 0.5], vel: [0.0, 0.0] },
        ];
        let gp = GpuParticles::new(&device, TextureFormat::Rgba8Unorm);
        let bufs = gp.make_buffers(&device, &parts);

        let target = device.create_texture(&wgpu::TextureDescriptor {
            label: Some("l0-particles-target"),
            size: wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
            mip_level_count: 1,
            sample_count: 1,
            dimension: wgpu::TextureDimension::D2,
            format: TextureFormat::Rgba8Unorm,
            usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC,
            view_formats: &[],
        });
        let view = target.create_view(&Default::default());
        let draw = DrawParams { point_size: 10.0, intensity: 0.45, tint: [0.4, 0.6, 1.0], speed_warm: 0.0 };
        let mut enc = device.create_command_encoder(&Default::default());
        gp.render(&device, &queue, &mut enc, &bufs, &view, &draw, false, w, h);

        // read back
        let bpp = 4u32;
        let unpadded = w * bpp;
        let align = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
        let padded = unpadded.div_ceil(align) * align;
        let readback = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("l0-particles-px-readback"),
            size: (padded * h) as u64,
            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
            mapped_at_creation: false,
        });
        enc.copy_texture_to_buffer(
            wgpu::TexelCopyTextureInfo { texture: &target, mip_level: 0, origin: wgpu::Origin3d::ZERO, aspect: wgpu::TextureAspect::All },
            wgpu::TexelCopyBufferInfo { buffer: &readback, layout: wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(padded), rows_per_image: Some(h) } },
            wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
        );
        queue.submit(Some(enc.finish()));
        let slice = readback.slice(..);
        let (tx, rx) = std::sync::mpsc::channel();
        slice.map_async(wgpu::MapMode::Read, move |r| { let _ = tx.send(r); });
        device.poll(wgpu::PollType::wait_indefinitely()).ok();
        rx.recv().unwrap().unwrap();
        let data = slice.get_mapped_range();
        let mut rgba = Vec::with_capacity((w * h * 4) as usize);
        for row in 0..h {
            let s = (row * padded) as usize;
            rgba.extend_from_slice(&data[s..s + unpadded as usize]);
        }
        drop(data);
        readback.unmap();

        let at = |x: u32, y: u32| -> u32 {
            let i = ((y * w + x) * 4) as usize;
            rgba[i] as u32 + rgba[i + 1] as u32 + rgba[i + 2] as u32
        };
        let lit = rgba.chunks_exact(4).filter(|p| p[0] as u32 + p[1] as u32 + p[2] as u32 > 0).count();
        assert!(lit > 0, "particles drew something ({lit} lit px)");
        let single = at(w / 4, h / 2); // under one particle
        let overlap = at(3 * w / 4, h / 2); // under two stacked particles
        assert!(single > 0, "single particle lit its centre ({single})");
        assert!(overlap > single + 10, "two stacked particles add brighter ({overlap}) than one ({single})");
    }

    /// INJECT-ASSERT: the WGSL exposes the entry points the pipelines name + the
    /// uniform blocks are the size the Rust structs declare.
    #[test]
    fn shader_entry_points_and_uniform_sizes() {
        assert!(PARTICLES_WGSL.contains("fn cs_step"));
        assert!(PARTICLES_WGSL.contains("fn cs_advect"));
        assert!(PARTICLES_WGSL.contains("fn sample_polyline"));
        assert!(PARTICLES_WGSL.contains("fn pt_vs"));
        assert!(PARTICLES_WGSL.contains("fn pt_fs"));
        assert!(PARTICLES_WGSL.contains("@workgroup_size(64)"));
        assert_eq!(std::mem::size_of::<Particle>(), 16);
        assert_eq!(std::mem::size_of::<SimUniforms>(), 48);
        assert_eq!(std::mem::size_of::<AdvectUniforms>(), 16);
        assert_eq!(std::mem::size_of::<DrawUniforms>(), 32);
    }

    /// INJECT-ASSERT (pure, no device): the CPU advect advances `s`, wraps it to
    /// `[0,1)`, and resamples the poly-line — deterministic (bit-identical twice)
    /// and it actually moves the particle along the edge (Δs > 0).
    #[test]
    fn advect_cpu_advances_and_wraps() {
        // Two straight edges of unit length in x.
        let polylines = vec![
            vec![[0.0f32, 0.0], [1.0, 0.0]],
            vec![[0.0f32, 0.5], [1.0, 0.5]],
        ];
        let part_edge = vec![0u32, 1u32];
        let init = vec![
            Particle { pos: [0.0, 0.0], vel: [0.9, 0.4] }, // fast, near the end → wraps
            Particle { pos: [0.0, 0.5], vel: [0.0, 0.1] }, // slow
        ];
        let a = advect_cpu(&init, &polylines, &part_edge, 0.5);
        let b = advect_cpu(&init, &polylines, &part_edge, 0.5);
        assert_eq!(a[0].pos[0].to_bits(), b[0].pos[0].to_bits(), "deterministic");
        // Particle 0: s = fract(0.9 + 0.4*0.5) = fract(1.1) = 0.1, x follows s.
        assert!((a[0].vel[0] - 0.1).abs() < 1e-5, "s wrapped: {}", a[0].vel[0]);
        assert!((a[0].pos[0] - 0.1).abs() < 1e-4, "pos rides the edge at s: {}", a[0].pos[0]);
        // Particle 1 sits on edge 1 (y = 0.5), moved by 0.05.
        assert!((a[1].pos[1] - 0.5).abs() < 1e-4, "stayed on its edge");
        assert!(a[1].pos[0] > init[1].pos[0], "advanced along the edge");
    }

    /// GPU↔CPU ADVECT PARITY PROOF (self-skips without a compute device / enough
    /// storage buffers): one GPU `advect` step matches [`advect_cpu`] within float
    /// tolerance and the particles actually moved (Δs > 0) — proving WGSL
    /// `cs_advect` and the CPU mirror are the same function (FC-7).
    #[test]
    fn gpu_advect_matches_cpu_reference() {
        let Some((device, queue)) = compute_device() else {
            eprintln!("[particles] no compute device — skipping GPU advect parity proof");
            return;
        };
        // The advect pass binds 5 storage buffers in one stage; skip if the device
        // can't offer that many (downlevel adapters cap at 4).
        if device.limits().max_storage_buffers_per_shader_stage < 5 {
            eprintln!("[particles] <5 storage buffers/stage — skipping GPU advect parity proof");
            return;
        }
        let polylines = vec![
            vec![[0.1f32, 0.2], [0.4, 0.3], [0.8, 0.25]],
            vec![[0.2f32, 0.8], [0.6, 0.6], [0.9, 0.7]],
        ];
        // Seed particles across the two edges with assorted s + speed.
        let init: Vec<Particle> = (0..24)
            .map(|i| {
                let e = (i % 2) as usize;
                let s0 = (i as f32 / 24.0).fract();
                let speed = 0.15 + 0.05 * (i % 3) as f32;
                let p0 = polylines[e][0];
                Particle { pos: p0, vel: [s0, speed] }
            })
            .collect();
        let part_edge: Vec<u32> = (0..24).map(|i| (i % 2) as u32).collect();
        let dt = 0.1f32;

        let gp = GpuParticles::new(&device, TextureFormat::Rgba16Float);
        let mut bufs = gp.make_buffers(&device, &init);
        let geo = FlowGeometry::new(&device, &polylines, &part_edge);
        let mut enc = device.create_command_encoder(&Default::default());
        gp.advect(&device, &queue, &mut enc, &mut bufs, &geo, dt);
        queue.submit(Some(enc.finish()));
        device.poll(wgpu::PollType::wait_indefinitely()).ok();

        let gpu = read_particles(&device, &queue, &bufs.a, bufs.count);
        let cpu = advect_cpu(&init, &polylines, &part_edge, dt);
        assert_eq!(gpu.len(), cpu.len());
        let mut moved = false;
        for (idx, (g, c)) in gpu.iter().zip(&cpu).enumerate() {
            assert!((g.pos[0] - c.pos[0]).abs() < 1e-3, "pos.x parity[{idx}]: gpu {} vs cpu {}", g.pos[0], c.pos[0]);
            assert!((g.pos[1] - c.pos[1]).abs() < 1e-3, "pos.y parity[{idx}]: gpu {} vs cpu {}", g.pos[1], c.pos[1]);
            assert!((g.vel[0] - c.vel[0]).abs() < 1e-3, "s parity[{idx}]");
            if (g.vel[0] - init[idx].vel[0]).abs() > 1e-4 {
                moved = true;
            }
        }
        assert!(moved, "advect actually advanced the arc param");
    }
}