resonance-engine/TRADE_LBM_DIFF_SPEC.md

trade_lbm_v1 \u2014 v5 \u2192 trade_v1 Kernel Diff-Spec

Purpose. Concrete, file-level diff from khra_gixx_1024_v5.cu to trade_lbm_v1.cu. No CUDA written yet; this is the spec the kernel author works against. Companion to TRADE_LBM_ARCHITECTURE.md v4.

Base file. D:\resonance-engine\cuda\khra_gixx_1024_v5.cu (1284 lines).

New file. D:\resonance-engine\cuda\trade_lbm_v1.cu (will start as a copy; sections marked KEEP retain identical code, ALTER replaces with new behaviour, ADD is brand new, REMOVE is deleted).

---

Section A \u2014 What is kept verbatim (KEEP)

These are the parts of v5 that are pure LBM mechanics and have no market dependency. Copy unchanged:

v5 lines	What	Why kept
32\u201337	M_PI, Q=9	unchanged
39\u201346	CUDA_CHECK macro	unchanged
48\u201352	d_cx, d_cy, d_w constants	D2Q9 lattice constants, identical
66\u201382	streaming_kernel	streaming step is the same regardless of grid meaning
121\u2013135	compute_rho	macroscopic density same
137\u2013175	compute_velocity_stress_v4	macroscopic ux/uy/stress identical
234\u2013410	host helpers: crc32_init, crc32_compute, save_checkpoint, load_checkpoint	infrastructure
413\u2013506	telemetry ring buffer	infrastructure
510\u2013575	ZMQ resilience + JSON helpers	infrastructure
747\u2013768	crash_handler / shutdown_handler / signal wiring	infrastructure

---

Section B \u2014 What changes (ALTER)

B.1 Grid dimensions (lines 35\u201336)

-#define NX 1024
-#define NY 1024
+#define NX 512   // price-tick columns (centred on rolling mid)
+#define NY 512   // time rows (row 0 = newest, NY-1 = oldest)

Per architecture \u00a72. Lower memory cost than 1024\u00b2 \u2014 NXNYQ*sizeof(float) = 9 MB vs 36 MB for v5.

B.2 Device-side parameters block (lines 54\u201355)

-__device__ float d_khra_amp = 0.03f;
-__device__ float d_gixx_amp = 0.008f;
+__device__ float d_oi_drive   = 0.0f;       // slow large-scale forcing amp; 0 = inactive
+__device__ float d_flow_drive = 0.0f;       // fast fine-scale forcing amp; 0 = inactive
+__device__ float d_tick_size  = 0.5f;       // USD per column; instrument-dependent
+__device__ float d_mid_price  = 0.0f;       // current mid (host updates on recenter)
+__device__ int   d_recenter_drift = 0;      // cumulative integer-column drift since boot
+
+// Per-column omega (replaces single global omega). Stored in device global memory,
+// not __constant__ because it's updated O(once per snapshot) which exceeds the
+// 64KB cbank limit budget.
+__device__ float d_omega_col[NX];
+
+// Per-column depth-weighted equilibrium target. Updated by host on book events.
+__device__ float d_rho_eq[NX];     // target density
+__device__ float d_ux_eq[NX];      // target u_x (bias from depth imbalance)
+// u_y_eq is a single scalar (constant upward streaming velocity), defined below.
+__device__ float d_uy_eq = 0.0f;   // typically 0 in v1; reserved for time-axis tuning

B.3 Wave forcing function (lines 57\u201364)

Rename and re-anchor to market-derived amplitudes. Wavelengths unchanged.

