Scruff AI
482 lines
25 KiB
Markdown
482 lines
25 KiB
Markdown
# 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)
|
||
|
||
```diff
|
||
-#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 NX*NY*Q*sizeof(float) = 9 MB vs 36 MB for v5.
|
||
|
||
### B.2 Device-side parameters block (lines 54\u201355)
|
||
|
||
```diff
|
||
-__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.
|
||
|
||
```diff
|
||
-__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)`.
|
||
|
||
```diff
|
||
-__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
|
||
|
||
```diff
|
||
-__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`
|
||
|
||
```diff
|
||
-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)
|
||
|
||
```diff
|
||
-"\"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":<f>,"side":"buy"|"sell","size":<f>,"aggressor":<bool>}` | host converts price\u2192column, signs the size, optionally multiplies by TAKER_WEIGHT, calls inject_trade_kernel |
|
||
| `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:
|
||
|
||
```c
|
||
// 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)
|
||
|
||
```c
|
||
// 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
|
||
|
||
```c
|
||
// 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:
|
||
|
||
```c
|
||
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)
|
||
|
||
```c
|
||
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:
|
||
|
||
```diff
|
||
- 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.
|