facett-graphview/src/gpu_layout.rs
//! **GPU-compute force-directed layout** — the elite engine's must-land core.
//!
//! A wgpu **compute-shader** n-body layout: per-node weights + CSR adjacency are
//! streamed from an Arrow edge batch straight into GPU **storage buffers**, and the
//! node positions then **iterate in VRAM** (ping-pong `src`→`dst`, no readback per
//! step) until they settle. The host reads the converged positions back as plain
//! `[f32; 2]` **DATA** — exactly the contract a headless test asserts on.
//!
//! The all-pairs repulsion is the same O(n²) compute pattern as the boids in
//! [`facett_core::render::gpu::particles`] (reused verbatim — storage ping-pong, a
//! `@workgroup_size(64)` kernel, the headless `request_adapter` device probe).
//!
//! # Measured, so the scale claim is checkable
//! `cargo run --release -p facett-graphview --features gpu --example scale_probe`
//! on an RTX 4090, 100 000 nodes / 199 367 edges, warm device:
//!
//! | lane | 1k | 10k | 100k |
//! |------------------------------|--------|---------|----------|
//! | [`step_cpu`], **one** step | 2.6 ms | 248 ms | 23.7 s |
//! | [`gpu::force_layout_gpu_on`], 50 steps | 49 ms | 89 ms | 983 ms |
//!
//! So the GPU lane is the only one that lays 100k out interactively; [`step_cpu`] at
//! that size is ~40 minutes for a 100-step converge and is a correctness reference,
//! not a fallback anyone should reach at scale. The O(n log n) Barnes–Hut + fork-join
//! CPU lane that IS usable at 100k lives in `facett-graphnav::layout` on top of
//! [`facett_core::barnes_hut`] — not in this crate.
//!
//! # Force model (Fruchterman-Reingold, velocity-integrated)
//! For node `i` with mass `wᵢ` (its streamed weight):
//! * **repulsion** from every other node `j`: a Fruchterman-Reingold `k/d` push,
//! `rep · wᵢ · wⱼ · (pᵢ − pⱼ) / dist²` — the `(pᵢ−pⱼ)` numerator carries one
//! `dist`, so the *magnitude* is `rep·wᵢ·wⱼ / dist` (heavier nodes shove harder
//! → important nodes claim space),
//! * **attraction** along each CSR neighbour `t`: a Hooke spring
//! `attr · (dist − ideal) · (pₜ − pᵢ)/dist` (pulls connected nodes to `ideal_len`),
//! * **gravity** `−grav · pᵢ` keeps the component centred,
//! * integrate `v ← (v + F/wᵢ·dt)·damp`, `p ← p + v·dt` (heavy nodes move slower).
//!
//! It is a **pure, deterministic** function of the seed + params + step count (FC-7):
//! no RNG (the seed is a golden-angle spiral), fixed loop order. The CPU
//! [`step_cpu`] is the **no-GPU fallback** *and* the parity reference the GPU
//! readback proof checks against — the WGSL `cs_layout` mirrors it line-for-line.
//!
//! # Roadmap (see `.nornir/elite-graph-engine.md`)
//! The O(n²) repulsion is the foundation; a GPU **Barnes-Hut** quadtree (or a Morton
//! / grid binning) drops it to O(n log n) for the 100k case — documented there, not
//! built here, because the foundation must land + test clean first.
use std::collections::BTreeSet;
/// CSR (compressed-sparse-row) adjacency + per-node weight — the **GPU-streamable**
/// graph. `offsets[i]..offsets[i+1]` indexes into `targets` for node `i`'s
/// neighbours; `weights[i]` is node `i`'s mass. The three `Vec`s map 1:1 onto the
/// three GPU storage buffers (`offsets`/`targets` read-only, `weights` packed into
/// the node buffer), so this *is* the upload format — no repack at the seam.
#[derive(Clone, Debug, PartialEq)]
pub struct GraphCsr {
pub n: usize,
/// Length `n + 1`. `offsets[i]..offsets[i+1]` = node `i`'s slice of `targets`.
pub offsets: Vec<u32>,
/// Neighbour node indices (undirected: each input edge contributes both ways).
pub targets: Vec<u32>,
/// Per-node mass / importance (default `1.0`). Streamed from an Arrow weight column.
pub weights: Vec<f32>,
}
impl GraphCsr {
/// Build CSR from a **directed** edge list (`from`,`to` node indices) over `n`
/// nodes. Layout forces are undirected, so each edge attracts **both** endpoints
/// (the `targets` row holds both directions); self-loops + out-of-range ids are
/// dropped, and parallel edges are de-duplicated. `weights` (per node) default to
/// `1.0` when `None` or short.
///
/// The `edges` are exactly the `(src, dst)` pairs an Arrow edge `RecordBatch`
/// exposes (`Int64` columns cast to `u32`); [`GraphCsr::from_arrow`] is the
/// zero-glue Arrow front door over this.
#[must_use]
pub fn from_edges(n: usize, edges: &[(u32, u32)], weights: Option<&[f32]>) -> Self {
let mut adj: Vec<BTreeSet<u32>> = vec![BTreeSet::new(); n];
for &(a, b) in edges {
let (a, b) = (a as usize, b as usize);
if a == b || a >= n || b >= n {
continue;
}
adj[a].insert(b as u32);
adj[b].insert(a as u32);
}
let mut offsets = Vec::with_capacity(n + 1);
let mut targets = Vec::new();
offsets.push(0u32);
for nbrs in &adj {
targets.extend(nbrs.iter().copied());
offsets.push(targets.len() as u32);
}
let weights = (0..n)
.map(|i| weights.and_then(|w| w.get(i)).copied().filter(|w| *w > 0.0).unwrap_or(1.0))
.collect();
Self { n, offsets, targets, weights }
}
/// The number of undirected adjacency entries (`= 2 × distinct edges`).
#[must_use]
pub fn degree_sum(&self) -> usize {
self.targets.len()
}
/// Node `i`'s neighbour slice.
#[must_use]
pub fn neighbours(&self, i: usize) -> &[u32] {
let (s, e) = (self.offsets[i] as usize, self.offsets[i + 1] as usize);
&self.targets[s..e]
}
}
#[cfg(feature = "arrow")]
impl GraphCsr {
/// **Stream a CSR straight from an Arrow edge batch** — `src`/`dst` are `Int64`
/// node-id columns, the optional `weight` a `Float64` per-node mass (indexed by
/// node id). Rows with a null endpoint, a self-loop, or an out-of-range id are
/// skipped. This is the "adjacency straight from Arrow into a GPU storage buffer"
/// path: the returned `Vec`s upload unchanged.
pub fn from_arrow(
n: usize,
src: &arrow_array::Int64Array,
dst: &arrow_array::Int64Array,
weight: Option<&arrow_array::Float64Array>,
) -> Self {
use arrow_array::Array;
let mut edges = Vec::with_capacity(src.len());
for i in 0..src.len().min(dst.len()) {
if src.is_null(i) || dst.is_null(i) {
continue;
}
let (a, b) = (src.value(i), dst.value(i));
if a < 0 || b < 0 {
continue;
}
edges.push((a as u32, b as u32));
}
let w: Option<Vec<f32>> = weight.map(|w| {
(0..n).map(|i| if i < w.len() && !w.is_null(i) { w.value(i) as f32 } else { 1.0 }).collect()
});
Self::from_edges(n, &edges, w.as_deref())
}
}
/// Force-layout tuning. [`Default`] is a calm, convergent spring system.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct LayoutParams {
/// Inverse-square repulsion strength (node spreading).
pub repulsion: f32,
/// Spring (edge attraction) stiffness.
pub attraction: f32,
/// Spring rest length — the distance a single edge settles to.
pub ideal_len: f32,
/// Pull toward the origin (keeps a component from drifting / exploding).
pub gravity: f32,
/// Per-step velocity retention `∈ (0,1)` — `< 1` bleeds energy so it settles.
pub damping: f32,
/// Integration timestep.
pub dt: f32,
/// Distance floor (clamps the repulsion singularity when two nodes coincide).
pub min_dist: f32,
}
impl Default for LayoutParams {
fn default() -> Self {
Self {
repulsion: 1.0,
attraction: 2.0,
ideal_len: 1.0,
gravity: 0.05,
damping: 0.9,
dt: 0.05,
min_dist: 0.05,
}
}
}
/// The mutable layout state: one position + velocity per node (the ping-pong the GPU
/// iterates in VRAM; here, the CPU mirror). Positions are the **DATA** a test reads.
#[derive(Clone, Debug, PartialEq)]
pub struct LayoutState {
pub pos: Vec<[f32; 2]>,
pub vel: Vec<[f32; 2]>,
}
impl LayoutState {
/// The **deterministic seed**: a golden-angle (sunflower) spiral of radius `scale`,
/// zero velocity. No RNG ⇒ the whole sim is reproducible (FC-7).
#[must_use]
pub fn seed(n: usize, scale: f32) -> Self {
const GOLDEN: f32 = 2.399_963_2; // π·(3−√5)
let pos = (0..n)
.map(|i| {
let r = scale * ((i as f32 + 0.5) / n.max(1) as f32).sqrt();
let a = i as f32 * GOLDEN;
[r * a.cos(), r * a.sin()]
})
.collect();
Self { pos, vel: vec![[0.0, 0.0]; n] }
}
/// Total kinetic energy `Σ |vᵢ|²` — the **convergence metric**: a settling layout's
/// energy decays toward zero (read it as DATA to prove it converged).
#[must_use]
pub fn kinetic_energy(&self) -> f32 {
self.vel.iter().map(|v| v[0] * v[0] + v[1] * v[1]).sum()
}
/// Total edge length over `g` — the **layout-quality metric** the FDEB / structure
/// tests read (connected nodes pulled together shrink this).
#[must_use]
pub fn total_edge_len(&self, g: &GraphCsr) -> f32 {
let mut sum = 0.0;
for i in 0..g.n {
for &t in g.neighbours(i) {
if (t as usize) > i {
let d = [self.pos[t as usize][0] - self.pos[i][0], self.pos[t as usize][1] - self.pos[i][1]];
sum += (d[0] * d[0] + d[1] * d[1]).sqrt();
}
}
}
sum
}
}
/// One **CPU** force step — the no-GPU fallback *and* the parity reference the WGSL
/// `cs_layout` mirrors exactly. Pure: `step_cpu(s)` depends only on `s`, `g`, `p`.
#[must_use]
pub fn step_cpu(s: &LayoutState, g: &GraphCsr, p: &LayoutParams) -> LayoutState {
let n = g.n;
let mut out = LayoutState { pos: Vec::with_capacity(n), vel: Vec::with_capacity(n) };
let min2 = p.min_dist * p.min_dist;
for i in 0..n {
let pi = s.pos[i];
let wi = g.weights[i];
let mut force = [0.0f32, 0.0f32];
// ── all-pairs inverse-square repulsion (the O(n²) GPU mirror) ──
for (j, pj) in s.pos.iter().enumerate() {
if j == i {
continue;
}
let d = [pi[0] - pj[0], pi[1] - pj[1]];
let dist2 = (d[0] * d[0] + d[1] * d[1]).max(min2);
let f = p.repulsion * wi * g.weights[j] / dist2;
force[0] += d[0] * f;
force[1] += d[1] * f;
}
// ── spring attraction along CSR neighbours ──
for &t in g.neighbours(i) {
let pt = s.pos[t as usize];
let d = [pt[0] - pi[0], pt[1] - pi[1]];
let dist = (d[0] * d[0] + d[1] * d[1]).sqrt().max(1e-4);
let f = p.attraction * (dist - p.ideal_len) / dist;
force[0] += d[0] * f;
force[1] += d[1] * f;
}
// ── centring gravity ──
force[0] -= pi[0] * p.gravity;
force[1] -= pi[1] * p.gravity;
// ── integrate (mass = wi: heavy nodes move slower) ──
let inv_m = 1.0 / wi.max(1e-4);
let vx = (s.vel[i][0] + force[0] * inv_m * p.dt) * p.damping;
let vy = (s.vel[i][1] + force[1] * inv_m * p.dt) * p.damping;
out.vel.push([vx, vy]);
out.pos.push([pi[0] + vx * p.dt, pi[1] + vy * p.dt]);
}
out
}
/// Run the CPU fallback to convergence: seed, then `iters` [`step_cpu`]s. Returns the
/// settled state (positions = DATA). Deterministic in `(g, p, iters)`.
#[must_use]
pub fn relax_cpu(g: &GraphCsr, p: &LayoutParams, iters: usize) -> LayoutState {
let mut s = LayoutState::seed(g.n, p.ideal_len * (g.n as f32).sqrt());
for _ in 0..iters {
s = step_cpu(&s, g, p);
}
s
}
// ─────────────────────────── GPU compute path (feature `gpu`) ───────────────────────────
#[cfg(feature = "gpu")]
pub(crate) mod gpu {
use super::{GraphCsr, LayoutParams, LayoutState};
// The box-fit that rescales the settled layout into the `[-0.9, 0.9]²` world box
// lives ONCE in core (`*1.8` centre-and-divide-by-larger-span) — shared with the
// graphnav/graphview CPU FR lanes so every layout backend occupies the same world.
use facett_core::barnes_hut::normalize_box;
use wgpu::util::DeviceExt;
/// Packed node for the storage buffer: `a = (pos.x, pos.y, vel.x, vel.y)`,
/// `b = (weight, 0, 0, 0)`. 32 bytes, 16-byte aligned — matches `LNode` in
/// `layout.wgsl`. Two storage buffers of this ping-pong (`src`→`dst`).
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
pub struct GpuNode {
pub a: [f32; 4],
pub b: [f32; 4],
}
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct Uniforms {
a: [f32; 4], // n, repulsion, attraction, ideal_len
b: [f32; 4], // gravity, damping, dt, min_dist
}
/// The layout compute shader source (`cs_layout`).
pub const LAYOUT_WGSL: &str = include_str!("layout.wgsl");
/// Pack a [`LayoutState`] + per-node weights into the GPU node array.
pub fn pack(s: &LayoutState, weights: &[f32]) -> Vec<GpuNode> {
(0..s.pos.len())
.map(|i| GpuNode { a: [s.pos[i][0], s.pos[i][1], s.vel[i][0], s.vel[i][1]], b: [weights[i], 0.0, 0.0, 0.0] })
.collect()
}
/// The compiled layout pipeline + the streamed CSR buffers. Positions iterate in
/// VRAM via the [`GpuLayout::step`] ping-pong; [`GpuLayout::read_positions`] copies
/// the settled positions back as `[f32; 2]` DATA.
pub struct GpuLayout {
pipeline: wgpu::ComputePipeline,
bgl: wgpu::BindGroupLayout,
uniforms: wgpu::Buffer,
// streamed-once, read-only adjacency
offsets: wgpu::Buffer,
targets: wgpu::Buffer,
// ping-pong node buffers
src: wgpu::Buffer,
dst: wgpu::Buffer,
count: u32,
/// The grid-pyramid repulsion lane, when this layout runs it (see [`GridLane`]).
grid: Option<GridLane>,
}
impl GpuLayout {
/// Compile `cs_layout` and **stream** the CSR (`offsets`/`targets`) + the seeded
/// nodes into VRAM. Requires a device with compute + ≥4 storage buffers/stage.
pub fn new(device: &wgpu::Device, g: &GraphCsr, seed: &LayoutState) -> Self {
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("graphview_layout"),
source: wgpu::ShaderSource::Wgsl(LAYOUT_WGSL.into()),
});
let storage = |ro: bool| wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: ro }, has_dynamic_offset: false, min_binding_size: None },
count: None,
};
let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("graphview_layout_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, ..storage(true) },
wgpu::BindGroupLayoutEntry { binding: 2, ..storage(false) },
wgpu::BindGroupLayoutEntry { binding: 3, ..storage(true) },
wgpu::BindGroupLayoutEntry { binding: 4, ..storage(true) },
],
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("graphview_layout_step"),
layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("graphview_layout_pll"),
bind_group_layouts: &[Some(&bgl)],
immediate_size: 0,
})),
module: &shader,
entry_point: Some("cs_layout"),
compilation_options: Default::default(),
cache: None,
});
let nodes = pack(seed, &g.weights);
let node_usage = wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST;
let src = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("graphview_layout_src"),
contents: bytemuck::cast_slice(&nodes),
usage: node_usage,
});
let dst = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("graphview_layout_dst"),
size: (std::mem::size_of::<GpuNode>() * nodes.len().max(1)) as u64,
usage: node_usage,
mapped_at_creation: false,
});
// CSR streamed read-only (pad to ≥1 elem so empty graphs still bind).
let off_data: &[u32] = if g.offsets.is_empty() { &[0] } else { &g.offsets };
let tgt_data: Vec<u32> = if g.targets.is_empty() { vec![0] } else { g.targets.clone() };
let offsets = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("graphview_layout_offsets"),
contents: bytemuck::cast_slice(off_data),
usage: wgpu::BufferUsages::STORAGE,
});
let targets = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("graphview_layout_targets"),
contents: bytemuck::cast_slice(&tgt_data),
usage: wgpu::BufferUsages::STORAGE,
});
let uniforms = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("graphview_layout_uniforms"),
size: std::mem::size_of::<Uniforms>() as u64,
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
Self { pipeline, bgl, uniforms, offsets, targets, src, dst, count: nodes.len() as u32, grid: None }
}
/// Iterate `iters` layout steps **entirely in VRAM** (ping-pong, no per-step
/// readback). After this the settled nodes are in `self.src`.
pub fn iterate(&mut self, device: &wgpu::Device, queue: &wgpu::Queue, g: &GraphCsr, p: &LayoutParams, iters: usize) {
self.submit_steps(device, queue, g, p, iters);
device.poll(wgpu::PollType::wait_indefinitely()).ok();
}
/// [`iterate`](Self::iterate) WITHOUT the wait: encode + submit `iters` steps and
/// return. This is the half a browser can call — `wait_indefinitely` on the
/// main thread is the thing WebGPU does not have, and the [`LayoutJob`] that
/// drives a layout across frames is built on this. Same encoder, same loop
/// order, same ping-pong, so the blocking and the non-blocking roads run the
/// identical sequence of dispatches (held by `layout_job_matches_the_blocking_lane`).
pub fn submit_steps(&mut self, device: &wgpu::Device, queue: &wgpu::Queue, g: &GraphCsr, p: &LayoutParams, iters: usize) {
if let Some(grid) = self.grid.as_mut() {
let mut enc = device.create_command_encoder(&Default::default());
grid.write_uniforms(queue, self.count, p);
for _ in 0..iters {
grid.encode_step(&mut enc, self.count);
std::mem::swap(&mut self.src, &mut self.dst);
}
queue.submit(Some(enc.finish()));
return;
}
let u = Uniforms {
a: [self.count as f32, p.repulsion, p.attraction, p.ideal_len],
b: [p.gravity, p.damping, p.dt, p.min_dist],
};
queue.write_buffer(&self.uniforms, 0, bytemuck::bytes_of(&u));
let _ = g; // (offsets/targets already streamed in `new`)
let mut enc = device.create_command_encoder(&Default::default());
for _ in 0..iters {
let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("graphview_layout_bind"),
layout: &self.bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: self.uniforms.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, resource: self.src.as_entire_binding() },
wgpu::BindGroupEntry { binding: 2, resource: self.dst.as_entire_binding() },
wgpu::BindGroupEntry { binding: 3, resource: self.offsets.as_entire_binding() },
wgpu::BindGroupEntry { binding: 4, resource: self.targets.as_entire_binding() },
],
});
{
let mut cp = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("graphview_layout_pass"), timestamp_writes: None });
cp.set_pipeline(&self.pipeline);
cp.set_bind_group(0, &bind, &[]);
cp.dispatch_workgroups(self.count.div_ceil(64), 1, 1);
}
std::mem::swap(&mut self.src, &mut self.dst);
}
queue.submit(Some(enc.finish()));
}
/// **The layout a caller should run for `g`**: the grid-pyramid repulsion kernel
/// ([`GridLane`], O(n log n) per step) at or above [`GRID_MIN_NODES`] on a device
/// with the seven storage buffers it binds, else the all-pairs kernel of
/// [`new`](Self::new) (O(n²) per step, and the exact `step_cpu` mirror).
///
/// MEASURED why (loki, 2026-09-28): the all-pairs step at 100 000 nodes is ~18 ms
/// on an RTX 5080 and the integrated Radeon of the same box is ~50× slower, so an
/// 80-step converge there is minutes. The grid step touches ~230 cells per node
/// instead of 100 000 nodes.
pub fn new_for(device: &wgpu::Device, g: &GraphCsr, seed: &LayoutState) -> Self {
if g.n >= GRID_MIN_NODES && let Some(gl) = Self::new_grid(device, g, seed) {
return gl;
}
Self::new(device, g, seed)
}
/// The grid-pyramid kernel at ANY node count — `None` when the device binds
/// fewer than the seven storage buffers it needs. [`new_for`](Self::new_for) is
/// what a caller runs; this exists so the grid lane can be held against the
/// all-pairs lane and the CPU lane on a graph small enough for all three
/// (`grid_all_pairs_and_cpu_agree_on_a_small_graph`) and benched below
/// [`GRID_MIN_NODES`] (the `layout_matrix` example).
pub fn new_grid(device: &wgpu::Device, g: &GraphCsr, seed: &LayoutState) -> Option<Self> {
if device.limits().max_storage_buffers_per_shader_stage < 7 {
return None;
}
let mut gl = Self::new(device, g, seed);
gl.grid = Some(GridLane::new(device, &gl));
Some(gl)
}
/// Which repulsion sum this layout runs: `"grid"` or `"all-pairs"`.
#[must_use]
pub fn kernel(&self) -> &'static str {
if self.grid.is_some() { "grid" } else { "all-pairs" }
}
/// Copy the settled positions out of VRAM as `[f32; 2]` **DATA**.
pub fn read_positions(&self, device: &wgpu::Device, queue: &wgpu::Queue) -> Vec<[f32; 2]> {
let size = (std::mem::size_of::<GpuNode>() as u32 * self.count) as u64;
let readback = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("graphview_layout_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(&self.src, 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().expect("wgpu: get_mapped_range on a mapped buffer");
let nodes: Vec<GpuNode> = bytemuck::cast_slice(&data).to_vec();
drop(data);
readback.unmap();
nodes.iter().map(|nd| [nd.a[0], nd.a[1]]).collect()
}
}
/// Node count at and above which [`GpuLayout::new_for`] runs the grid kernel.
/// Below it the all-pairs step is cheap on any GPU and exact.
pub const GRID_MIN_NODES: usize = 16_384;
/// The grid-pyramid kernel source (`layout_grid.wgsl`).
pub const LAYOUT_GRID_WGSL: &str = include_str!("layout_grid.wgsl");
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct GridUniforms {
a: [f32; 4], // n, repulsion, attraction, ideal_len
b: [f32; 4], // gravity, damping, dt, min_dist
c: [f32; 4], // G, K, 0, 0
}
/// Stride of one level entry in the dynamic-offset uniform (WebGPU's
/// `minUniformBufferOffsetAlignment` default).
const LEVEL_STRIDE: u64 = 256;
/// **The grid-pyramid repulsion lane** of a [`GpuLayout`] — the passes of
/// `layout_grid.wgsl`, encoded per step: clear, bounds, bin, finest cells, K
/// reductions, force. Every buffer is allocated once, at construction; a step
/// allocates nothing (both ping-pong bind groups are prebuilt).
pub(crate) struct GridLane {
bounds_p: wgpu::ComputePipeline,
bin_p: wgpu::ComputePipeline,
fine_p: wgpu::ComputePipeline,
reduce_p: wgpu::ComputePipeline,
force_p: wgpu::ComputePipeline,
uniforms: wgpu::Buffer,
bounds: wgpu::Buffer,
fine: wgpu::Buffer,
_cells: wgpu::Buffer,
_levels: wgpu::Buffer,
levels_bg: wgpu::BindGroup,
/// `[src→dst, dst→src]`, selected by `flip`.
bg: [wgpu::BindGroup; 2],
flip: bool,
side: u32,
k: u32,
}
impl GridLane {
/// The finest grid side for `n` nodes: a power of two with about two nodes per
/// cell, clamped to `8 ..= 512`.
fn side_for(n: usize) -> u32 {
let want = ((n as f64 / 2.0).sqrt().ceil() as u32).max(8);
want.next_power_of_two().min(512)
}
fn new(device: &wgpu::Device, gl: &GpuLayout) -> Self {
let side = Self::side_for(gl.count as usize);
let k = side.trailing_zeros();
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("graphview_layout_grid"),
source: wgpu::ShaderSource::Wgsl(LAYOUT_GRID_WGSL.into()),
});
let entry = |binding: u32, ty: wgpu::BindingType| wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
ty,
count: None,
};
let storage = |ro: bool| wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: ro },
has_dynamic_offset: false,
min_binding_size: None,
};
let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("graphview_layout_grid_bgl"),
entries: &[
entry(
0,
wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: false,
min_binding_size: None,
},
),
entry(1, storage(true)),
entry(2, storage(false)),
entry(3, storage(true)),
entry(4, storage(true)),
entry(5, storage(false)),
entry(6, storage(false)),
entry(7, storage(false)),
],
});
let lvl_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("graphview_layout_grid_lvl_bgl"),
entries: &[entry(
0,
wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: true,
min_binding_size: wgpu::BufferSize::new(16),
},
)],
});
let pl0 = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("graphview_layout_grid_pl0"),
bind_group_layouts: &[Some(&bgl)],
immediate_size: 0,
});
let pl1 = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("graphview_layout_grid_pl1"),
bind_group_layouts: &[Some(&bgl), Some(&lvl_bgl)],
immediate_size: 0,
});
let pipe = |layout: &wgpu::PipelineLayout, name: &str| {
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some(name),
layout: Some(layout),
module: &shader,
entry_point: Some(name),
compilation_options: Default::default(),
cache: None,
})
};
let bounds_p = pipe(&pl0, "cs_bounds");
let bin_p = pipe(&pl0, "cs_bin");
let fine_p = pipe(&pl0, "cs_fine");
let reduce_p = pipe(&pl1, "cs_reduce");
let force_p = pipe(&pl0, "cs_force");
let buf = |label: &str, size: u64, usage: wgpu::BufferUsages| {
device.create_buffer(&wgpu::BufferDescriptor { label: Some(label), size: size.max(16), usage, mapped_at_creation: false })
};
let st = wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST;
let uniforms = buf(
"graphview_layout_grid_u",
std::mem::size_of::<GridUniforms>() as u64,
wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
);
let bounds = buf("graphview_layout_grid_bounds", 16, st);
let fine = buf("graphview_layout_grid_fine", u64::from(side) * u64::from(side) * 3 * 4, st);
let cells_len = ((1u64 << (2 * (k + 1))) - 1) / 3;
let cells = buf("graphview_layout_grid_cells", cells_len * 16, st);
let mut lv = vec![0u8; (LEVEL_STRIDE * u64::from(k.max(1))) as usize];
for l in 0..k {
let at = (LEVEL_STRIDE * u64::from(l)) as usize;
lv[at..at + 4].copy_from_slice(&l.to_le_bytes());
}
let levels = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("graphview_layout_grid_levels"),
contents: &lv,
usage: wgpu::BufferUsages::UNIFORM,
});
let levels_bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("graphview_layout_grid_levels_bg"),
layout: &lvl_bgl,
entries: &[wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
buffer: &levels,
offset: 0,
size: wgpu::BufferSize::new(16),
}),
}],
});
let bind = |src: &wgpu::Buffer, dst: &wgpu::Buffer| {
device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("graphview_layout_grid_bg"),
layout: &bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: uniforms.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, resource: src.as_entire_binding() },
wgpu::BindGroupEntry { binding: 2, resource: dst.as_entire_binding() },
wgpu::BindGroupEntry { binding: 3, resource: gl.offsets.as_entire_binding() },
wgpu::BindGroupEntry { binding: 4, resource: gl.targets.as_entire_binding() },
wgpu::BindGroupEntry { binding: 5, resource: bounds.as_entire_binding() },
wgpu::BindGroupEntry { binding: 6, resource: fine.as_entire_binding() },
wgpu::BindGroupEntry { binding: 7, resource: cells.as_entire_binding() },
],
})
};
let bg = [bind(&gl.src, &gl.dst), bind(&gl.dst, &gl.src)];
Self {
bounds_p,
bin_p,
fine_p,
reduce_p,
force_p,
uniforms,
bounds,
fine,
_cells: cells,
_levels: levels,
levels_bg,
bg,
flip: false,
side,
k,
}
}
fn write_uniforms(&self, queue: &wgpu::Queue, n: u32, p: &LayoutParams) {
let u = GridUniforms {
a: [n as f32, p.repulsion, p.attraction, p.ideal_len],
b: [p.gravity, p.damping, p.dt, p.min_dist],
c: [self.side as f32, self.k as f32, 0.0, 0.0],
};
queue.write_buffer(&self.uniforms, 0, bytemuck::bytes_of(&u));
}
/// Encode ONE layout step. The caller swaps its `src`/`dst` after, exactly as the
/// all-pairs lane does; `flip` tracks which prebuilt bind group matches.
fn encode_step(&mut self, enc: &mut wgpu::CommandEncoder, n: u32) {
enc.clear_buffer(&self.bounds, 0, None);
enc.clear_buffer(&self.fine, 0, None);
let bg = &self.bg[usize::from(self.flip)];
{
let mut cp = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("graphview_layout_grid"), timestamp_writes: None });
cp.set_bind_group(0, bg, &[]);
cp.set_pipeline(&self.bounds_p);
cp.dispatch_workgroups(n.div_ceil(256), 1, 1);
cp.set_pipeline(&self.bin_p);
cp.dispatch_workgroups(n.div_ceil(256), 1, 1);
cp.set_pipeline(&self.fine_p);
cp.dispatch_workgroups((self.side * self.side).div_ceil(256), 1, 1);
cp.set_pipeline(&self.reduce_p);
for l in (0..self.k).rev() {
cp.set_bind_group(1, &self.levels_bg, &[(LEVEL_STRIDE * u64::from(l)) as u32]);
cp.dispatch_workgroups((1u32 << (2 * l)).div_ceil(256), 1, 1);
}
cp.set_pipeline(&self.force_p);
cp.dispatch_workgroups(n.div_ceil(64), 1, 1);
}
self.flip = !self.flip;
}
}
/// Acquire a headless compute `(Device, Queue)` — `None` when no usable wgpu adapter
/// exists (headless CI / no GPU), so the caller falls back to the CPU lane. Requires
/// compute shaders + ≥4 storage buffers per stage (the layout pipeline's bind group).
pub 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::HighPerformance,
force_fallback_adapter: false,
compatible_surface: None,
apply_limit_buckets: false,
}))
.ok()?;
if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
return None;
}
if adapter.limits().max_storage_buffers_per_shader_stage < 4 {
return None;
}
pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
label: Some("facett-graphview-gpu-layout"),
required_features: wgpu::Features::empty(),
required_limits: adapter.limits(),
memory_hints: wgpu::MemoryHints::Performance,
experimental_features: wgpu::ExperimentalFeatures::disabled(),
trace: wgpu::Trace::Off,
}))
.ok()
}
/// **Runtime GPU probe** for the layout lane: `true` iff a wgpu compute device with
/// the layout pipeline's requirements can actually be created right now. The
/// two-lane selector calls this once and routes to the GPU lane when it holds.
#[must_use]
pub fn gpu_layout_available() -> bool {
compute_device().is_some()
}
/// **GPU lane front door** — lay out `n` nodes with directed `edges` on the GPU and
/// return normalised `[-1, 1]²` positions, or `None` if no usable adapter exists.
/// On `None` the caller picks its own fallback; this crate's only CPU layout is the
/// O(n²) [`step_cpu`]/[`relax_cpu`] reference, which is not usable above ~30k. The
/// Barnes–Hut fork-join lane is `facett-graphnav::layout::force_directed`.
/// Streams a CSR into VRAM, iterates
/// the `cs_layout` compute kernel `iters` times **entirely in VRAM** (ping-pong, no
/// per-step readback), reads the settled positions back as DATA, then fits them to
/// the box. Deterministic in `(n, edges, iters)` — the `cs_layout` mirrors
/// [`step_cpu`] line-for-line (FC-7 parity, proven by `gpu_layout_matches_cpu_reference`).
#[must_use]
pub fn force_layout_gpu(n: usize, edges: &[(usize, usize)], iters: usize) -> Option<Vec<[f32; 2]>> {
if n == 0 {
return Some(Vec::new());
}
let (device, queue) = compute_device()?;
Some(force_layout_gpu_on(&device, &queue, n, edges, iters))
}
/// The **warm-device** GPU layout — lay out on an ALREADY-CREATED `(device, queue)`
/// and return normalised `[-1, 1]²` positions. This is the realistic per-converge
/// cost in a wgpu host (egui-wgpu / vello already holds a live device, so the ~one-
/// time adapter/device init is NOT paid per layout): stream the CSR, iterate
/// `cs_layout` in VRAM `iters` times, read back, fit to the box. [`force_layout_gpu`]
/// is the one-shot front door that creates a device then calls this.
#[must_use]
pub fn force_layout_gpu_on(
device: &wgpu::Device,
queue: &wgpu::Queue,
n: usize,
edges: &[(usize, usize)],
iters: usize,
) -> Vec<[f32; 2]> {
if n == 0 {
return Vec::new();
}
let e32: Vec<(u32, u32)> = edges
.iter()
.filter(|&&(a, b)| a < n && b < n)
.map(|&(a, b)| (a as u32, b as u32))
.collect();
let g = GraphCsr::from_edges(n, &e32, None);
let p = LayoutParams::default();
let seed = LayoutState::seed(n, p.ideal_len * (n as f32).sqrt());
let mut gl = GpuLayout::new_for(device, &g, &seed);
gl.iterate(device, queue, &g, &p, iters);
let mut pos = gl.read_positions(device, queue);
normalize_box(&mut pos);
pos
}
// ── THE HOST'S DEVICE, BORROWED ──────────────────────────────────────────
//
// `force_layout_gpu` makes its own instance, adapter and device with
// `pollster::block_on`. Natively that is merely wasteful (the egui-wgpu host
// already holds a device). In a browser it is impossible: blocking the main thread
// on a promise is a deadlock, not a slow path — which is why, until 2026-09-18, the
// wasm wheel shipped this whole compute shader and could never dispatch it.
//
// The seam is the same shape every other GPU lane here uses: the HOST installs
// once at startup, from the `RenderState` it already has, and the lane finds what
// was installed. `thread_local`, not a `static`: the UI thread is the only thread
// that installs or lays out, wasm has exactly one thread, and a thread-local asks
// nothing of `Device: Sync` on a target where that is a feature flag.
thread_local! {
static HOST_DEVICE: std::cell::RefCell<Option<(wgpu::Device, wgpu::Queue)>> =
const { std::cell::RefCell::new(None) };
}
/// **Lend the layout lane the host's `(Device, Queue)`** — call once at startup with
/// the clones out of the host's `egui_wgpu::RenderState`. Returns `false` (and
/// installs nothing) when the device cannot run the layout pipeline, so a host on a
/// downlevel adapter keeps the CPU lane instead of failing at first dispatch.
/// Idempotent: a second install replaces the first.
pub fn install_host_device(device: &wgpu::Device, queue: &wgpu::Queue) -> bool {
// The same two requirements `compute_device` checks on an adapter, read off
// the DEVICE — the host made it, so its limits are what we actually get.
if device.limits().max_storage_buffers_per_shader_stage < 4 {
return false;
}
HOST_DEVICE.with(|h| *h.borrow_mut() = Some((device.clone(), queue.clone())));
true
}
/// Whether a host device is installed on THIS thread.
#[must_use]
pub fn host_device_installed() -> bool {
HOST_DEVICE.with(|h| h.borrow().is_some())
}
/// Run `f` with the installed host device, if any.
pub fn with_host_device<R>(f: impl FnOnce(&wgpu::Device, &wgpu::Queue) -> R) -> Option<R> {
HOST_DEVICE.with(|h| h.borrow().as_ref().map(|(d, q)| f(d, q)))
}
/// TEST/teardown hook: forget the installed device.
pub fn clear_host_device() {
HOST_DEVICE.with(|h| *h.borrow_mut() = None);
}
/// How many `cs_layout` steps one [`LayoutJob::poll`] submits. Small enough that a
/// 100k-node all-pairs chunk stays well under a browser's GPU watchdog, large
/// enough that a typical job finishes in a handful of frames. The total is what
/// decides the layout, not the chunking — the dispatches are the same sequence.
pub const JOB_STEPS_PER_POLL: usize = 10;
/// What a [`LayoutJob::poll`] answers.
#[derive(Debug, Clone, PartialEq)]
pub enum JobState {
/// Still iterating, or the readback has not landed yet. Poll again next frame.
Pending,
/// Settled: normalised `[-1, 1]²` positions, exactly what
/// [`force_layout_gpu_on`] returns for the same `(n, edges, iters)`.
Ready(Vec<[f32; 2]>),
/// The job could not produce a layout, and WHY. The caller falls back to its CPU
/// lane and carries the reason into its own readout — a fallback that does not
/// say why is indistinguishable, from outside, from the lane never being tried.
Failed(&'static str),
}
/// **A GPU layout that never blocks** — the browser-shaped twin of
/// [`force_layout_gpu_on`].
///
/// `start` streams the CSR and returns immediately. Each `poll` does a bounded
/// amount of work: submit the next [`JOB_STEPS_PER_POLL`] steps, or — once every
/// step is submitted — kick the readback, or check whether it landed. No call
/// waits on the GPU, so it is safe on a browser's main thread, and natively it
/// simply spreads the same work over a few frames.
///
/// Deterministic in `(n, edges, iters)` like the blocking lane: same seed, same
/// params, same dispatch sequence (`layout_job_matches_the_blocking_lane`).
pub struct LayoutJob {
layout: GpuLayout,
g: GraphCsr,
p: LayoutParams,
remaining: usize,
readback: Option<wgpu::Buffer>,
/// 0 = mapping pending, 1 = mapped, 2 = failed. Set by wgpu's map callback.
mapped: std::sync::Arc<std::sync::atomic::AtomicU8>,
done: bool,
/// Fit the settled positions to the `[-0.9, 0.9]²` box ([`start`](Self::start)),
/// or hand them back in layout units ([`start_with`](Self::start_with)).
normalize: bool,
}
impl LayoutJob {
/// Stream the graph into VRAM and return — nothing is dispatched yet.
#[must_use]
pub fn start(device: &wgpu::Device, n: usize, edges: &[(usize, usize)], iters: usize) -> Self {
let e32: Vec<(u32, u32)> = edges
.iter()
.filter(|&&(a, b)| a < n && b < n)
.map(|&(a, b)| (a as u32, b as u32))
.collect();
let g = GraphCsr::from_edges(n, &e32, None);
let p = LayoutParams::default();
let seed = LayoutState::seed(n, p.ideal_len * (n as f32).sqrt());
Self::start_with(device, g, p, &seed, iters, true)
}
/// [`start`](Self::start) for a caller with its OWN params and seed — the GPU
/// twin of `seed → iters × step_cpu` ([`super::relax_cpu`] when `seed` is its
/// seed). `normalize: false` hands the settled positions back in layout units,
/// so a host whose CPU lane scaled `relax_cpu`'s output keeps the same world on
/// the GPU (facett-graphview's falkor browser). The kernel follows the size as in
/// [`GpuLayout::new_for`]: all-pairs below [`GRID_MIN_NODES`], the grid above.
#[must_use]
pub fn start_with(
device: &wgpu::Device,
g: GraphCsr,
p: LayoutParams,
seed: &LayoutState,
iters: usize,
normalize: bool,
) -> Self {
let layout = GpuLayout::new_for(device, &g, seed);
Self {
layout,
g,
p,
remaining: iters,
readback: None,
mapped: std::sync::Arc::new(std::sync::atomic::AtomicU8::new(0)),
done: false,
normalize,
}
}
/// Which repulsion kernel this job runs: `"grid"` or `"all-pairs"`.
#[must_use]
pub fn kernel(&self) -> &'static str {
self.layout.kernel()
}
/// Steps not yet submitted — a host can show progress off this.
#[must_use]
pub fn remaining_steps(&self) -> usize {
self.remaining
}
/// Advance the job by one bounded slice of work. Never waits on the GPU.
pub fn poll(&mut self, device: &wgpu::Device, queue: &wgpu::Queue) -> JobState {
use std::sync::atomic::Ordering;
if self.done {
return JobState::Failed("job already consumed");
}
if self.remaining > 0 {
let k = self.remaining.min(JOB_STEPS_PER_POLL);
self.layout.submit_steps(device, queue, &self.g, &self.p, k);
self.remaining -= k;
return JobState::Pending;
}
let Some(readback) = &self.readback else {
// Every step is submitted: copy the settled nodes out and ask for the map.
let size = (std::mem::size_of::<GpuNode>() as u32 * self.layout.count) as u64;
let rb = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("graphview_layout_job_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(&self.layout.src, 0, &rb, 0, size);
queue.submit(Some(enc.finish()));
let flag = std::sync::Arc::clone(&self.mapped);
rb.slice(..).map_async(wgpu::MapMode::Read, move |r| {
flag.store(if r.is_ok() { 1 } else { 2 }, Ordering::SeqCst);
});
self.readback = Some(rb);
return JobState::Pending;
};
// Natively this is what fires the map callback; on WebGPU the browser's
// event loop does and this is a no-op. Either way it does not wait.
device.poll(wgpu::PollType::Poll).ok();
match self.mapped.load(Ordering::SeqCst) {
0 => JobState::Pending,
1 => {
let data = readback.slice(..).get_mapped_range().expect("wgpu: get_mapped_range on a mapped buffer");
let nodes: Vec<GpuNode> = bytemuck::cast_slice(&data).to_vec();
drop(data);
readback.unmap();
let mut pos: Vec<[f32; 2]> = nodes.iter().map(|nd| [nd.a[0], nd.a[1]]).collect();
if self.normalize {
normalize_box(&mut pos);
}
self.done = true;
JobState::Ready(pos)
}
_ => {
self.done = true;
JobState::Failed("readback map was rejected (device lost, or a validation error upstream)")
}
}
}
}
}
#[cfg(feature = "gpu")]
pub use gpu::{
clear_host_device, compute_device, force_layout_gpu, force_layout_gpu_on, gpu_layout_available,
host_device_installed, install_host_device, with_host_device, GpuLayout, GpuNode, JobState, LayoutJob,
GRID_MIN_NODES, JOB_STEPS_PER_POLL, LAYOUT_GRID_WGSL, LAYOUT_WGSL,
};
#[cfg(test)]
mod tests {
use super::*;
/// A small fixture: two triangles (0-1-2, 3-4-5) bridged by a single edge 2-3.
fn two_triangles() -> GraphCsr {
let edges = [(0, 1), (1, 2), (2, 0), (3, 4), (4, 5), (5, 3), (2, 3)];
GraphCsr::from_edges(6, &edges, None)
}
#[test]
fn csr_is_symmetric_and_dedups() {
// Parallel + reversed edges collapse; adjacency is undirected.
let g = GraphCsr::from_edges(3, &[(0, 1), (1, 0), (0, 1), (1, 2)], None);
assert_eq!(g.neighbours(0), &[1]);
assert_eq!(g.neighbours(1), &[0, 2]);
assert_eq!(g.neighbours(2), &[1]);
assert_eq!(g.degree_sum(), 4, "two undirected edges → 4 directed entries");
// self-loops + out-of-range dropped
let g2 = GraphCsr::from_edges(2, &[(0, 0), (0, 5), (0, 1)], None);
assert_eq!(g2.neighbours(0), &[1]);
}
#[test]
fn seed_is_deterministic_and_bounded() {
let a = LayoutState::seed(50, 3.0);
let b = LayoutState::seed(50, 3.0);
assert_eq!(a, b, "no RNG: the seed is reproducible");
for q in &a.pos {
assert!(q[0].is_finite() && q[1].is_finite());
assert!((q[0] * q[0] + q[1] * q[1]).sqrt() <= 3.0 + 1e-3, "inside the seed disc");
}
}
#[test]
fn step_cpu_is_deterministic() {
let g = two_triangles();
let p = LayoutParams::default();
let s = LayoutState::seed(g.n, 2.0);
let a = step_cpu(&s, &g, &p);
let b = step_cpu(&s, &g, &p);
for (x, y) in a.pos.iter().zip(&b.pos) {
assert_eq!(x[0].to_bits(), y[0].to_bits(), "bit-identical pos.x");
assert_eq!(x[1].to_bits(), y[1].to_bits(), "bit-identical pos.y");
}
}
#[test]
fn layout_converges_energy_decays() {
let g = two_triangles();
let p = LayoutParams::default();
let mut s = LayoutState::seed(g.n, p.ideal_len * (g.n as f32).sqrt());
// skip the first few steps (the system spikes as springs engage), then watch
// the kinetic energy fall as it settles — the convergence proof, as DATA.
for _ in 0..15 {
s = step_cpu(&s, &g, &p);
}
let e_early = s.kinetic_energy();
for _ in 0..200 {
s = step_cpu(&s, &g, &p);
}
let e_late = s.kinetic_energy();
assert!(e_late < e_early * 0.1, "energy decayed an order of magnitude (converged): {e_early} → {e_late}");
assert!(e_late < 0.05, "settled to near-rest: {e_late}");
for q in &s.pos {
assert!(q[0].is_finite() && q[1].is_finite(), "no NaN/inf blow-up");
}
}
#[test]
fn springs_settle_an_edge_at_the_analytic_equilibrium() {
// A single edge 0—1 balances the FR repulsion (magnitude rep/d) against the
// Hooke spring (attr·(d−ideal)): rep/d = attr·(d−ideal) ⇒ for the defaults
// (rep=1, attr=2, ideal=1) ⇒ 2d²−2d−1=0 ⇒ d = (1+√3)/2 ≈ 1.366.
let g = GraphCsr::from_edges(2, &[(0, 1)], None);
let p = LayoutParams { gravity: 0.0, ..LayoutParams::default() };
let s = relax_cpu(&g, &p, 600);
let d = ((s.pos[1][0] - s.pos[0][0]).powi(2) + (s.pos[1][1] - s.pos[0][1]).powi(2)).sqrt();
let expected = (1.0 + 3.0_f32.sqrt()) / 2.0;
assert!((d - expected).abs() < 0.02, "edge rests at the force equilibrium {expected} (got {d})");
}
#[test]
fn relax_is_deterministic_positions_are_data() {
let g = two_triangles();
let p = LayoutParams::default();
let a = relax_cpu(&g, &p, 120);
let b = relax_cpu(&g, &p, 120);
assert_eq!(a.pos, b.pos, "same params → identical positions (reproducible DATA)");
// The bridged triangles separate: the cross-bridge edge 2—3 is the graph's
// longest single edge by construction (its endpoints anchor opposite clusters).
let len = |i: usize, j: usize| ((a.pos[j][0] - a.pos[i][0]).powi(2) + (a.pos[j][1] - a.pos[i][1]).powi(2)).sqrt();
let within = len(0, 1).max(len(1, 2)).max(len(3, 4)).max(len(4, 5));
assert!(len(2, 3) > within * 0.9, "the two clusters spread apart along the bridge");
}
#[test]
fn heavier_nodes_move_less_per_step() {
// Same geometry, but node 0 is 50× heavier → its first-step displacement is
// far smaller (mass in the integrator).
let edges = [(0, 1), (0, 2)];
let light = GraphCsr::from_edges(3, &edges, None);
let heavy = GraphCsr::from_edges(3, &edges, Some(&[50.0, 1.0, 1.0]));
let p = LayoutParams::default();
let s = LayoutState::seed(3, 2.0);
let dl = step_cpu(&s, &light, &p);
let dh = step_cpu(&s, &heavy, &p);
let disp = |st: &LayoutState| ((st.pos[0][0] - s.pos[0][0]).powi(2) + (st.pos[0][1] - s.pos[0][1]).powi(2)).sqrt();
assert!(disp(&dh) < disp(&dl) * 0.5, "the heavy node barely moves ({} vs {})", disp(&dh), disp(&dl));
}
// ─────────────────────────── GPU parity proof (feature `gpu`) ───────────────────────────
/// Headless compute device with the adapter's real limits; `None` ⇒ self-skip.
#[cfg(feature = "gpu")]
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,
apply_limit_buckets: false,
}))
.ok()?;
if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
return None;
}
// Name the adapter in the test output: a box with a CPU rasteriser next to its
// GPU (llvmpipe beside a Radeon) must show WHICH one these proofs ran on.
let info = adapter.get_info();
eprintln!("[gpu_layout tests] adapter: {} ({:?}, {:?})", info.name, info.device_type, info.backend);
if adapter.limits().max_storage_buffers_per_shader_stage < 4 {
return None;
}
pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
label: Some("graphview-layout-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()
}
/// GPU↔CPU PARITY PROOF (self-skips without a compute device): the same CSR seeded
/// the same way and stepped N times **in VRAM** lands on the same positions as the
/// CPU fallback (within float tolerance) — proving `cs_layout` and `step_cpu` are
/// one deterministic function, and that positions read back as real DATA.
#[cfg(feature = "gpu")]
#[test]
fn gpu_layout_matches_cpu_reference() {
// THE GPU TURN, held before the bring-up and across the whole body —
// `facett_core::render::gputurn`, the same one facett-core's probes and
// facett-map's kittest files take (LAW 5: never a second lock). This
// hand-rolled `compute_device()` took no turn, which is the shape that
// SIGSEGV'd `facett_graph3d`'s two kittest tests 17 times in 30 runs on
// 2026-09-03 and deadlocked `facett-core --lib` for 13 min on 2026-08-31.
let _gpu = facett_core::render::gputurn::hold();
let Some((device, queue)) = compute_device() else {
eprintln!("[gpu_layout] no compute device — skipping GPU parity proof");
return;
};
let g = two_triangles();
let p = LayoutParams::default();
let iters = 50;
let seed = LayoutState::seed(g.n, p.ideal_len * (g.n as f32).sqrt());
let mut gl = GpuLayout::new(&device, &g, &seed);
gl.iterate(&device, &queue, &g, &p, iters);
let gpu = gl.read_positions(&device, &queue);
let mut cpu = seed.clone();
for _ in 0..iters {
cpu = step_cpu(&cpu, &g, &p);
}
assert_eq!(gpu.len(), cpu.pos.len());
let mut moved = false;
for (gp, cp) in gpu.iter().zip(&cpu.pos) {
assert!((gp[0] - cp[0]).abs() < 1e-2, "pos.x parity: gpu {} vs cpu {}", gp[0], cp[0]);
assert!((gp[1] - cp[1]).abs() < 1e-2, "pos.y parity: gpu {} vs cpu {}", gp[1], cp[1]);
assert!(gp[0].is_finite() && gp[1].is_finite(), "readback is finite DATA");
if (gp[0] - seed.pos[0][0]).abs() > 1e-3 {
moved = true;
}
}
assert!(moved, "the VRAM iteration actually advanced the layout");
}
/// **GPU LANE FRONT DOOR (self-skips headless):** `force_layout_gpu` returns
/// normalised `[-1, 1]²` positions that are finite DATA and topology-preserving — the
/// spring attraction actually contracts edges, so the mean *edge* length is well
/// below the mean length of an equal number of *random* node pairs. A broken wiring
/// (no iteration, un-normalised, or springs not pulling) trips one of these. This is
/// the GPU↔CPU equivalence at the layout level (bbox/topology, not byte-equality),
/// complementing the step-level `gpu_layout_matches_cpu_reference` parity proof.
#[cfg(feature = "gpu")]
#[test]
fn gpu_front_door_lays_out_in_box_and_contracts_edges() {
// THE GPU TURN: `gpu_layout_available` and `force_layout_gpu` each bring a device up.
let _gpu = facett_core::render::gputurn::hold();
if !gpu_layout_available() {
eprintln!("[gpu_front_door] no compute device — skipping");
return;
}
// A ring of triangles, well above the fan-out threshold.
let n = 6_000usize;
let mut edges = Vec::new();
for t in 0..n / 3 {
let (a, b, c) = (t * 3, t * 3 + 1, t * 3 + 2);
edges.push((a, b));
edges.push((b, c));
edges.push((c, a));
edges.push((a, (a + 3) % n));
}
let pos = force_layout_gpu(n, &edges, 120).expect("adapter present");
assert_eq!(pos.len(), n, "every node placed");
for p in &pos {
assert!(p[0].is_finite() && p[1].is_finite(), "finite DATA");
assert!(p[0].abs() <= 1.01 && p[1].abs() <= 1.01, "fit to [-1,1]² box: {p:?}");
}
let dist = |i: usize, j: usize| {
((pos[i][0] - pos[j][0]).powi(2) + (pos[i][1] - pos[j][1]).powi(2)).sqrt()
};
let mean_edge: f32 = edges.iter().map(|&(a, b)| dist(a, b)).sum::<f32>() / edges.len() as f32;
// A deterministic spread of non-adjacent pairs as the "random" baseline.
let mut sum = 0.0f32;
let m = 4_000usize;
for s in 0..m {
sum += dist((s * 7) % n, (s * 4099 + 1) % n);
}
let mean_pair = sum / m as f32;
assert!(mean_edge < mean_pair * 0.6, "springs contracted edges: edge {mean_edge} vs pair {mean_pair}");
}
/// INJECT-ASSERT: the WGSL exposes the entry point + workgroup size the pipeline
/// names, and `GpuNode` is the 32-byte two-vec4 the shader's `LNode` expects.
/// **The non-blocking job lays out EXACTLY what the blocking lane does.**
///
/// The job exists for the browser, where `wait_indefinitely` does not: it chunks
/// the same dispatches across polls and reads back through a callback. "Same
/// dispatches" is a claim, and this is where it is held — bit-for-bit equality
/// with [`force_layout_gpu_on`] on one device, with an `iters` that is NOT a
/// multiple of the chunk so the ragged last slice is exercised too. If the
/// chunking ever reordered the ping-pong, the two layouts would diverge here
/// rather than in a browser nobody is profiling.
#[cfg(feature = "gpu")]
#[test]
fn layout_job_matches_the_blocking_lane() {
let _gpu = facett_core::render::gputurn::hold();
let Some((device, queue)) = compute_device() else {
eprintln!("no wgpu adapter — skipping");
return;
};
let n = 300usize;
let edges: Vec<(usize, usize)> = (0..n).map(|i| (i, (i * 7 + 1) % n)).collect();
let iters = super::JOB_STEPS_PER_POLL * 3 + 4;
let blocking = force_layout_gpu_on(&device, &queue, n, &edges, iters);
let mut job = super::LayoutJob::start(&device, n, &edges, iters);
let mut polls = 0usize;
let got = loop {
polls += 1;
assert!(polls < 10_000, "the job never finished — a poll that cannot complete is a hang");
match job.poll(&device, &queue) {
super::JobState::Pending => continue,
super::JobState::Ready(p) => break p,
super::JobState::Failed(why) => panic!("the job failed on a live device: {why}"),
}
};
assert!(polls > 4, "finished in {polls} poll(s): the work was not chunked, so a browser would get one giant submit");
assert_eq!(got.len(), n);
assert_eq!(got, blocking, "the chunked job and the blocking lane laid the same graph out differently");
// A consumed job does not hand the layout out twice.
assert!(matches!(job.poll(&device, &queue), super::JobState::Failed(_)));
}
/// A community graph above [`super::GRID_MIN_NODES`]: rings of ~sqrt(n) nodes with
/// a local chord each, plus one long-range edge per 16 nodes.
#[cfg(feature = "gpu")]
fn community_graph(n: usize) -> Vec<(usize, usize)> {
let c = (n as f64).sqrt().ceil() as usize;
let per = n.div_ceil(c);
let mut e = Vec::new();
for k in 0..c {
let (lo, hi) = (k * per, ((k + 1) * per).min(n));
if hi <= lo + 3 {
continue;
}
let len = hi - lo;
for i in lo..hi {
e.push((i, lo + (i - lo + 1) % len));
e.push((i, lo + (i - lo + 7) % len));
}
}
for i in (0..n).step_by(16) {
e.push((i, (i * 7919 + 13) % n));
}
e
}
/// **The grid kernel pushes the way the all-pairs kernel pushes.** One step from the
/// same seed on the same device, with NO edges (springs would dominate one step and
/// hide the repulsion — measured: a sign-flipped far field still passed with the
/// community graph's edges in): the per-node displacement of the grid lane must
/// point where the exact lane's points (mean cosine > 0.95) with a similar length
/// (median ratio within 0.75..1.33). A pyramid that dropped a level, double-counted
/// the near field or mis-signed the far field fails one of the two.
#[cfg(feature = "gpu")]
#[test]
fn grid_kernel_step_tracks_the_all_pairs_step() {
let _gpu = facett_core::render::gputurn::hold();
let Some((device, queue)) = compute_device() else {
eprintln!("no wgpu adapter — skipping");
return;
};
let n = super::GRID_MIN_NODES + 3_000;
let g = GraphCsr::from_edges(n, &[], None);
let p = LayoutParams::default();
let seed = LayoutState::seed(n, p.ideal_len * (n as f32).sqrt());
let mut exact = super::GpuLayout::new(&device, &g, &seed);
let mut grid = super::GpuLayout::new_for(&device, &g, &seed);
assert_eq!(exact.kernel(), "all-pairs");
assert_eq!(grid.kernel(), "grid", "{n} nodes is above GRID_MIN_NODES");
exact.iterate(&device, &queue, &g, &p, 1);
grid.iterate(&device, &queue, &g, &p, 1);
let a = exact.read_positions(&device, &queue);
let b = grid.read_positions(&device, &queue);
let mut cos_sum = 0.0f64;
let mut ratios = Vec::with_capacity(n);
for i in 0..n {
let da = [a[i][0] - seed.pos[i][0], a[i][1] - seed.pos[i][1]];
let db = [b[i][0] - seed.pos[i][0], b[i][1] - seed.pos[i][1]];
let la = (da[0] * da[0] + da[1] * da[1]).sqrt();
let lb = (db[0] * db[0] + db[1] * db[1]).sqrt();
assert!(lb.is_finite(), "grid step produced a non-finite position at {i}");
if la > 1e-6 && lb > 1e-6 {
cos_sum += f64::from((da[0] * db[0] + da[1] * db[1]) / (la * lb));
ratios.push(lb / la);
}
}
let cos = cos_sum / ratios.len().max(1) as f64;
ratios.sort_by(|x, y| x.partial_cmp(y).unwrap());
let med = ratios[ratios.len() / 2];
assert!(cos > 0.95, "grid displacement direction diverges from all-pairs: mean cosine {cos:.3}");
assert!((0.75..1.33).contains(&med), "grid displacement length off: median ratio {med:.3}");
}
/// **The grid lane lays a large graph out** (bbox + topology, like the front-door
/// test for the all-pairs lane), and the non-blocking job reproduces the blocking
/// lane bit for bit — the pyramid sums integer atomics and reduces in a fixed order,
/// so the grid kernel is deterministic too.
#[cfg(feature = "gpu")]
#[test]
fn grid_lane_lays_out_and_the_job_matches_it() {
let _gpu = facett_core::render::gputurn::hold();
let Some((device, queue)) = compute_device() else {
eprintln!("no wgpu adapter — skipping");
return;
};
let n = super::GRID_MIN_NODES + 1_000;
let edges = community_graph(n);
let iters = 60;
let pos = force_layout_gpu_on(&device, &queue, n, &edges, iters);
assert_eq!(pos.len(), n);
for q in &pos {
assert!(q[0].is_finite() && q[1].is_finite() && q[0].abs() <= 1.01 && q[1].abs() <= 1.01, "{q:?}");
}
let dist = |i: usize, j: usize| ((pos[i][0] - pos[j][0]).powi(2) + (pos[i][1] - pos[j][1]).powi(2)).sqrt();
let mean_edge: f32 = edges.iter().map(|&(a, b)| dist(a, b)).sum::<f32>() / edges.len() as f32;
let m = 4_000usize;
let mean_pair: f32 = (0..m).map(|s| dist((s * 7) % n, (s * 4099 + 1) % n)).sum::<f32>() / m as f32;
assert!(mean_edge < mean_pair * 0.6, "springs contracted edges: edge {mean_edge} vs pair {mean_pair}");
let mut job = super::LayoutJob::start(&device, n, &edges, iters);
assert_eq!(job.kernel(), "grid");
let got = loop {
match job.poll(&device, &queue) {
super::JobState::Pending => continue,
super::JobState::Ready(p) => break p,
super::JobState::Failed(why) => panic!("grid job failed: {why}"),
}
};
assert_eq!(got, pos, "the chunked grid job and the blocking grid lane differ");
}
/// **Three lanes, one small graph, one answer.** The CPU reference (`step_cpu`), the
/// all-pairs GPU kernel (`cs_layout`) and the grid-pyramid GPU kernel
/// (`layout_grid.wgsl`, forced below [`super::GRID_MIN_NODES`] with
/// [`super::GpuLayout::new_grid`]) run the same seed:
///
/// * all-pairs GPU vs CPU, 30 steps: the same function, so every position within
/// `ALL_PAIRS_TOL` (absolute, world units; the layout spans ~2·sqrt(n));
/// * grid GPU vs CPU, ONE step: the grid sums the far field approximately (θ ≤ 1), so
/// its per-node displacement must point where the exact one points (mean cosine >
/// 0.95) with a similar length (median ratio 0.75..1.33) — the bars of
/// `grid_kernel_step_tracks_the_all_pairs_step`;
/// * grid GPU vs CPU, 30 steps: positions are NOT compared — a force layout amplifies a
/// per-step approximation (MEASURED t14s Radeon 890M 2026-09-29: max |Δ| 10.3 of a
/// 28.7 span, mean cosine 0.968 after 30 steps), so the settled picture is held by
/// its shape instead: mean edge length and span within `GRID_SHAPE_TOL` of the CPU's.
///
/// It also prints one timing row per lane (`[parity-timing] …`), the per-step cost on
/// this box, which is what `.nornir/facett-benchmarks.md` quotes for the small graph.
#[cfg(feature = "gpu")]
#[test]
fn grid_all_pairs_and_cpu_agree_on_a_small_graph() {
const ALL_PAIRS_TOL: f32 = 1e-2;
const GRID_SHAPE_TOL: f32 = 0.2;
let _gpu = facett_core::render::gputurn::hold();
let Some((device, queue)) = compute_device() else {
eprintln!("no wgpu adapter — skipping");
return;
};
let n = 400usize;
let e: Vec<(u32, u32)> = community_graph(n).iter().map(|&(a, b)| (a as u32, b as u32)).collect();
let g = GraphCsr::from_edges(n, &e, None);
let p = LayoutParams::default();
let iters = 30usize;
let seed = LayoutState::seed(n, p.ideal_len * (n as f32).sqrt());
let cpu1 = step_cpu(&seed, &g, &p);
let t = std::time::Instant::now();
let mut cpu = seed.clone();
for _ in 0..iters {
cpu = step_cpu(&cpu, &g, &p);
}
let cpu_ms = t.elapsed().as_secs_f64() * 1e3;
// Pipelines compile in `new`/`new_grid`, outside the clock.
let run = |mut gl: super::GpuLayout, steps: usize| {
let t = std::time::Instant::now();
gl.iterate(&device, &queue, &g, &p, steps);
let pos = gl.read_positions(&device, &queue);
(pos, t.elapsed().as_secs_f64() * 1e3)
};
let grid_new = || super::GpuLayout::new_grid(&device, &g, &seed).expect("a compute device with 7 storage buffers");
let (exact, exact_ms) = run(super::GpuLayout::new(&device, &g, &seed), iters);
assert_eq!(grid_new().kernel(), "grid", "new_grid forces the grid kernel at {n} nodes");
let (grid, grid_ms) = run(grid_new(), iters);
let (grid1, _) = run(grid_new(), 1);
assert_eq!(exact.len(), n);
assert_eq!(grid.len(), n);
// all-pairs vs CPU after `iters` steps: position parity.
let mut exact_max = 0.0f32;
for i in 0..n {
let (a, c, b) = (exact[i], cpu.pos[i], grid[i]);
assert!(a[0].is_finite() && a[1].is_finite() && b[0].is_finite() && b[1].is_finite(), "non-finite at {i}");
exact_max = exact_max.max((a[0] - c[0]).abs()).max((a[1] - c[1]).abs());
}
// grid vs CPU, one step: displacement direction + length.
let mut cos_sum = 0.0f64;
let mut ratios = Vec::with_capacity(n);
for (i, g1) in grid1.iter().enumerate().take(n) {
let dc = [cpu1.pos[i][0] - seed.pos[i][0], cpu1.pos[i][1] - seed.pos[i][1]];
let dg = [g1[0] - seed.pos[i][0], g1[1] - seed.pos[i][1]];
let (lc, lg) = ((dc[0] * dc[0] + dc[1] * dc[1]).sqrt(), (dg[0] * dg[0] + dg[1] * dg[1]).sqrt());
if lc > 1e-6 && lg > 1e-6 {
cos_sum += f64::from((dc[0] * dg[0] + dc[1] * dg[1]) / (lc * lg));
ratios.push(lg / lc);
}
}
let cos1 = cos_sum / ratios.len().max(1) as f64;
ratios.sort_by(|x, y| x.partial_cmp(y).unwrap());
let med1 = ratios[ratios.len() / 2];
// grid vs CPU after `iters` steps: the shape.
let shape = |pos: &[[f32; 2]]| {
let (mut lo, mut hi) = ([f32::MAX; 2], [f32::MIN; 2]);
for q in pos {
for k in 0..2 {
lo[k] = lo[k].min(q[k]);
hi[k] = hi[k].max(q[k]);
}
}
let edge: f32 = e
.iter()
.map(|&(a, b)| {
let (a, b) = (pos[a as usize], pos[b as usize]);
((a[0] - b[0]).powi(2) + (a[1] - b[1]).powi(2)).sqrt()
})
.sum::<f32>()
/ e.len() as f32;
(edge, (hi[0] - lo[0]).max(hi[1] - lo[1]))
};
let (edge_c, span_c) = shape(&cpu.pos);
let (edge_g, span_g) = shape(&grid);
let per = |ms: f64| ms / iters as f64;
eprintln!(
"[parity-timing] n={n} edges={} steps={iters} cpu-1 {:.3} ms/step · gpu all-pairs {:.3} ms/step · gpu grid {:.3} ms/step \
(incl. readback) · max|all-pairs−cpu| {exact_max:.2e} · grid 1-step cos {cos1:.4} median ratio {med1:.3} · \
{iters}-step mean edge grid {edge_g:.3} / cpu {edge_c:.3}, span grid {span_g:.2} / cpu {span_c:.2}",
e.len(),
per(cpu_ms),
per(exact_ms),
per(grid_ms),
);
assert!(exact_max < ALL_PAIRS_TOL, "all-pairs GPU left the CPU reference: max |Δ| {exact_max}");
assert!(cos1 > 0.95, "grid step direction diverges from the CPU step: mean cosine {cos1:.4}");
assert!((0.75..1.33).contains(&med1), "grid step length off: median ratio {med1:.3}");
assert!((edge_g / edge_c - 1.0).abs() < GRID_SHAPE_TOL, "grid mean edge {edge_g} vs cpu {edge_c}");
assert!((span_g / span_c - 1.0).abs() < GRID_SHAPE_TOL, "grid span {span_g} vs cpu {span_c}");
}
/// The host-device seam: nothing installed means nothing borrowed, an install is
/// visible on this thread, and `clear` really forgets it.
#[cfg(feature = "gpu")]
#[test]
fn the_host_device_seam_installs_and_clears() {
let _gpu = facett_core::render::gputurn::hold();
super::clear_host_device();
assert!(!super::host_device_installed());
assert!(super::with_host_device(|_, _| ()).is_none());
let Some((device, queue)) = compute_device() else {
eprintln!("no wgpu adapter — skipping");
return;
};
assert!(super::install_host_device(&device, &queue));
assert!(super::host_device_installed());
assert_eq!(super::with_host_device(|d, _| d.limits().max_storage_buffers_per_shader_stage >= 4), Some(true));
super::clear_host_device();
assert!(!super::host_device_installed());
}
#[cfg(feature = "gpu")]
#[test]
fn layout_shader_entry_and_node_size() {
assert!(LAYOUT_WGSL.contains("fn cs_layout"));
assert!(LAYOUT_WGSL.contains("@workgroup_size(64)"));
assert_eq!(std::mem::size_of::<GpuNode>(), 32);
}
#[cfg(feature = "arrow")]
#[test]
fn csr_streams_from_an_arrow_edge_batch() {
use arrow_array::{Float64Array, Int64Array};
let src = Int64Array::from(vec![0, 1, 2, 0]);
let dst = Int64Array::from(vec![1, 2, 0, 2]);
let w = Float64Array::from(vec![3.0, 1.0, 1.0]);
let g = GraphCsr::from_arrow(3, &src, &dst, Some(&w));
assert_eq!(g.neighbours(0), &[1, 2]);
assert_eq!(g.weights[0], 3.0, "weight column streamed in");
// It lays out to finite, deterministic DATA.
let s = relax_cpu(&g, &LayoutParams::default(), 80);
assert!(s.pos.iter().all(|q| q[0].is_finite() && q[1].is_finite()));
}
}