-__device__ float khra_gixx_wave_1024(int x, int y, int cycle) {
-    float khra = sinf(2.0f*M_PI*x/128.0f + cycle*0.025f) *
-                 cosf(2.0f*M_PI*y/128.0f + cycle*0.015f) * d_khra_amp;
-    float gixx = sinf(2.0f*M_PI*x/8.0f + cycle*0.4f) *
-                 cosf(2.0f*M_PI*y/8.0f + cycle*0.35f) * d_gixx_amp;
-    float asymmetry_factor = 1.0f + sinf(cycle * 0.05f) * 0.5f;
-    return khra + gixx * asymmetry_factor;
-}
+__device__ float market_wave(int x, int y, int cycle) {
+    // Large-scale OI-driven wave: \u03bb = NX/4 = 128 columns (matches v5 khra \u03bb).
+    float oi   = sinf(2.0f*M_PI*x/(NX/4.0f) + cycle*0.025f) *
+                 cosf(2.0f*M_PI*y/(NX/4.0f) + cycle*0.015f) * d_oi_drive;
+    // Fine-scale flow-driven wave: \u03bb = 8 columns (matches v5 gixx \u03bb).
+    float flow = sinf(2.0f*M_PI*x/8.0f + cycle*0.4f) *
+                 cosf(2.0f*M_PI*y/8.0f + cycle*0.35f) * d_flow_drive;
+    // Drop the asymmetry_factor envelope \u2014 it was a v5 substrate artefact;
+    // market amplitudes already carry their own time variation via host updates.
+    return oi + flow;
+}

Note: ky = kx / 2 (the 2:1 spatial anisotropy) is removed. That was a physics-lattice geometric choice; on a price-time grid it would inject unwanted bias. The collision kernel below applies the wave as a scalar density nudge, not as separately-anisotropic ux/uy nudges.

B.4 Collision kernel (lines 83\u2013119) \u2014 largest change

This is the heart of the diff. Two structural changes:

1. Per-column omega (read from d_omega_col[x]) replaces scalar omega argument.
2. Equilibrium target uses (d_rho_eq[x], d_ux_eq[x], d_uy_eq) instead of the cell's local (rho, ux, uy).

-__global__ void collide_kernel_khragixx(float* f, float omega, int cycle) {
+__global__ void collide_kernel_trade(float* f, int cycle) {
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     if (x >= NX || y >= NY) return;
     int idx = (y * NX + x) * Q;
     float rho = 0.0f, ux = 0.0f, uy = 0.0f;
     for (int i = 0; i < Q; i++) {
         rho += f[idx + i];
         ux += f[idx + i] * d_cx[i];
         uy += f[idx + i] * d_cy[i];
     }
     if (rho < 0.1f) rho = 0.1f;
     if (rho > 10.0f) rho = 10.0f;
     ux /= rho;
     uy /= rho;
-    float u_mag = sqrtf(ux*ux + uy*uy);
-    if (u_mag > 0.25f) { ux = ux * 0.25f / u_mag; uy = uy * 0.25f / u_mag; }
-    float kx = khra_gixx_wave_1024(x, y, cycle);
-    float ky = kx * 0.5f;
-    ux += kx;
-    uy += ky;
-    u_mag = sqrtf(ux*ux + uy*uy);
-    if (u_mag > 0.25f) { ux = ux * 0.25f / u_mag; uy = uy * 0.25f / u_mag; }
-    float u_sq = ux*ux + uy*uy;
-    for (int i = 0; i < Q; i++) {
-        float eu = d_cx[i]*ux + d_cy[i]*uy;
-        float feq = d_w[i] * rho * (1.0f + 3.0f*eu + 4.5f*eu*eu - 1.5f*u_sq);
-        f[idx + i] = f[idx + i] - omega * (f[idx + i] - feq);
-        if (isnan(f[idx + i]) || isinf(f[idx + i])) f[idx + i] = d_w[i];
-    }
+    // Mach safety on observed velocity
+    float u_mag = sqrtf(ux*ux + uy*uy);
+    if (u_mag > 0.25f) { ux = ux * 0.25f / u_mag; uy = uy * 0.25f / u_mag; }
+
+    // Apply market wave as a scalar density nudge (not a velocity push).
+    // CRITICAL: the wave modulates the EQUILIBRIUM TARGET, not the current cell
+    // state. It is added to d_rho_eq[x] to produce rho_target. Do NOT add it to
+    // the locally-observed `rho` above — that would push the field instead of
+    // shifting where the field wants to relax to, and would destroy the
+    // book-as-attractor property that justifies this whole architecture.
+    float wave = market_wave(x, y, cycle);
+    float rho_target = d_rho_eq[x] + wave;
+    if (rho_target < 0.1f) rho_target = 0.1f;
+    if (rho_target > 10.0f) rho_target = 10.0f;
+
+    // Equilibrium target from depth-weighted book.
+    float ux_eq = d_ux_eq[x];
+    float uy_eq = d_uy_eq;          // global, typically 0
+    float ueq_sq = ux_eq*ux_eq + uy_eq*uy_eq;
+
+    // Per-column relaxation rate (from spread tier).
+    float omega_x = d_omega_col[x];
+    if (omega_x < 0.05f) omega_x = 0.05f;
+    if (omega_x > 1.99f) omega_x = 1.99f;
+
+    for (int i = 0; i < Q; i++) {
+        float eu  = d_cx[i]*ux_eq + d_cy[i]*uy_eq;
+        float feq = d_w[i] * rho_target * (1.0f + 3.0f*eu + 4.5f*eu*eu - 1.5f*ueq_sq);
+        f[idx + i] = f[idx + i] - omega_x * (f[idx + i] - feq);
+        if (isnan(f[idx + i]) || isinf(f[idx + i])) f[idx + i] = d_w[i] * rho_target;
+    }

Subtle but important: feq is now computed against the book-derived (rho_target, ux_eq, uy_eq), not against the local cell's observed (rho, ux, uy). That is what makes the book itself the relaxation target. The local rho/ux/uy are only computed to satisfy Mach safety and to detect numerical blow-ups; they do not feed the equilibrium.

B.5 Inject kernel (lines 180\u2013203) \u2014 renamed + price-aware

-__global__ void inject_density_kernel(float* f,
-                                      float cx, float cy,
-                                      float sigma, float strength) {
+// inject_trade_kernel: a trade event at (column, row=0, sigma_x ticks, signed_size)
+// Note: trades always land at row 0 (newest). Sigma now in ticks along X only;
+// Y spread is one row (the trade happens at one wallclock instant).
+__global__ void inject_trade_kernel(float* f,
+                                    int    col_centre,
+                                    float  sigma_x,
+                                    float  signed_dRho)
+{
     int x = blockIdx.x * blockDim.x + threadIdx.x;
-    int y = blockIdx.y * blockDim.y + threadIdx.y;
     if (x >= NX || y >= NY) return;
+    int y = 0;   // always row 0
-    float dx = (float)x - cx;
-    float dy = (float)y - cy;
-    if (dx >  NX * 0.5f) dx -= NX;
-    if (dx < -NX * 0.5f) dx += NX;
-    if (dy >  NY * 0.5f) dy -= NY;
-    if (dy < -NY * 0.5f) dy += NY;
-    float r2 = dx * dx + dy * dy;
-    float inv_2sig2 = 1.0f / (2.0f * sigma * sigma);
-    float delta_rho = strength * expf(-r2 * inv_2sig2);
+    // X-axis: clamp distance (no wrap \u2014 price columns are NOT periodic).
+    float dx = (float)x - (float)col_centre;
+    float inv_2sig2 = 1.0f / (2.0f * sigma_x * sigma_x);
+    float delta_rho = signed_dRho * expf(-dx * dx * inv_2sig2);
     int f_idx = (y * NX + x) * Q;
     for (int i = 0; i < Q; i++) f[f_idx + i] += d_w[i] * delta_rho;
 }

Launch this kernel with grid = (NX/blockDim.x, 1) instead of full 2D \u2014 only one row is touched. Periodicity is dropped on X because price columns are bounded; trades outside the window go to off_grid_mass (host accumulator).

B.6 Asymmetry helper (lines 205\u2013214) \u2014 rebased on rho_eq

-float calculate_asymmetry_magnifying(float* h_rho) {
+// Asymmetry now measured relative to the per-column rho_eq (depth target),
+// not against a uniform 1.0 background. This makes the metric meaningful
+// against a non-uniform equilibrium.
+float calculate_asymmetry_magnifying(float* h_rho, const float* h_rho_eq) {
     float sum_sq = 0.0f;
-    for (int i = 0; i < NX * NY; i++) {
-        float dev = h_rho[i] - 1.0f;
+    for (int idx = 0; idx < NX * NY; idx++) {
+        int x = idx % NX;
+        float dev = h_rho[idx] - h_rho_eq[x];
         sum_sq += dev * dev;
     }
     return (sum_sq / (NX * NY)) * 100.0f;
 }

B.7 Telemetry JSON (line ~1136)

-"\"coherence\":%.4f,\"asymmetry\":%.4f,"
+"\"coherence\":%.4f,\"asymmetry\":%.4f,\"regime_product\":%.4f,"

regime_product = coherence * asymmetry computed host-side at telemetry emit. Adds one float multiply and ~20 bytes per snapshot.

B.8 Command handler (line 595, especially lines 617\u2013640 and 703\u2013740)

Replace set_khra_amp / set_gixx_amp / inject_density with new commands:

New command	JSON shape	What it does
set_oi_drive	{"cmd":"set_oi_drive","value":<float>}	sets d_oi_drive, clamp [0, 0.2]
set_flow_drive	{"cmd":"set_flow_drive","value":<float>}	sets d_flow_drive, clamp [0, 0.1]
set_omega_profile	{"cmd":"set_omega_profile","values":[NX floats]}	bulk write d_omega_col[], each clamped [0.05, 1.99]
set_omega_col	{"cmd":"set_omega_col","col":<int>,"value":<float>}	single-column update for cheap incremental refresh
set_book	{"cmd":"set_book","bid":[NX floats],"ask":[NX floats]}	bulk write depths; host derives d_rho_eq and d_ux_eq and uploads
inject_trade	`{"cmd":"inject_trade","price":,"side":"buy"	"sell","size":,"aggressor":}`
set_mid	{"cmd":"set_mid","price":<f>}	manual recenter
set_tick_size	{"cmd":"set_tick_size","value":<f>}	rare; restart-equivalent

Keep verbatim: pause, resume, reset, snapshot, health_check, reset_equilibrium, stress_snapshot_now, export_timeseries.

Remove entirely: set_khra_amp, set_gixx_amp, inject_density.

B.9 Main loop (line 770+ and the cycle % 10 telemetry block at line 1050)

Minor changes:

1. Call collide_kernel_trade<<<...>>>(d_f, cycle) instead of collide_kernel_khragixx<<<...>>>(d_f, h_omega, cycle).
2. Drop h_omega host scalar.
3. Add h_rho_eq[NX], h_ux_eq[NX], h_omega_col[NX], h_bid_depth[NX], h_ask_depth[NX] host arrays and a small helper recompute_equilibrium_from_book() that runs whenever set_book arrives. Formula is fixed (see B.9.1 below); kernel author MUST NOT improvise.
4. Add off_grid_mass host accumulator updated when inject_trade finds column out of range.

B.9.1 recompute_equilibrium_from_book() formula — fixed

The book-to-equilibrium mapping is part of the architecture, not a kernel-author choice. Use exactly:

// Constants (compile-time)
#define RHO_EQ_MIN        0.5f     // floor for rho_eq band
#define RHO_EQ_MAX        2.0f     // ceiling for rho_eq band
#define RHO_EQ_NOMINAL    1.0f     // target mean across columns
#define MAX_UX_BIAS       0.10f    // velocity-bias cap; well below Mach safety 0.25
#define DEPTH_EPS         1e-6f    // divide-by-zero guard for empty columns

static void recompute_equilibrium_from_book(void) {
    // 1. Mean total depth across columns — used as the normalisation constant
    //    so that rho_eq[x] is centred near RHO_EQ_NOMINAL regardless of
    //    absolute book size.
    double total_sum = 0.0;
    for (int x = 0; x < NX; x++) {
        total_sum += (double)(h_bid_depth[x] + h_ask_depth[x]);
    }
    float mean_total = (float)(total_sum / (double)NX);
    if (mean_total < DEPTH_EPS) mean_total = DEPTH_EPS;
    float norm = mean_total / RHO_EQ_NOMINAL;   // divisor that maps mean->1.0

    // 2. Per-column rho_eq (clamped) and ux_eq (imbalance-driven).
    for (int x = 0; x < NX; x++) {
        float total = h_bid_depth[x] + h_ask_depth[x];
        float rho_eq = total / norm;
        if (rho_eq < RHO_EQ_MIN) rho_eq = RHO_EQ_MIN;
        if (rho_eq > RHO_EQ_MAX) rho_eq = RHO_EQ_MAX;
        h_rho_eq[x] = rho_eq;

        float denom = total > DEPTH_EPS ? total : DEPTH_EPS;
        float imbalance = (h_bid_depth[x] - h_ask_depth[x]) / denom;  // [-1, 1]
        h_ux_eq[x] = imbalance * MAX_UX_BIAS;
    }

    // 3. Upload both arrays in one DMA pair.
    CUDA_CHECK(cudaMemcpyToSymbol(d_rho_eq, h_rho_eq, sizeof(h_rho_eq)));
    CUDA_CHECK(cudaMemcpyToSymbol(d_ux_eq,  h_ux_eq,  sizeof(h_ux_eq)));
}

Why these constants.

• RHO_EQ_NOMINAL = 1.0 matches the D2Q9 weight normalisation; equilibrium density of 1.0 is the natural LBM rest state.
• RHO_EQ_MIN/MAX = 0.5 / 2.0 is a ~2× dynamic range — enough that a 4×-deeper-than-average column is visible, but small enough that no column ever pushes the BGK collision outside well-tested regimes.
• MAX_UX_BIAS = 0.10 is 0.4× the Mach safety ceiling 0.25. A 100% one-sided book (bid only) produces ux_eq = 0.10. The collision kernel's Mach clamp at 0.25 is then a hard fault detector, not a routine clip.
• The normalisation uses mean total depth across columns, not a fixed constant, so the equilibrium band auto-rescales to whatever absolute book size the venue is showing. This means BTC at 200 contracts/level and ETH at 2000 contracts/level both produce the same equilibrium magnitudes.

This function runs on the host on every set_book command, which architecture §3.5 caps at ≤1 Hz. Cost: O(NX) = O(512) floats per call — negligible.

---

Section C \u2014 What is added (ADD)

C.1 Recentre routine (host)

// Called after set_mid or auto-trigger when |cumulative drift| > NX/16.
// Shifts the entire f array by `shift` integer columns (positive = field
// scrolls right). Implemented as a CUDA memcpy with offset; out-of-range
// columns re-initialised to local equilibrium.
static void recentre_field(float* d_f, int shift);

Trigger logic: host tracks mid_price and mid_price_at_grid_centre. When |price - centre| > RECENTRE_THRESHOLD * tick_size, call recentre_field, update d_mid_price, set recenter_event telemetry flag for one snapshot.

C.2 Per-column profile telemetry kernel

// Reduces NY rows down to per-column scalars for the profile observables
// listed in architecture \u00a75. One block per column, reduction across Y.
__global__ void compute_column_profiles(
    const float* rho, const float* ux, const float* uy,
    const float* sxx, const float* sxy,
    float* density_profile,    // [NX]
    float* vorticity_profile,  // [NX]
    float* divergence_profile, // [NX]
    float* stress_xx_profile); // [NX]

Host DMA's the four NX-length arrays back, includes in telemetry as compact JSON arrays (every Kth column with significant activity, K chosen so JSON stays under ~4KB per snapshot).

Telemetry byte budget — confirm before first run. Naive size at NX=512 is:

Component	Size
4 profile arrays × 512 floats × ~8 chars per JSON float (-1.234e-2,)	~16 KB
Global scalars (rho_total, asymmetry, coherence, regime_product, ...)	~0.5 KB
Event detector output (typically 0–10 events, struct ≈ 80 bytes JSON)	~1 KB max

Unsubsampled ≈ 17 KB per snapshot. Even at K=2 (every other column) we are at ~9 KB — well above the 4 KB target. Mitigations, applied in order until budget met:

1. Subsample columns with K=4 (every 4th column) when no event detector has fired in the last N snapshots → ~4 KB.
2. Emit profiles as float16 hex-encoded (4 chars/value) instead of decimal JSON → halves the dominant cost.
3. Switch profile transport from JSON to ZMQ multipart binary frames (frame 0 = JSON scalars + events, frame 1 = raw float32 binary blob of profiles, 8 KB exactly).

ZMQ socket configuration. v5 telemetry is scalar-only (~1 KB/snapshot at 10 Hz = 10 KB/s). v1 telemetry is at least 5–16× larger. Before first run, the kernel author MUST verify on the publisher (d_telemetry_pub):

• ZMQ_SNDHWM ≥ 1024 (default in v5 should be checked).
• ZMQ_SNDBUF ≥ 256 KB (raise from OS default if needed).
• Snapshot cadence stays at cycle % 10 == 0 initially — do not raise frequency until profile size is bounded.

If any subscriber (Fractonaut, dashboard) starts dropping or lagging, that is the size budget being exceeded. Fall back to mitigation 3 (binary frames) before tuning anything else.

C.3 Event detectors (host)

Functions, not kernels \u2014 cheap on NX=512:

typedef struct {
    int   col;         // column index of detected feature
    float price;       // absolute price at col (translated via d_mid_price + d_tick_size)
    float magnitude;   // z-score
    int   persisted_rows;
} structural_event;

int detect_breakouts      (const float* divergence_profile,
                           const float* density_profile,
                           structural_event* out, int max_out);
int detect_consolidation  (const float* vorticity_profile,
                           const float* divergence_profile,
                           structural_event* out, int max_out);
int detect_support_levels (const float* density_profile_history, // [NY][NX] ring
                           structural_event* out, int max_out);
int detect_resistance     (const float* density_profile_history, // [NY][NX] ring
                           structural_event* out, int max_out);

Detector thresholds (z-score floors, persistence row counts) start as compile-time defaults and become runtime-tunable in a later pass.

C.4 last_trade_age_s tracking (host)

static double last_trade_unix = 0.0;
// In telemetry assembly:
double now = ...;
int last_trade_age_s = (int)(now - last_trade_unix);
// Snippet added to JSON: "last_trade_age_s":%d

Fractonaut prompt rule (\u00a79 staleness handling): if age > 60, treat as drifting toward equilibrium, structural-only.

C.5 New ZMQ ports

Change all ZMQ bind ports from 5556\u20135560 to 5566\u20135570 so trade lattice coexists with physics lattice. Host-side: update endpoint strings in main(). No kernel impact.

---

Section D \u2014 What is removed (REMOVE)

• __device__ float d_khra_amp and d_gixx_amp (replaced).
• khra_gixx_wave_1024 (replaced by market_wave).
• set_khra_amp, set_gixx_amp, inject_density command handlers.
• The asymmetry_factor = 1.0f + sinf(cycle * 0.05f) * 0.5f envelope: it was a v5 visualisation trick, has no market analogue.
• The ky = kx * 0.5f anisotropy in collide_kernel_khragixx: rejected per architecture \u00a72.
• The cy argument to inject_density_kernel (trades always land at y=0).
• The periodic-distance code in inject (if (dx > NX * 0.5f) dx -= NX; etc): X is not periodic on price grid.

---

Section E \u2014 Build and deploy

Build command unchanged from v5 (assumes the existing CMake / nvcc setup discovers the new .cu file automatically). If the build script targets v5 by name:

- nvcc -O3 -arch=sm_75 khra_gixx_1024_v5.cu -o khra_gixx_1024_v5 \
+ nvcc -O3 -arch=sm_75 trade_lbm_v1.cu       -o trade_lbm_v1 \
       -lzmq -lnvidia-ml -rdynamic

-arch=sm_75 should match the actual GPU (the v5 build uses this on Beast; confirm at compile time).

Output binary: /mnt/d/Resonance_Engine/beast-build/trade_lbm_v1 (parallels the existing khra_gixx_1024_v5 location). Launch with --agent-owner=RESONANCE per the BEAST ledger.

---

Section F \u2014 What is intentionally NOT in v1

• Multi-instrument support. One symbol per binary.
• GPU-side event detection. CPU is enough at NX=512.
• Live trading hooks. Trade lattice output goes to Fractonaut + JSON only.
• Adaptive tick sizing. Constant per startup.
• Auto-tuning of TAKER_WEIGHT, DEPTH_WEIGHT, BIAS_WEIGHT. All compile-time in v1; tuned during validation in \u00a710.

---

Section G \u2014 Verification checklist for the kernel author

Before declaring trade_lbm_v1 ready for the \u00a710 validation gate:

1. set_oi_drive value=0.0 + set_flow_drive value=0.0 + no set_book calls (so d_rho_eq is zero-initialised) \u2192 lattice should sit at zero everywhere, regime_product ~ 0. If non-zero, equilibrium reads aren't wired.
2. set_book with uniform bid=ask=1.0 across all columns \u2192 lattice should relax to regime_product close to the trade-lattice equivalent of the v5 "normal" attractor. Magnitude TBD by calibration but must be flat across columns.
3. set_book with one heavy column \u2192 corresponding density_profile[that column] must build a persistent peak; relaxation rate matches expected 1/\u03c9 timescale for that column's omega.
4. inject_trade with price out of range \u2192 off_grid_mass increments; nothing else.
5. set_mid larger than RECENTRE_THRESHOLD away \u2192 recenter_event fires; density_profile shifts by the expected number of columns; old peaks stay at their absolute price (verified by structural-event detector reporting same prices, different columns).
6. Send a synthetic trade at known price with known size \u2192 verify density_profile peak at the correct column with the correct magnitude (size * TAKER_WEIGHT if aggressor flag set). Note (added 2026-06-08 after first-run verification): Test 6 must run with omega_profile <= 0.1 across all columns to observe one-shot injection impact. Default omega=1.0 correctly absorbs single trades in one cycle — this is correct BGK behaviour when the book is the attractor, not a bug. Without lowering omega before injecting, the next telemetry snapshot will show the book-derived equilibrium (flat) regardless of the injection. If all six pass, kernel is ready for replay against April HL data per \u00a710.

---

Section H \u2014 Estimated kernel-author effort

Item	Approx LOC delta vs v5
ALTER changes (collide kernel, inject kernel, wave function, telemetry JSON, command handler)	~250 lines edited
ADD: recentre_field	~40 lines
ADD: compute_column_profiles	~50 lines
ADD: 4 event detectors (host)	~150 lines
ADD: recompute_equilibrium_from_book (formula in B.9.1) + book/omega/etc host scratch buffers	~100 lines (formula is ~40 of these, copy verbatim from B.9.1)
ADD: 8 new command-handler branches	~150 lines
REMOVE	~80 lines

Net target file size: ~1600 LOC (vs v5's 1284). Well within single-file scope.

The most error-prone item is B.4 (collision kernel using book-derived equilibrium). The most tedious item is C.3 (event detectors). The most subtle item is C.1 (recentre logic, especially making sure the absolute-price reports stay correct after a shift).

---

End of diff-spec. Ready for kernel author review and the answers to Section G verification before any binary is shipped to validation.