docs(per-horizon-cfc): kickoff — spec + plan + historical record
Spec: docs/superpowers/specs/2026-05-21-per-horizon-cfc-inference-design.md Plan: docs/superpowers/plans/2026-05-21-per-horizon-cfc-inference-plan.md Spec went through 2 critical-review passes (32 total findings, all resolved). Bucketing source: CfC.tau (per-channel, trained, log-uniform init at HIDDEN_DIM=128). Atomic refactor reverts MTER scaffolding. 5 ISV-driven controllers. All-on-device transition (no bulk DtoH). Single fused per-branch kernel. Compact ragged heads_w_skip. Validated via 3 Argo smokes (training stability, CRT.diag inference differentiation, fxt-backtest end-to-end). Also committing historical record: superseded MTER spec/plan, intervention A plan, ISV λ controller spec/plan, GPU log ring spec/plan. Per spec §5.3 these stay as audit trail; not in build path. Plan has 18 tasks, full TDD with bite-sized steps, ready for subagent-driven-development execution. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,389 @@
|
||||
# CRT.train ISV-Driven λ Controller — Implementation Plan
|
||||
|
||||
> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development. Each task is a checkbox list of steps.
|
||||
|
||||
**Goal:** Replace static `SMOOTHNESS_LAMBDA_RATIO` with a signal-driven λ controller anchored on observed h30 jitter, with permanent floor.
|
||||
|
||||
**Architecture:** New `smoothness_lambda_controller.cu` kernel reads `raw_per_h` from `output_smoothness` (already emitted), maintains a Wiener-α-floor-0.5 EMA of jitter per horizon, derives per-horizon target as `jitter_ema[h30] × (K_h30 / K_h)`, and emits new `λ[h]` for next step's `output_smoothness` launch. λ floor = 1e-4 (permanent floor pattern). Launch order: ... → BCE → output_smoothness → smoothness_lambda_controller → backward K-loop.
|
||||
|
||||
**Tech Stack:** Same as the predecessor (Rust 1.85+, cudarc 0.19, CUDA 12.4, pre-compiled cubin).
|
||||
|
||||
**Predecessor commits (already shipped):** `bf5ecd109..70ecfc0fe` (the 10-commit static-λ chain).
|
||||
|
||||
---
|
||||
|
||||
## File structure (locked decisions)
|
||||
|
||||
| File | Action |
|
||||
|------|--------|
|
||||
| `crates/ml-alpha/cuda/smoothness_lambda_controller.cu` | CREATE |
|
||||
| `crates/ml-alpha/build.rs` | Add `"smoothness_lambda_controller"` to `KERNELS`, bump v12→v13 |
|
||||
| `crates/ml-alpha/src/heads.rs` | DELETE `SMOOTHNESS_LAMBDA_RATIO` |
|
||||
| `crates/ml-alpha/src/trainer/perception.rs` | Add `jitter_ema_d`, `jitter_first_obs_d`, controller fn handle + module Arc; change λ init from `base × ratio[h]` to `[LAMBDA_FLOOR; 5]`; launch controller after `output_smoothness` in `step_batched`; remove `SMOOTHNESS_LAMBDA_RATIO` import |
|
||||
| `crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs` | CREATE — 4 invariant tests |
|
||||
|
||||
---
|
||||
|
||||
## Task 1: Create kernel `smoothness_lambda_controller.cu`
|
||||
|
||||
- [ ] **1.1** Write `crates/ml-alpha/cuda/smoothness_lambda_controller.cu`:
|
||||
|
||||
```cuda
|
||||
// smoothness_lambda_controller.cu — ISV-driven per-horizon λ for the
|
||||
// output_smoothness regularizer.
|
||||
//
|
||||
// Reads `raw_per_h[5]` (emitted by output_smoothness_loss_and_grad),
|
||||
// maintains a per-horizon Wiener-α-floor EMA of observed jitter,
|
||||
// derives per-horizon target by anchoring on observed h30 jitter
|
||||
// scaled by HORIZONS[0]/HORIZONS[h], and emits next-step λ[h] with a
|
||||
// permanent floor.
|
||||
//
|
||||
// Per `pearl_controller_anchors_isv_driven`: target is signal-derived,
|
||||
// not a constant. Per `pearl_first_observation_bootstrap`: first
|
||||
// observation replaces EMA directly. Per
|
||||
// `pearl_blend_formulas_must_have_permanent_floor`: λ has a permanent
|
||||
// floor so the controller never fully self-disables. Per
|
||||
// `pearl_wiener_alpha_floor_for_nonstationary`: α floored at 0.5 since
|
||||
// the controller's target drifts as the policy co-adapts.
|
||||
//
|
||||
// Per `feedback_no_atomicadd`: 5 threads, single block, single writer
|
||||
// per (h) slot; no atomics. Per `pearl_no_host_branches_in_captured_graph`:
|
||||
// no host branching; thread-id gating only.
|
||||
|
||||
#define SLC_N_HORIZONS 5
|
||||
#define SLC_LAMBDA_FLOOR 1.0e-4f
|
||||
#define SLC_TARGET_EPS 1.0e-9f
|
||||
#define SLC_ALPHA_FLOOR 0.5f
|
||||
|
||||
// HORIZONS = {30, 100, 300, 1000, 6000}.
|
||||
// Ratios HORIZONS[0]/HORIZONS[h] = {1, 0.3, 0.1, 0.03, 0.005}.
|
||||
// Constant array known at compile time.
|
||||
__device__ __constant__ float TARGET_K_RATIO[SLC_N_HORIZONS] = {
|
||||
1.0f,
|
||||
30.0f / 100.0f,
|
||||
30.0f / 300.0f,
|
||||
30.0f / 1000.0f,
|
||||
30.0f / 6000.0f,
|
||||
};
|
||||
|
||||
extern "C" __global__ void smoothness_lambda_controller(
|
||||
const float* __restrict__ raw_per_h, // [5] emitted by output_smoothness
|
||||
float* __restrict__ jitter_ema, // [5] in/out — EMA state
|
||||
int* __restrict__ first_obs, // [1] in/out — sentinel
|
||||
float base_lambda, // scalar — amplitude knob
|
||||
float* __restrict__ lambda_out // [5] output — λ for next step
|
||||
) {
|
||||
const int h = threadIdx.x;
|
||||
if (h >= SLC_N_HORIZONS) return;
|
||||
|
||||
__shared__ float s_jitter_after[SLC_N_HORIZONS];
|
||||
|
||||
// Pass 1: EMA update with sentinel bootstrap.
|
||||
const float raw_h = raw_per_h[h];
|
||||
const int sentinel = first_obs[0];
|
||||
float jitter_h;
|
||||
if (sentinel == 0) {
|
||||
jitter_h = raw_h;
|
||||
} else {
|
||||
jitter_h = (1.0f - SLC_ALPHA_FLOOR) * jitter_ema[h] + SLC_ALPHA_FLOOR * raw_h;
|
||||
}
|
||||
jitter_ema[h] = jitter_h;
|
||||
s_jitter_after[h] = jitter_h;
|
||||
|
||||
// Single-writer of sentinel — thread h=0 only.
|
||||
if (h == 0 && sentinel == 0) {
|
||||
first_obs[0] = 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// Pass 2: derive target and update λ.
|
||||
// target[h] = jitter_ema[0] * TARGET_K_RATIO[h]
|
||||
// = jitter_ema[0] for h=0 (self-target)
|
||||
// < jitter_ema[0] for h>0
|
||||
const float jitter_h0 = s_jitter_after[0];
|
||||
const float target_h = jitter_h0 * TARGET_K_RATIO[h];
|
||||
|
||||
// Excess controller: ratio - 1, clamped at zero (only push UP).
|
||||
// When observed > target: excess > 0 → λ grows
|
||||
// When observed ≤ target: excess = 0 → λ relaxes toward base_lambda × 1 = base_lambda
|
||||
// Floor: λ ≥ LAMBDA_FLOOR.
|
||||
const float safe_target = fmaxf(target_h, SLC_TARGET_EPS);
|
||||
const float excess_ratio = fmaxf(0.0f, jitter_h / safe_target - 1.0f);
|
||||
const float lambda_new = base_lambda * (1.0f + excess_ratio);
|
||||
lambda_out[h] = fmaxf(SLC_LAMBDA_FLOOR, lambda_new);
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **1.2** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/cuda/smoothness_lambda_controller.cu
|
||||
git commit -m "feat(crt-train): add ISV-driven smoothness_lambda_controller kernel"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 2: Register kernel in `build.rs`
|
||||
|
||||
- [ ] **2.1** Edit `crates/ml-alpha/build.rs`. In the `KERNELS` array, after `"output_smoothness"`, add `"smoothness_lambda_controller"` with comment. Bump cache-bust v12→v13:
|
||||
```rust
|
||||
"output_smoothness", // CRT.train: per-horizon adjacent-position prob-jitter penalty
|
||||
"smoothness_lambda_controller", // CRT.train: ISV-driven λ controller anchored on h30 jitter
|
||||
];
|
||||
```
|
||||
And replace the v12 comment with `// Cache bust v13 (2026-05-21): smoothness_lambda_controller.cu added — ISV-driven λ.`
|
||||
|
||||
- [ ] **2.2** Verify:
|
||||
```bash
|
||||
SQLX_OFFLINE=true CARGO_FEATURE_CUDA=1 cargo build -p ml-alpha 2>&1 | grep -E "smoothness_lambda_controller|^error" | head
|
||||
```
|
||||
Expected: `compiled cuda/smoothness_lambda_controller.cu -> .../smoothness_lambda_controller.cubin`. No errors.
|
||||
|
||||
- [ ] **2.3** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/build.rs
|
||||
git commit -m "build(crt-train): register smoothness_lambda_controller.cu in build.rs"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 3: Delete `SMOOTHNESS_LAMBDA_RATIO` from `heads.rs`
|
||||
|
||||
- [ ] **3.1** Edit `crates/ml-alpha/src/heads.rs`. DELETE the entire `pub const SMOOTHNESS_LAMBDA_RATIO: [f32; N_HORIZONS] = [ ... ];` declaration (added in Task 3 of the predecessor plan, around lines 27-39). Also delete its doc-comment block.
|
||||
|
||||
- [ ] **3.2** This will break `perception.rs` import — that's expected; Task 4 below replaces the import.
|
||||
|
||||
- [ ] **3.3** Do NOT commit yet — Task 4 ships in the same commit (atomic refactor per `feedback_no_partial_refactor`).
|
||||
|
||||
---
|
||||
|
||||
## Task 4: Wire controller into `PerceptionTrainer` (perception.rs)
|
||||
|
||||
This is the largest task. It includes: new fields, new cubin/handle, new λ init pattern, controller launch in `step_batched`, removal of the static-ratio code.
|
||||
|
||||
- [ ] **4.1** In `crates/ml-alpha/src/trainer/perception.rs`:
|
||||
|
||||
a. **Remove** the `SMOOTHNESS_LAMBDA_RATIO` from the `use crate::heads::{...}` line (the import added in predecessor Task 5).
|
||||
|
||||
b. **Add** the controller cubin constant adjacent to `SMOOTHNESS_CUBIN`:
|
||||
```rust
|
||||
const SMOOTHNESS_CONTROLLER_CUBIN: &[u8] = include_bytes!(
|
||||
concat!(env!("OUT_DIR"), "/smoothness_lambda_controller.cubin")
|
||||
);
|
||||
```
|
||||
|
||||
c. **Add** new struct fields adjacent to the existing `smoothness_*` fields (after `_smoothness_module`):
|
||||
```rust
|
||||
/// Per-horizon EMA of `raw_per_h` from output_smoothness. Updated
|
||||
/// by `smoothness_lambda_controller` each step. Drives the
|
||||
/// controller's target derivation.
|
||||
smoothness_jitter_ema_d: CudaSlice<f32>,
|
||||
/// First-observation sentinel (per pearl_first_observation_bootstrap).
|
||||
/// 0 → next step bootstraps EMA from raw_per_h; 1 → Wiener-α update.
|
||||
smoothness_jitter_first_obs_d: CudaSlice<i32>,
|
||||
/// Cached handle for `smoothness_lambda_controller`.
|
||||
smoothness_controller_fn: CudaFunction,
|
||||
_smoothness_controller_module: Arc<CudaModule>,
|
||||
```
|
||||
|
||||
d. **In the constructor**, after the existing smoothness module load, add the controller module load:
|
||||
```rust
|
||||
let smoothness_controller_module = ctx
|
||||
.load_cubin(SMOOTHNESS_CONTROLLER_CUBIN.to_vec())
|
||||
.context("load smoothness_lambda_controller cubin")?;
|
||||
let smoothness_controller_fn = smoothness_controller_module
|
||||
.load_function("smoothness_lambda_controller")
|
||||
.context("load smoothness_lambda_controller")?;
|
||||
```
|
||||
|
||||
e. **In the constructor**, **REPLACE** the existing λ init block (which computes `base × ratio[h]` per the predecessor Task 5) with a uniform-floor init:
|
||||
```rust
|
||||
// ── CRT.train: output-smoothness regularizer state ──
|
||||
// λ is now ISV-driven by smoothness_lambda_controller — initialised
|
||||
// to LAMBDA_FLOOR per horizon and updated each step inside the
|
||||
// captured graph. The base_lambda amplitude knob enters the
|
||||
// kernel as a scalar arg at launch time.
|
||||
const LAMBDA_FLOOR_INIT: f32 = 1.0e-4;
|
||||
let smoothness_lambda_host: Vec<f32> = vec![LAMBDA_FLOOR_INIT; N_HORIZONS];
|
||||
let smoothness_lambda_d = {
|
||||
let staging = unsafe { MappedF32Buffer::new(N_HORIZONS) }
|
||||
.map_err(|e| anyhow::anyhow!("smoothness λ staging: {e}"))?;
|
||||
staging.write_from_slice(&smoothness_lambda_host);
|
||||
let mut dst = stream.alloc_zeros::<f32>(N_HORIZONS)
|
||||
.context("smoothness_lambda_d alloc")?;
|
||||
let nbytes = N_HORIZONS * std::mem::size_of::<f32>();
|
||||
unsafe {
|
||||
let (dst_ptr, _g) = dst.device_ptr_mut(&stream);
|
||||
cudarc::driver::result::memcpy_dtod_async(
|
||||
dst_ptr, staging.dev_ptr, nbytes, stream.cu_stream(),
|
||||
).context("smoothness λ DtoD")?;
|
||||
}
|
||||
dst
|
||||
};
|
||||
let smoothness_loss_d = stream.alloc_zeros::<f32>(1)
|
||||
.context("smoothness_loss_d alloc")?;
|
||||
let smoothness_loss_host_d = unsafe { MappedF32Buffer::new(1) }
|
||||
.map_err(|e| anyhow::anyhow!("smoothness_loss_host_d: {e}"))?;
|
||||
let smoothness_loss_per_horizon_d = stream.alloc_zeros::<f32>(N_HORIZONS)
|
||||
.context("smoothness_loss_per_horizon_d alloc")?;
|
||||
let smoothness_loss_per_horizon_host_d = unsafe { MappedF32Buffer::new(N_HORIZONS) }
|
||||
.map_err(|e| anyhow::anyhow!("smoothness_loss_per_horizon_host_d: {e}"))?;
|
||||
// Controller state buffers.
|
||||
let smoothness_jitter_ema_d = stream.alloc_zeros::<f32>(N_HORIZONS)
|
||||
.context("smoothness_jitter_ema_d alloc")?;
|
||||
let smoothness_jitter_first_obs_d = stream.alloc_zeros::<i32>(1)
|
||||
.context("smoothness_jitter_first_obs_d alloc")?;
|
||||
```
|
||||
|
||||
f. **In the `Ok(Self { ... })` block**, add the new controller fields (alphabetically near the other `smoothness_*` fields):
|
||||
```rust
|
||||
smoothness_jitter_ema_d,
|
||||
smoothness_jitter_first_obs_d,
|
||||
smoothness_controller_fn,
|
||||
_smoothness_controller_module: smoothness_controller_module,
|
||||
```
|
||||
|
||||
g. **In `step_batched`** (inside the captured-graph region), find the smoothness launch added in predecessor Task 6 (around line 1991-2008). IMMEDIATELY AFTER the smoothness launch block's closing brace, insert the controller launch:
|
||||
```rust
|
||||
|
||||
// ── 5d. ISV-driven λ controller ──
|
||||
// Reads the raw per-horizon mean-sq-diff just emitted by
|
||||
// the output_smoothness kernel, updates jitter_ema, derives
|
||||
// per-horizon target anchored on h30, and produces λ for
|
||||
// the NEXT step's output_smoothness launch. Same captured
|
||||
// graph; one-step delay between observation and effect
|
||||
// (standard closed-loop pattern).
|
||||
let base_lambda = self.cfg.smoothness_base_lambda;
|
||||
let smooth_ctrl_cfg = LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (N_HORIZONS as u32, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
{
|
||||
let mut launch = self.stream.launch_builder(&self.smoothness_controller_fn);
|
||||
launch
|
||||
.arg(&self.smoothness_loss_per_horizon_d)
|
||||
.arg(&mut self.smoothness_jitter_ema_d)
|
||||
.arg(&mut self.smoothness_jitter_first_obs_d)
|
||||
.arg(&base_lambda)
|
||||
.arg(&mut self.smoothness_lambda_d);
|
||||
unsafe { launch.launch(smooth_ctrl_cfg).context("smoothness_lambda_controller launch")?; }
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **4.2** Verify:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets 2>&1 | grep -E "^error" | head
|
||||
```
|
||||
Expected: clean.
|
||||
|
||||
- [ ] **4.3** If you've removed `SMOOTHNESS_LAMBDA_RATIO` import successfully and the build still compiles, commit BOTH file changes atomically:
|
||||
```bash
|
||||
git add crates/ml-alpha/src/heads.rs crates/ml-alpha/src/trainer/perception.rs
|
||||
git commit -m "feat(crt-train): wire ISV-driven λ controller; remove SMOOTHNESS_LAMBDA_RATIO"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 5: Write controller invariant tests
|
||||
|
||||
- [ ] **5.1** Create `crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs`. Use the same upload/download helpers as `output_smoothness_grad_finite_diff.rs` (mapped-pinned + memcpy_dtod_async).
|
||||
|
||||
Tests to implement (4 #[test] functions):
|
||||
|
||||
1. `first_observation_bootstraps_ema_and_sets_floor`:
|
||||
- Init `first_obs=0`, `jitter_ema=[0;5]`, raw_per_h=[0.5, 0.3, 0.1, 0.05, 0.01], base_lambda=0.001
|
||||
- After kernel: `jitter_ema == raw_per_h`, `first_obs==1`, `λ[h] == LAMBDA_FLOOR` for all h (because excess_ratio = 0 on the bootstrap step: target = raw_per_h[0]*ratio[h], and jitter_ema[h] = raw_per_h[h] which is at-target when raw_per_h IS scaled by ratio; assert λ near LAMBDA_FLOOR or base_lambda, NOT runaway)
|
||||
- (Actually on first observation, jitter_ema[h] = raw_per_h[h] for all h. target[h] = raw_per_h[0]*ratio[h]. excess_ratio[h] = (raw_per_h[h] / (raw_per_h[0]*ratio[h])) - 1. If the test setup has raw_per_h[h] = raw_per_h[0]*ratio[h], excess_ratio = 0 → λ[h] = max(LAMBDA_FLOOR, base_lambda). Use a test setup that satisfies this for clean assertion.)
|
||||
|
||||
2. `steady_state_at_target_yields_base_lambda`:
|
||||
- Pre-set `first_obs=1`, `jitter_ema = [0.5, 0.15, 0.05, 0.015, 0.0025]` (= 0.5 × [1, 0.3, 0.1, 0.03, 0.005])
|
||||
- raw_per_h same as above
|
||||
- base_lambda = 0.01
|
||||
- Expected: jitter_ema stays at target (Wiener-α update at α=0.5: same value), excess_ratio = 0, λ[h] = max(LAMBDA_FLOOR, 0.01) = 0.01 for all h
|
||||
|
||||
3. `excess_at_h6000_lifts_lambda_proportionally`:
|
||||
- Pre-set `first_obs=1`, `jitter_ema = [0.5, 0.15, 0.05, 0.015, 0.025]` (h6000 at 10× target)
|
||||
- raw_per_h same
|
||||
- base_lambda = 0.01
|
||||
- After kernel: jitter_ema[h6000] EMA update → 0.5*0.025 + 0.5*0.025 = 0.025 (raw matches old EMA, no change)
|
||||
- Wait: if raw == old_ema, jitter_ema doesn't change. Use raw_per_h[h6000] = 0.025 (matches the pre-set ema)
|
||||
- target[h6000] = jitter_ema_new[0]*0.005 = 0.5*0.005 = 0.0025
|
||||
- excess_ratio = 0.025/0.0025 - 1 = 9.0
|
||||
- λ[h6000] = 0.01 * (1 + 9) = 0.10
|
||||
- Assert λ[h6000] is within 0.1 ± 0.005 (allowing for ema update on h0 if it shifts)
|
||||
|
||||
4. `lambda_floor_when_base_lambda_zero`:
|
||||
- Pre-set `first_obs=1`, `jitter_ema = [0.5; 5]`, raw matching
|
||||
- base_lambda = 0.0
|
||||
- Expected: λ[h] = LAMBDA_FLOOR = 1e-4 for all h (max(1e-4, 0 * (1 + excess)) = 1e-4)
|
||||
|
||||
- [ ] **5.2** Build:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test smoothness_lambda_controller_invariants --no-run 2>&1 | tail
|
||||
```
|
||||
Expected: clean.
|
||||
|
||||
If local CUDA available:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test smoothness_lambda_controller_invariants -- --nocapture 2>&1 | tail -30
|
||||
```
|
||||
Expected: 4 tests pass.
|
||||
|
||||
- [ ] **5.3** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs
|
||||
git commit -m "test(crt-train): invariant tests for smoothness_lambda_controller"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 6: Push + validate
|
||||
|
||||
- [ ] **6.1** Run full check + perception_overfit:
|
||||
```bash
|
||||
cd /home/jgrusewski/Work/foxhunt
|
||||
SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets 2>&1 | tail
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test perception_overfit -- --nocapture 2>&1 | tail -30 # if local CUDA
|
||||
```
|
||||
Expected: clean check; perception_overfit 9/9 (smoothness=0 base_lambda + LAMBDA_FLOOR = essentially-no-op).
|
||||
|
||||
- [ ] **6.2** Push:
|
||||
```bash
|
||||
git push origin ml-alpha-phase-a
|
||||
```
|
||||
|
||||
- [ ] **6.3** Submit validation retrain at `smoothness-base-lambda=0.001` (the amplitude knob; the controller scales adaptively from there):
|
||||
```bash
|
||||
argo submit -n foxhunt --from=wftmpl/alpha-perception \
|
||||
-p commit-sha=$(git rev-parse HEAD) \
|
||||
-p git-branch=ml-alpha-phase-a \
|
||||
-p gpu-pool=ci-training-l40s \
|
||||
-p smoothness-base-lambda=0.001 \
|
||||
-p n-train-seqs=8000 \
|
||||
-p n-val-seqs=1000 \
|
||||
-p epochs=5
|
||||
```
|
||||
|
||||
Watch logs for `final_smooth_loss_per_horizon` in the summary JSON. Expected: h30 ≈ h6000_target × 200, h6000 ≈ h30 / 200 (controller pulls h6000 to its 200× lower target).
|
||||
|
||||
---
|
||||
|
||||
## Self-Review
|
||||
|
||||
1. Spec extension §3 mapped to Task 1 (kernel), Task 4 (wiring). ✓
|
||||
2. Spec extension §3.4 removals mapped to Task 3 (heads.rs) + Task 4 (perception.rs constructor). ✓
|
||||
3. Spec extension §5 invariants mapped to Task 5 (4 tests). ✓
|
||||
4. No placeholders. No TODO markers. ✓
|
||||
5. Atomic refactor: Task 4 ships heads.rs deletion + perception.rs wiring in ONE commit (per `feedback_no_partial_refactor`). ✓
|
||||
6. Memory rules cited:
|
||||
- `feedback_no_atomicadd` ✓ (5 threads, single writer per slot)
|
||||
- `feedback_no_htod_htoh_only_mapped_pinned` ✓ (λ init still uses MappedF32Buffer)
|
||||
- `feedback_no_nvrtc` ✓ (build.rs registration)
|
||||
- `pearl_no_host_branches_in_captured_graph` ✓ (controller launch has no host branches)
|
||||
- `pearl_first_observation_bootstrap` ✓ (sentinel pattern in kernel)
|
||||
- `pearl_blend_formulas_must_have_permanent_floor` ✓ (LAMBDA_FLOOR is permanent)
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` ✓ (α=0.5 floor)
|
||||
- `pearl_controller_anchors_isv_driven` ✓ (target anchored on observed h30 jitter)
|
||||
- `feedback_no_partial_refactor` ✓ (atomic delete+wire in Task 4)
|
||||
|
||||
Ready to execute.
|
||||
795
docs/superpowers/plans/2026-05-21-gpu-log-ring-plan.md
Normal file
795
docs/superpowers/plans/2026-05-21-gpu-log-ring-plan.md
Normal file
@@ -0,0 +1,795 @@
|
||||
# GPU Log Ring — Implementation Plan
|
||||
|
||||
> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development. Each task is checkbox-tracked.
|
||||
|
||||
**Goal:** Unified high-performance in-kernel diagnostic logging for ml-alpha CUDA kernels via mapped-pinned ring buffer with deterministic slot reservation (no atomicAdd), header-last commit, and async host drain.
|
||||
|
||||
**Predecessor spec:** `docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md`
|
||||
|
||||
**Architecture in one paragraph.** A 32768-slot × 64 B mapped-pinned ring buffer holds `LogRecord` entries. A 1-thread tick kernel runs first in each captured-graph replay, incrementing a device `step_counter`. Every logging-capable kernel takes two new device-pointer args (`LogRing*`, `const int* step_counter`) and emits records via a `log_record()` device helper that derives its slot from `(step * N_RPS + kid * N_RT + rt) & MASK`, single-writer per slot, header-last commit. CPU drainer is a tokio task that volatile-reads the ring at 500 ms cadence, dispatching to per-(kernel_id, record_type) decoders that emit structured `tracing::info!` events. All gated by a `cuda-diag-log` Cargo feature.
|
||||
|
||||
**Tech Stack:** Rust 1.85+, cudarc 0.19, CUDA 12.4 (sm_89 L40S, sm_80 RTX 3050 Ti local), `tokio`, `tracing`, pre-compiled cubins.
|
||||
|
||||
---
|
||||
|
||||
## File structure (locked-in decisions)
|
||||
|
||||
| File | Action |
|
||||
|------|--------|
|
||||
| `crates/ml-alpha/cuda/gpu_log_ids.h` | NEW — kernel_id + record_type macros (single source of truth) |
|
||||
| `crates/ml-alpha/cuda/gpu_log_ring.cu` | NEW — `LogRecord`, `LogRing`, `log_record()` device helper, `gpu_log_tick` kernel |
|
||||
| `crates/ml-alpha/build.rs` | Add `"gpu_log_ring"` to KERNELS; bump cache-bust v13 → v14 |
|
||||
| `crates/ml-alpha/Cargo.toml` | Add feature flag `cuda-diag-log = []` |
|
||||
| `crates/ml-alpha/src/pinned_mem.rs` | Add generic `MappedRecordBuffer<T>`; `MappedF32Buffer` becomes alias `MappedRecordBuffer<f32>` |
|
||||
| `crates/ml-alpha/src/gpu_log.rs` | NEW — ring allocator, tokio drain task, decoder registry |
|
||||
| `crates/ml-alpha/src/lib.rs` | `#[cfg(feature = "cuda-diag-log")] pub mod gpu_log;` |
|
||||
| `crates/ml-alpha/tests/gpu_log_ring_invariants.rs` | NEW — 6 unit tests + 1 integration test |
|
||||
| `crates/ml-alpha/cuda/smoothness_lambda_controller.cu` | Add `g_log_ring`, `g_step_counter` args; emit RT_INPUT/STATE/OUTPUT |
|
||||
| `crates/ml-alpha/src/trainer/perception.rs` | Allocate ring + step counter at construction (feature-gated); wire pointers into smoothness controller launch |
|
||||
|
||||
---
|
||||
|
||||
## Task 1: Central IDs header
|
||||
|
||||
- [ ] **1.1** Create `crates/ml-alpha/cuda/gpu_log_ids.h`:
|
||||
|
||||
```c
|
||||
// gpu_log_ids.h — single source of truth for log-ring kernel/record IDs.
|
||||
//
|
||||
// Both kernel-side (#include from .cu files) and Rust-side decoder reference
|
||||
// these symbols. Drift between them is checked by `build.rs` (Task 2).
|
||||
//
|
||||
// kernel_id (uint8) — which kernel emitted the record.
|
||||
// record_type (uint8) — sub-type within the kernel (input/state/output/debug).
|
||||
// payload layout — defined per (kernel_id, record_type) in decoder.
|
||||
|
||||
#ifndef GPU_LOG_IDS_H
|
||||
#define GPU_LOG_IDS_H
|
||||
|
||||
#define GPU_LOG_N_KERNELS 8
|
||||
#define GPU_LOG_N_RT_PER_KERNEL 4
|
||||
#define GPU_LOG_N_RECORDS_PER_STEP (GPU_LOG_N_KERNELS * GPU_LOG_N_RT_PER_KERNEL)
|
||||
#define GPU_LOG_N_SLOTS 32768
|
||||
#define GPU_LOG_RECORD_BYTES 64
|
||||
#define GPU_LOG_PAYLOAD_F32 12
|
||||
#define GPU_LOG_MAGIC 0xF0F0A710u
|
||||
|
||||
// Kernel IDs (uint8). Add new kernels at the bottom; never reorder.
|
||||
#define KID_SMOOTHNESS_CONTROLLER 0
|
||||
#define KID_OUTPUT_SMOOTHNESS 1
|
||||
#define KID_BCE_MULTI_HORIZON 2
|
||||
#define KID_HORIZON_LAMBDA 3
|
||||
#define KID_HEADS_GRN_FWD 4
|
||||
#define KID_HEADS_GRN_BWD 5
|
||||
#define KID_CFC_STEP 6
|
||||
#define KID_RESERVED_7 7
|
||||
|
||||
// Record types within a kernel (uint8). Reuse RT_INPUT/STATE/OUTPUT/DEBUG
|
||||
// across kernels — payload layout differs per (kid, rt) pair but the four
|
||||
// roles are conventional.
|
||||
#define RT_INPUT 0
|
||||
#define RT_STATE 1
|
||||
#define RT_OUTPUT 2
|
||||
#define RT_DEBUG 3
|
||||
|
||||
#endif // GPU_LOG_IDS_H
|
||||
```
|
||||
|
||||
- [ ] **1.2** Commit:
|
||||
```bash
|
||||
cd /home/jgrusewski/Work/foxhunt
|
||||
git add crates/ml-alpha/cuda/gpu_log_ids.h
|
||||
git commit -m "feat(gpu-log): add gpu_log_ids.h — single source of truth for ring IDs"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 2: Device-side ring + tick kernel
|
||||
|
||||
- [ ] **2.1** Create `crates/ml-alpha/cuda/gpu_log_ring.cu`:
|
||||
|
||||
```c
|
||||
// gpu_log_ring.cu — unified GPU diagnostic log ring (device-side).
|
||||
//
|
||||
// Provides `log_record()` (a device __forceinline__ helper) and the
|
||||
// `gpu_log_tick` kernel that increments the step counter once per
|
||||
// captured-graph step.
|
||||
//
|
||||
// Constraints:
|
||||
// - No atomicAdd (slot reservation by deterministic step/kid/rt math).
|
||||
// - No host branches; thread-id gating only.
|
||||
// - Header-last commit (magic field written after payload + other header).
|
||||
// - Mapped-pinned ring; host reads via volatile, no DtoH memcpy.
|
||||
//
|
||||
// Per `feedback_nvidia_grade_perf_for_kernels`: single-thread writer for
|
||||
// the slot's 64 B cacheline, one coalesced store.
|
||||
|
||||
#include "gpu_log_ids.h"
|
||||
|
||||
#include <cstdint>
|
||||
|
||||
struct LogHeader {
|
||||
uint32_t magic;
|
||||
uint32_t step;
|
||||
uint8_t kernel_id;
|
||||
uint8_t record_type;
|
||||
uint8_t payload_words;
|
||||
uint8_t reserved;
|
||||
uint32_t _padding;
|
||||
};
|
||||
|
||||
struct LogRecord {
|
||||
LogHeader header;
|
||||
float payload[GPU_LOG_PAYLOAD_F32];
|
||||
};
|
||||
|
||||
struct LogRing {
|
||||
LogRecord records[GPU_LOG_N_SLOTS];
|
||||
};
|
||||
|
||||
// Tick kernel — 1 thread, 1 block. Increments step counter; runs FIRST in
|
||||
// every captured-graph replay so all subsequent kernels see the new step.
|
||||
extern "C" __global__ void gpu_log_tick(int* step_counter) {
|
||||
if (threadIdx.x == 0 && blockIdx.x == 0) {
|
||||
step_counter[0] = step_counter[0] + 1;
|
||||
}
|
||||
}
|
||||
|
||||
// log_record — emit a diagnostic record from a kernel.
|
||||
//
|
||||
// Single-writer per (step, kernel_id, record_type) slot. Thread (0,0) of
|
||||
// block 0 writes; all other threads no-op. Header-last commit ensures
|
||||
// torn-record reads on the host side are detectable via magic mismatch.
|
||||
//
|
||||
// Caller responsibility:
|
||||
// - g_log_ring == nullptr → no-op (logging disabled).
|
||||
// - g_step_counter is the device step counter (incremented by gpu_log_tick).
|
||||
// - payload_words ≤ GPU_LOG_PAYLOAD_F32 (12); kernel must respect this.
|
||||
__device__ __forceinline__ void log_record(
|
||||
LogRing* g_log_ring,
|
||||
const int* g_step_counter,
|
||||
uint8_t kernel_id,
|
||||
uint8_t record_type,
|
||||
const float* payload,
|
||||
int payload_words
|
||||
) {
|
||||
if (g_log_ring == nullptr) return;
|
||||
if (blockIdx.x != 0 || threadIdx.x != 0) return;
|
||||
|
||||
const int step = g_step_counter[0];
|
||||
const int slot = (step * GPU_LOG_N_RECORDS_PER_STEP
|
||||
+ (int)kernel_id * GPU_LOG_N_RT_PER_KERNEL
|
||||
+ (int)record_type) & (GPU_LOG_N_SLOTS - 1);
|
||||
LogRecord* rec = &g_log_ring->records[slot];
|
||||
|
||||
const int n = payload_words < GPU_LOG_PAYLOAD_F32 ? payload_words : GPU_LOG_PAYLOAD_F32;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < GPU_LOG_PAYLOAD_F32; ++i) {
|
||||
rec->payload[i] = (i < n) ? payload[i] : 0.0f;
|
||||
}
|
||||
rec->header.step = (uint32_t)step;
|
||||
rec->header.kernel_id = kernel_id;
|
||||
rec->header.record_type = record_type;
|
||||
rec->header.payload_words = (uint8_t)n;
|
||||
rec->header.reserved = 0;
|
||||
rec->header._padding = 0;
|
||||
|
||||
__threadfence(); // make payload + header fields visible …
|
||||
rec->header.magic = GPU_LOG_MAGIC; // … BEFORE the commit sentinel.
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **2.2** Register the kernel in `crates/ml-alpha/build.rs`. After `"smoothness_lambda_controller"` in the `KERNELS` array, add:
|
||||
```rust
|
||||
"smoothness_lambda_controller", // CRT.train: ISV-driven λ controller anchored on h30 jitter
|
||||
"gpu_log_ring", // GPU diagnostic log ring — tick kernel + log_record helper
|
||||
];
|
||||
```
|
||||
Also bump the cache-bust comment from v13 → v14:
|
||||
```rust
|
||||
// Cache bust v14 (2026-05-21): gpu_log_ring.cu added — unified in-kernel diagnostic logging.
|
||||
```
|
||||
|
||||
- [ ] **2.3** Verify:
|
||||
```bash
|
||||
SQLX_OFFLINE=true CARGO_FEATURE_CUDA=1 cargo build -p ml-alpha 2>&1 | grep -E "gpu_log|^error" | head
|
||||
```
|
||||
Expected: `compiled cuda/gpu_log_ring.cu -> .../gpu_log_ring.cubin`. No errors.
|
||||
|
||||
- [ ] **2.4** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/cuda/gpu_log_ring.cu crates/ml-alpha/build.rs
|
||||
git commit -m "feat(gpu-log): add gpu_log_ring.cu device helpers + tick kernel"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 3: Cargo feature flag
|
||||
|
||||
- [ ] **3.1** In `crates/ml-alpha/Cargo.toml`, locate the `[features]` section (or add one if missing — grep first). Add the feature:
|
||||
```toml
|
||||
[features]
|
||||
default = []
|
||||
cuda = []
|
||||
cuda-diag-log = []
|
||||
```
|
||||
(Adjust to merge with existing features. If `cuda` already exists with content, preserve it; only ADD `cuda-diag-log`.)
|
||||
|
||||
- [ ] **3.2** Verify:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo build -p ml-alpha --features cuda-diag-log 2>&1 | tail -5
|
||||
```
|
||||
Expected: builds cleanly (feature flag declared, not yet referenced anywhere).
|
||||
|
||||
- [ ] **3.3** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/Cargo.toml
|
||||
git commit -m "feat(gpu-log): add cuda-diag-log Cargo feature flag"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 4: Generic `MappedRecordBuffer<T>` in pinned_mem.rs
|
||||
|
||||
- [ ] **4.1** In `crates/ml-alpha/src/pinned_mem.rs`, locate the existing `MappedF32Buffer` impl. Extract its inner logic into a generic `MappedRecordBuffer<T: Copy>` and keep `MappedF32Buffer` as a type alias for backward compat. Concrete steps:
|
||||
|
||||
a. Add the generic struct + impl above the existing F32 type:
|
||||
```rust
|
||||
/// Mapped-pinned host buffer of `T` records. Allocated via cudaHostAlloc
|
||||
/// with the Mapped flag; both host and device see the same physical
|
||||
/// memory. Device pointer obtained via cudaHostGetDevicePointer.
|
||||
/// Per `feedback_no_htod_htoh_only_mapped_pinned`: the ONLY legitimate
|
||||
/// channel for cross-direction transfer in the captured-graph hot path.
|
||||
pub struct MappedRecordBuffer<T: Copy> {
|
||||
pub host_ptr: *mut T,
|
||||
pub dev_ptr: u64,
|
||||
pub len: usize,
|
||||
}
|
||||
|
||||
impl<T: Copy> MappedRecordBuffer<T> {
|
||||
pub unsafe fn new(n: usize) -> Result<Self, String> {
|
||||
// … existing allocation logic from MappedF32Buffer::new, generalised …
|
||||
}
|
||||
pub fn write_record(&self, idx: usize, val: T) {
|
||||
assert!(idx < self.len);
|
||||
unsafe { std::ptr::write_volatile(self.host_ptr.add(idx), val); }
|
||||
}
|
||||
pub fn read_record(&self, idx: usize) -> T {
|
||||
assert!(idx < self.len);
|
||||
unsafe { std::ptr::read_volatile(self.host_ptr.add(idx)) }
|
||||
}
|
||||
pub fn read_all(&self) -> Vec<T> {
|
||||
(0..self.len).map(|i| self.read_record(i)).collect()
|
||||
}
|
||||
}
|
||||
|
||||
impl<T: Copy> Drop for MappedRecordBuffer<T> {
|
||||
fn drop(&mut self) {
|
||||
// … existing cudaFreeHost logic …
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
b. Make `MappedF32Buffer` a type alias:
|
||||
```rust
|
||||
pub type MappedF32Buffer = MappedRecordBuffer<f32>;
|
||||
```
|
||||
|
||||
c. Preserve any F32-specific helpers (`write_from_slice(&[f32])`, etc.) by adding an `impl MappedRecordBuffer<f32>` block.
|
||||
|
||||
- [ ] **4.2** Verify nothing else broke:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets 2>&1 | grep -E "^error" | head
|
||||
```
|
||||
Expected: no errors. Existing callers of `MappedF32Buffer::new()` still work via the type alias.
|
||||
|
||||
- [ ] **4.3** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/src/pinned_mem.rs
|
||||
git commit -m "refactor(gpu-log): generic MappedRecordBuffer<T>; MappedF32Buffer as alias"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 5: Rust-side log ring + drain task (`gpu_log.rs`)
|
||||
|
||||
- [ ] **5.1** Create `crates/ml-alpha/src/gpu_log.rs`:
|
||||
|
||||
```rust
|
||||
//! GPU diagnostic log ring — host-side allocator, drainer, decoder registry.
|
||||
//!
|
||||
//! See `docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md`.
|
||||
//!
|
||||
//! Behaviour: when the `cuda-diag-log` feature is enabled, the trainer
|
||||
//! allocates a mapped-pinned ring buffer + a 1-element device step counter,
|
||||
//! passes their pointers into logging-capable kernels, and starts a tokio
|
||||
//! drain task that polls the ring at 500 ms cadence and emits structured
|
||||
//! tracing events.
|
||||
|
||||
use std::sync::Arc;
|
||||
use std::sync::atomic::{AtomicBool, Ordering};
|
||||
use std::time::Duration;
|
||||
use anyhow::{Context, Result};
|
||||
use cudarc::driver::{CudaSlice, CudaStream};
|
||||
use tokio::task::JoinHandle;
|
||||
use tracing::{info, warn, debug};
|
||||
|
||||
use crate::pinned_mem::MappedRecordBuffer;
|
||||
|
||||
pub const N_KERNELS: usize = 8;
|
||||
pub const N_RT_PER_KERNEL: usize = 4;
|
||||
pub const N_RECORDS_PER_STEP: usize = N_KERNELS * N_RT_PER_KERNEL;
|
||||
pub const N_SLOTS: usize = 32768;
|
||||
pub const PAYLOAD_F32: usize = 12;
|
||||
pub const MAGIC: u32 = 0xF0F0_A710;
|
||||
|
||||
#[repr(C)]
|
||||
#[derive(Clone, Copy, Debug)]
|
||||
pub struct LogHeader {
|
||||
pub magic: u32,
|
||||
pub step: u32,
|
||||
pub kernel_id: u8,
|
||||
pub record_type: u8,
|
||||
pub payload_words: u8,
|
||||
pub reserved: u8,
|
||||
pub _padding: u32,
|
||||
}
|
||||
|
||||
#[repr(C)]
|
||||
#[derive(Clone, Copy)]
|
||||
pub struct LogRecord {
|
||||
pub header: LogHeader,
|
||||
pub payload: [f32; PAYLOAD_F32],
|
||||
}
|
||||
|
||||
impl Default for LogRecord {
|
||||
fn default() -> Self {
|
||||
Self {
|
||||
header: LogHeader {
|
||||
magic: 0, step: 0,
|
||||
kernel_id: 0, record_type: 0, payload_words: 0, reserved: 0,
|
||||
_padding: 0,
|
||||
},
|
||||
payload: [0.0; PAYLOAD_F32],
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
pub struct LogRing {
|
||||
pub buffer: MappedRecordBuffer<LogRecord>,
|
||||
}
|
||||
|
||||
impl LogRing {
|
||||
pub fn alloc() -> Result<Self> {
|
||||
let buffer = unsafe { MappedRecordBuffer::<LogRecord>::new(N_SLOTS) }
|
||||
.map_err(|e| anyhow::anyhow!("log ring alloc: {e}"))?;
|
||||
// Zero-initialise so magic is 0 everywhere (no false-positive reads).
|
||||
let zero = LogRecord::default();
|
||||
for i in 0..N_SLOTS {
|
||||
buffer.write_record(i, zero);
|
||||
}
|
||||
Ok(Self { buffer })
|
||||
}
|
||||
|
||||
/// Device pointer to the ring (for kernel arg).
|
||||
pub fn dev_ptr(&self) -> u64 {
|
||||
self.buffer.dev_ptr
|
||||
}
|
||||
}
|
||||
|
||||
/// Allocate the step counter as a device buffer initialised to 0.
|
||||
pub fn alloc_step_counter(stream: &Arc<CudaStream>) -> Result<CudaSlice<i32>> {
|
||||
stream.alloc_zeros::<i32>(1).context("step_counter alloc")
|
||||
}
|
||||
|
||||
/// Type for a per-(kernel_id, record_type) decoder function.
|
||||
pub type DecoderFn = fn(step: u32, payload: &[f32]);
|
||||
|
||||
/// Per-(kernel_id, record_type) decoder registry. Default is the
|
||||
/// generic debug-print decoder. Specific decoders override entries.
|
||||
fn default_decoder(step: u32, payload: &[f32]) {
|
||||
debug!(target: "gpu_log", step, ?payload, "unknown record");
|
||||
}
|
||||
|
||||
fn smoothness_input(step: u32, payload: &[f32]) {
|
||||
if payload.len() >= 10 {
|
||||
info!(target: "gpu_log",
|
||||
step,
|
||||
raw_h30=payload[0], raw_h100=payload[1], raw_h300=payload[2],
|
||||
raw_h1000=payload[3], raw_h6000=payload[4],
|
||||
jitter_in_h30=payload[5], jitter_in_h100=payload[6],
|
||||
jitter_in_h300=payload[7], jitter_in_h1000=payload[8],
|
||||
jitter_in_h6000=payload[9],
|
||||
"smoothness_ctrl RT_INPUT");
|
||||
}
|
||||
}
|
||||
|
||||
fn smoothness_state(step: u32, payload: &[f32]) {
|
||||
if payload.len() >= 10 {
|
||||
info!(target: "gpu_log",
|
||||
step,
|
||||
ema_h30=payload[0], ema_h100=payload[1], ema_h300=payload[2],
|
||||
ema_h1000=payload[3], ema_h6000=payload[4],
|
||||
target_h30=payload[5], target_h100=payload[6], target_h300=payload[7],
|
||||
target_h1000=payload[8], target_h6000=payload[9],
|
||||
"smoothness_ctrl RT_STATE");
|
||||
}
|
||||
}
|
||||
|
||||
fn smoothness_output(step: u32, payload: &[f32]) {
|
||||
if payload.len() >= 10 {
|
||||
info!(target: "gpu_log",
|
||||
step,
|
||||
excess_h30=payload[0], excess_h100=payload[1], excess_h300=payload[2],
|
||||
excess_h1000=payload[3], excess_h6000=payload[4],
|
||||
lambda_h30=payload[5], lambda_h100=payload[6], lambda_h300=payload[7],
|
||||
lambda_h1000=payload[8], lambda_h6000=payload[9],
|
||||
"smoothness_ctrl RT_OUTPUT");
|
||||
}
|
||||
}
|
||||
|
||||
/// Static dispatch table indexed by [kernel_id][record_type].
|
||||
/// Add entries here as new kernels opt into the ring.
|
||||
fn decode(kernel_id: u8, record_type: u8, step: u32, payload: &[f32]) {
|
||||
match (kernel_id, record_type) {
|
||||
(0, 0) => smoothness_input(step, payload), // KID_SMOOTHNESS_CONTROLLER, RT_INPUT
|
||||
(0, 1) => smoothness_state(step, payload),
|
||||
(0, 2) => smoothness_output(step, payload),
|
||||
_ => default_decoder(step, payload),
|
||||
}
|
||||
}
|
||||
|
||||
/// Spawn the drain task. Polls the ring at 500 ms cadence; for each new
|
||||
/// step in [last_drained .. step_counter], walks N_RECORDS_PER_STEP slots,
|
||||
/// validates magic + step match, dispatches to decoder.
|
||||
/// Returns a JoinHandle + a stop signal.
|
||||
pub fn spawn_drain_task(
|
||||
ring: Arc<LogRing>,
|
||||
step_counter_host_shadow: Arc<MappedRecordBuffer<i32>>,
|
||||
stop: Arc<AtomicBool>,
|
||||
) -> JoinHandle<()> {
|
||||
tokio::spawn(async move {
|
||||
let mut last_drained_step: u32 = 0;
|
||||
let interval = Duration::from_millis(500);
|
||||
loop {
|
||||
if stop.load(Ordering::Relaxed) { break; }
|
||||
tokio::time::sleep(interval).await;
|
||||
|
||||
// Read device step counter via mapped-pinned shadow.
|
||||
let current_step = step_counter_host_shadow.read_record(0) as u32;
|
||||
if current_step <= last_drained_step { continue; }
|
||||
|
||||
let max_drain = (N_SLOTS / N_RECORDS_PER_STEP) as u32; // step window
|
||||
let oldest_safe_step = current_step.saturating_sub(max_drain);
|
||||
if last_drained_step < oldest_safe_step {
|
||||
warn!(target: "gpu_log",
|
||||
last_drained_step, current_step,
|
||||
skipped=oldest_safe_step.saturating_sub(last_drained_step),
|
||||
"drainer fell behind — records dropped");
|
||||
last_drained_step = oldest_safe_step;
|
||||
}
|
||||
|
||||
for step in (last_drained_step + 1)..=current_step {
|
||||
for kid in 0..N_KERNELS as u8 {
|
||||
for rt in 0..N_RT_PER_KERNEL as u8 {
|
||||
let slot = ((step as usize) * N_RECORDS_PER_STEP
|
||||
+ (kid as usize) * N_RT_PER_KERNEL
|
||||
+ rt as usize) & (N_SLOTS - 1);
|
||||
let rec = ring.buffer.read_record(slot);
|
||||
if rec.header.magic != MAGIC { continue; }
|
||||
if rec.header.step != step { continue; }
|
||||
let n = rec.header.payload_words as usize;
|
||||
let n = n.min(PAYLOAD_F32);
|
||||
decode(rec.header.kernel_id, rec.header.record_type,
|
||||
rec.header.step, &rec.payload[..n]);
|
||||
}
|
||||
}
|
||||
}
|
||||
last_drained_step = current_step;
|
||||
}
|
||||
})
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **5.2** Expose the module in `crates/ml-alpha/src/lib.rs`. Add (preserving existing `pub mod` style):
|
||||
```rust
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
pub mod gpu_log;
|
||||
```
|
||||
|
||||
- [ ] **5.3** Verify:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo build -p ml-alpha --features cuda-diag-log 2>&1 | tail -10
|
||||
SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets 2>&1 | grep -E "^error" | head
|
||||
```
|
||||
Expected: both clean.
|
||||
|
||||
- [ ] **5.4** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/src/gpu_log.rs crates/ml-alpha/src/lib.rs
|
||||
git commit -m "feat(gpu-log): Rust-side LogRing, drain task, decoder registry"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 6: Invariant tests
|
||||
|
||||
- [ ] **6.1** Create `crates/ml-alpha/tests/gpu_log_ring_invariants.rs`. The tests don't depend on a real kernel writing records — they exercise the Rust-side allocator and the host-visible round-trip by writing records FROM the host to the mapped-pinned buffer and verifying the layout matches what a kernel would write (`#[repr(C)]` ensures binary compatibility).
|
||||
|
||||
Six unit tests:
|
||||
|
||||
1. `ring_alloc_zeros_magic` — alloc ring, every slot's magic = 0.
|
||||
2. `host_write_round_trip` — host writes a record at a slot, reads it back, all fields match (asserts `#[repr(C)]` layout).
|
||||
3. `magic_validation_skips_torn` — write record with magic=0; drainer ignores. Write record with magic=MAGIC and step mismatch; drainer also ignores. Write record with magic=MAGIC and matching step; drainer reads.
|
||||
4. `ring_wrap_modulo_arithmetic` — for step values N_SLOTS/N_RECORDS_PER_STEP and beyond, slot indices wrap to start; verify no out-of-bounds.
|
||||
5. `step_gap_detection` — emit records at step 100 and step 200 with gap of 100 between; verify drainer detects the gap and emits warn.
|
||||
6. `decoder_dispatch_unknown_falls_through` — write a record with kernel_id=99 (unregistered); verify no panic, default decoder hit.
|
||||
|
||||
Test 7 (integration, gated by `--features cuda-diag-log` AND CUDA present):
|
||||
|
||||
7. `kernel_writes_visible_to_host` — launch the gpu_log_tick kernel + a tiny test kernel that calls log_record() with a known payload; drain; verify host sees the record.
|
||||
|
||||
Use `tokio::test` for any test exercising the drain task. Use `serial_test` if multiple tests share the ring (otherwise allocate fresh per test).
|
||||
|
||||
- [ ] **6.2** Build + run:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --features cuda-diag-log --test gpu_log_ring_invariants -- --nocapture 2>&1 | tail -30
|
||||
```
|
||||
Expected: 6 unit tests pass; test 7 passes if CUDA available, skips otherwise.
|
||||
|
||||
- [ ] **6.3** Commit:
|
||||
```bash
|
||||
git add crates/ml-alpha/tests/gpu_log_ring_invariants.rs
|
||||
git commit -m "test(gpu-log): invariant tests for ring allocator + drainer"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 7: Wire smoothness_lambda_controller as the first user
|
||||
|
||||
This is the ATOMIC refactor (per `feedback_no_partial_refactor`): the kernel signature changes; the trainer launch changes; the controller emits records. Single commit covers all three.
|
||||
|
||||
- [ ] **7.1** Modify `crates/ml-alpha/cuda/smoothness_lambda_controller.cu`:
|
||||
|
||||
a. Add `#include "gpu_log_ids.h"` and a forward declaration of `log_record()` (or factor `log_record` into a separate `gpu_log_helpers.cuh` header and include it). Decide once: header-only inclusion is simpler — extract `log_record` and the `LogRecord`/`LogRing`/`LogHeader` structs into `cuda/gpu_log_helpers.cuh`, then `gpu_log_ring.cu` includes the header to define the tick kernel; smoothness_lambda_controller.cu ALSO includes the header to use `log_record`.
|
||||
|
||||
b. Add two new args to the kernel signature:
|
||||
```c
|
||||
extern "C" __global__ void smoothness_lambda_controller(
|
||||
const float* __restrict__ raw_per_h,
|
||||
float* __restrict__ jitter_ema,
|
||||
int* __restrict__ first_obs,
|
||||
float base_lambda,
|
||||
float* __restrict__ lambda_out,
|
||||
LogRing* g_log_ring, // NEW (nullable)
|
||||
const int* g_step_counter // NEW
|
||||
)
|
||||
```
|
||||
|
||||
c. After the lambda update at end of kernel, emit three records (RT_INPUT, RT_STATE, RT_OUTPUT). Thread 0 of the single block reads the relevant arrays into stack variables for the payload, calls `log_record()` three times:
|
||||
```c
|
||||
if (h == 0) {
|
||||
// RT_INPUT: raw_per_h + jitter_ema_in (10 floats)
|
||||
float in_payload[10] = {
|
||||
raw_per_h[0], raw_per_h[1], raw_per_h[2], raw_per_h[3], raw_per_h[4],
|
||||
/* jitter_ema_in was the prior value — but we've already overwritten it.
|
||||
Capture BEFORE the EMA update via a temp register. */
|
||||
jitter_in_h30, jitter_in_h100, jitter_in_h300, jitter_in_h1000, jitter_in_h6000,
|
||||
};
|
||||
log_record(g_log_ring, g_step_counter, KID_SMOOTHNESS_CONTROLLER, RT_INPUT, in_payload, 10);
|
||||
|
||||
// RT_STATE: jitter_ema_out + target (10 floats)
|
||||
float state_payload[10] = {
|
||||
s_jitter_after[0], s_jitter_after[1], s_jitter_after[2], s_jitter_after[3], s_jitter_after[4],
|
||||
target_0, target_1, target_2, target_3, target_4,
|
||||
};
|
||||
log_record(g_log_ring, g_step_counter, KID_SMOOTHNESS_CONTROLLER, RT_STATE, state_payload, 10);
|
||||
|
||||
// RT_OUTPUT: excess_ratio + lambda_out (10 floats)
|
||||
float out_payload[10] = {
|
||||
excess_0, excess_1, excess_2, excess_3, excess_4,
|
||||
lambda_out[0], lambda_out[1], lambda_out[2], lambda_out[3], lambda_out[4],
|
||||
};
|
||||
log_record(g_log_ring, g_step_counter, KID_SMOOTHNESS_CONTROLLER, RT_OUTPUT, out_payload, 10);
|
||||
}
|
||||
```
|
||||
d. The kernel's per-thread excess_ratio + target are computed per-h; thread (h=0) needs to gather all 5. Use a shared array `__shared__ float s_target[5]; __shared__ float s_excess[5];` populated by each thread (per-(h) writes its own slot), barrier, then thread 0 reads.
|
||||
|
||||
- [ ] **7.2** Modify `crates/ml-alpha/src/trainer/perception.rs`:
|
||||
|
||||
a. Behind `#[cfg(feature = "cuda-diag-log")]`, add new fields to `PerceptionTrainer`:
|
||||
```rust
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
log_ring: std::sync::Arc<crate::gpu_log::LogRing>,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
log_step_counter_d: CudaSlice<i32>,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
log_step_counter_host_shadow: std::sync::Arc<crate::pinned_mem::MappedRecordBuffer<i32>>,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
gpu_log_tick_fn: CudaFunction,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
_gpu_log_module: std::sync::Arc<cudarc::driver::CudaModule>,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
log_drain_stop: std::sync::Arc<std::sync::atomic::AtomicBool>,
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
log_drain_handle: Option<tokio::task::JoinHandle<()>>,
|
||||
```
|
||||
|
||||
b. In the constructor (behind `#[cfg(feature = "cuda-diag-log")]`): allocate the ring, step counter, host shadow; load the `gpu_log_tick` cubin and function; spawn drain task.
|
||||
|
||||
c. In `step_batched`: BEFORE the existing kernel launches in the captured graph, launch `gpu_log_tick`:
|
||||
```rust
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
{
|
||||
let tick_cfg = LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (1, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
let mut launch = self.stream.launch_builder(&self.gpu_log_tick_fn);
|
||||
launch.arg(&mut self.log_step_counter_d);
|
||||
unsafe { launch.launch(tick_cfg).context("gpu_log_tick launch")?; }
|
||||
}
|
||||
```
|
||||
|
||||
d. In the smoothness_lambda_controller launch (the existing block from Task 4 of the predecessor plan), add two more args at the end:
|
||||
```rust
|
||||
launch
|
||||
.arg(&self.smoothness_loss_per_horizon_d)
|
||||
.arg(&mut self.smoothness_jitter_ema_d)
|
||||
.arg(&mut self.smoothness_jitter_first_obs_d)
|
||||
.arg(&base_lambda)
|
||||
.arg(&mut self.smoothness_lambda_d)
|
||||
// NEW (cfg-gated):
|
||||
.arg(&log_ring_dev_ptr) // u64 device pointer; null when feature off
|
||||
.arg(&log_step_counter_dev_ptr);
|
||||
```
|
||||
When feature is off, pass null pointers (the kernel branches on `g_log_ring == nullptr`). When feature is on, pass `self.log_ring.dev_ptr()` and the step counter pointer.
|
||||
|
||||
Since the launch args change between feature-on and feature-off builds, this is a `#[cfg]`-gated build of the launch block. The cleanest approach: keep the call site uniform, but conditionally compute the pointer values. Concrete:
|
||||
```rust
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
let log_ring_ptr: u64 = self.log_ring.dev_ptr();
|
||||
#[cfg(not(feature = "cuda-diag-log"))]
|
||||
let log_ring_ptr: u64 = 0;
|
||||
|
||||
#[cfg(feature = "cuda-diag-log")]
|
||||
let log_step_ptr: u64 = {
|
||||
let (p, _g) = self.log_step_counter_d.device_ptr(&self.stream);
|
||||
p
|
||||
};
|
||||
#[cfg(not(feature = "cuda-diag-log"))]
|
||||
let log_step_ptr: u64 = 0;
|
||||
```
|
||||
And the kernel signature ALWAYS has the two args (kernel always takes them; checks for null at top). So the kernel signature change is permanent; the values are conditional.
|
||||
|
||||
e. In drop / shutdown: set `log_drain_stop` and await the drain handle.
|
||||
|
||||
- [ ] **7.3** Verify both builds:
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo build -p ml-alpha 2>&1 | tail -5 # feature off
|
||||
SQLX_OFFLINE=true cargo build -p ml-alpha --features cuda-diag-log 2>&1 | tail -5 # feature on
|
||||
SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets 2>&1 | grep -E "^error" | head
|
||||
```
|
||||
Expected: all three clean.
|
||||
|
||||
Also run the existing controller invariant test (since the kernel signature changed, the test must also pass the two new args — likely null pointers for feature-off):
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test smoothness_lambda_controller_invariants -- --nocapture 2>&1 | tail
|
||||
```
|
||||
Expected: 4/4 pass. If the test breaks because the kernel signature changed, update the test in this same commit to pass null pointers (atomic refactor).
|
||||
|
||||
- [ ] **7.4** Commit (atomic):
|
||||
```bash
|
||||
git add crates/ml-alpha/cuda/smoothness_lambda_controller.cu \
|
||||
crates/ml-alpha/cuda/gpu_log_helpers.cuh \
|
||||
crates/ml-alpha/src/trainer/perception.rs \
|
||||
crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs
|
||||
git commit -m "feat(gpu-log): wire smoothness_lambda_controller as first ring producer"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 8: Push + smoke + retrain decision
|
||||
|
||||
- [ ] **8.1** Push:
|
||||
```bash
|
||||
git push origin ml-alpha-phase-a
|
||||
```
|
||||
|
||||
- [ ] **8.2** Local CUDA smoke (RTX 3050 Ti):
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --features cuda-diag-log --test gpu_log_ring_invariants -- --nocapture 2>&1 | tail -30
|
||||
FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml-alpha --features cuda-diag-log --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -30
|
||||
```
|
||||
Expected: ring tests pass; smoke runs without panic; tracing emits structured `gpu_log` records.
|
||||
|
||||
- [ ] **8.3** Decision gate based on CRT.diag (`lob-backtest-sweep-sqkst`) WIN-gate result:
|
||||
- **WIN passed at base_lambda=0.001**: the log ring is valuable infra but not blocking; leave the cuda-diag-log feature off by default in CI; revisit if a subsequent controller needs diagnostics.
|
||||
- **WIN failed at base_lambda=0.001**: enable `cuda-diag-log` for the next retrain at base_lambda=0.01-0.02 and use the per-step `lambda_h6000` trajectory to confirm mechanism #1 (amplitude underflow) vs mechanism #2 (shared trunk leakage).
|
||||
|
||||
Update the project memory file `project_crt_diag_findings.md` with the gate outcome and the diagnostic plan.
|
||||
|
||||
- [ ] **8.4** (Conditional) Submit Argo retrain at base_lambda=0.01 with diagnostic logging:
|
||||
```bash
|
||||
argo submit -n foxhunt --from=wftmpl/alpha-perception \
|
||||
-p commit-sha=$(git rev-parse HEAD) \
|
||||
-p git-branch=ml-alpha-phase-a \
|
||||
-p gpu-pool=ci-training-l40s \
|
||||
-p smoothness-base-lambda=0.01 \
|
||||
-p n-train-seqs=8000 -p n-val-seqs=1000 -p epochs=5
|
||||
# Note: the Argo template doesn't currently expose --features. To enable
|
||||
# cuda-diag-log on the cluster build, add a workflow parameter
|
||||
# `cargo-features` and thread it into `cargo build --release --features
|
||||
# "$CARGO_FEATURES" -p ml-alpha --example alpha_train` in the
|
||||
# ensure-binary step. Task 9 below covers that wiring if needed.
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Task 9: (Optional) Plumb feature flag through Argo template
|
||||
|
||||
Only execute if Task 8.3 decides we need diagnostics on the cluster (WIN failed branch).
|
||||
|
||||
- [ ] **9.1** In `infra/k8s/argo/alpha-perception-template.yaml`, add a workflow parameter:
|
||||
```yaml
|
||||
- name: cargo-features
|
||||
value: ""
|
||||
```
|
||||
|
||||
- [ ] **9.2** In the `ensure-binary` step's cargo build invocation:
|
||||
```yaml
|
||||
cargo build --release \
|
||||
${CARGO_FEATURES:+--features "$CARGO_FEATURES"} \
|
||||
-p ml-alpha --example alpha_train
|
||||
```
|
||||
Where `CARGO_FEATURES="{{workflow.parameters.cargo-features}}"` is set from the parameter.
|
||||
|
||||
- [ ] **9.3** Commit + push:
|
||||
```bash
|
||||
git add infra/k8s/argo/alpha-perception-template.yaml
|
||||
git commit -m "feat(gpu-log): plumb --features through alpha-perception template"
|
||||
git push origin ml-alpha-phase-a
|
||||
```
|
||||
|
||||
- [ ] **9.4** Submit the retrain with diagnostics enabled:
|
||||
```bash
|
||||
argo submit -n foxhunt --from=wftmpl/alpha-perception \
|
||||
-p commit-sha=$(git rev-parse HEAD) \
|
||||
-p git-branch=ml-alpha-phase-a \
|
||||
-p gpu-pool=ci-training-l40s \
|
||||
-p smoothness-base-lambda=0.01 \
|
||||
-p cargo-features=cuda-diag-log \
|
||||
-p n-train-seqs=8000 -p n-val-seqs=1000 -p epochs=5
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Self-Review
|
||||
|
||||
1. **Spec coverage:**
|
||||
- §3.1 (record format) → Task 1 (header) + Task 2 (struct in .cu) ✓
|
||||
- §3.2 (slot reservation) → Task 2 (`log_record`) ✓
|
||||
- §3.3 (header-last commit) → Task 2 ✓
|
||||
- §3.4 (mapped-pinned host visibility + drain) → Task 4 (MappedRecordBuffer) + Task 5 (drain task) ✓
|
||||
- §3.5 (kernel-author API) → Task 1 (header) + Task 7 (first user) ✓
|
||||
- §3.6 (Rust decoder) → Task 5 ✓
|
||||
- §3.7 (wiring lifecycle) → Task 7 ✓
|
||||
- §4 (memory rule compliance) → enforced throughout; verified per-task ✓
|
||||
- §5 (validation gates) → Task 6 (6 unit tests + 1 integration) ✓
|
||||
2. **No placeholders.** All snippets are complete code; no TODOs in shipping content. ✓
|
||||
3. **Type consistency.** `LogRecord`, `LogHeader`, `LogRing` named identically in Tasks 1-5. `MappedRecordBuffer<T>` named identically in Tasks 4-5. ✓
|
||||
4. **Atomic-refactor compliance.** Task 7 ships kernel sig change + trainer wiring + test update in ONE commit. ✓
|
||||
5. **Memory rule citations.**
|
||||
- `feedback_no_atomicadd` → Task 2 (slot reservation by deterministic math) ✓
|
||||
- `feedback_no_htod_htoh_only_mapped_pinned` → Tasks 4, 5 (mapped-pinned ring + step counter shadow) ✓
|
||||
- `feedback_no_nvrtc` → Task 2 (build.rs registration) ✓
|
||||
- `pearl_no_host_branches_in_captured_graph` → Task 7 (only device pointer args change between captures) ✓
|
||||
- `pearl_cudarc_disable_event_tracking_for_graph_capture` → Tasks 5, 7 (no host malloc inside captured graph; ring + step counter pre-allocated at construction) ✓
|
||||
- `feedback_no_stubs`, `feedback_no_todo_fixme`, `feedback_no_hiding` → enforced ✓
|
||||
- `feedback_no_partial_refactor` → Task 7 single commit ✓
|
||||
- `feedback_single_source_of_truth_no_duplicates` → Task 1 single header ✓
|
||||
- `feedback_wire_everything_up` → Tasks 7-8 wire allocate → tick → record → drain → tracing ✓
|
||||
- `pearl_build_rs_rerun_if_env_changed` → no new env reads ✓
|
||||
|
||||
Ready to execute when CRT.diag's WIN-gate decision lands.
|
||||
2611
docs/superpowers/plans/2026-05-21-per-horizon-cfc-inference-plan.md
Normal file
2611
docs/superpowers/plans/2026-05-21-per-horizon-cfc-inference-plan.md
Normal file
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,360 @@
|
||||
# CRT.train Intervention B — Multi-Timescale EMA Readout (Design v5)
|
||||
|
||||
**Date:** 2026-05-21 (v5 — addresses all concerns from v4 review)
|
||||
**Branch context:** `ml-alpha-phase-a` (HEAD `6d65423e3`)
|
||||
**Predecessor specs:**
|
||||
- `2026-05-20-crt-train-output-smoothness-design.md` (intervention A)
|
||||
- `2026-05-21-crt-train-isv-driven-lambda-controller.md` (controller refinement)
|
||||
**Predecessor pearl:** [pearl_training_smoothness_does_not_transfer_to_inference](memory/pearl_training_smoothness_does_not_transfer_to_inference.md)
|
||||
**Status:** v5 — pending review
|
||||
|
||||
---
|
||||
|
||||
## v5 changelog from v4
|
||||
|
||||
| Concern | v4 | v5 fix |
|
||||
|---|---|---|
|
||||
| **A** Inference path changes unspecified | Mentioned only in passing | **NEW §4: Inference path design** — `forward_step_into` MTER launch, per-instance h_view state lifecycle, LOB backtest harness integration |
|
||||
| **B** Multiple golden tests need regeneration, not just heads_bit_equiv | Only `heads_bit_equiv.rs` mentioned | §6.4 enumerates ALL affected golden tests: `heads_bit_equiv.rs`, `forward_step_golden.rs`, `perception_forward_golden.rs`, `forward_graph_capture.rs` |
|
||||
| **C** Val warmup bias mitigation undecided | Two options listed | **Decided**: post-warmup AUC measurement in val loop. AUC accumulator only counts samples after `events_since_val_start ≥ τ_max`. Cheaper than 5×ing val. |
|
||||
| **D** Smoothness slot orphaning in gpu_log_ids.h | Not addressed | **Decided**: reclaim slot 0 for `KID_MTER_READOUT`. `KID_SMOOTHNESS_CONTROLLER` removed entirely. Decoders renamed. Other kernel IDs shift down by 1. |
|
||||
| **E** Commit-1 smoke validation implicit | Not explicit | **§5.4 commit 1 explicitly requires** sequential-vs-random smoke before commit 2 proceeds. Embedded pre-validation. |
|
||||
| **F** Heads kernel LOC underestimated (250) | +250/-80 | **Bumped to +350-400/-100**. Honest. |
|
||||
| **G** Multi-fold trainer reuse silent-bug risk | Not addressed | **Added `reset_for_new_fold()` method** + explicit assert in trainer construction that fold transitions reset state. |
|
||||
| **H** 75-min wall estimate too high | "~75 min" | **Corrected to ~36 min**: 5 build + 18 train + 13 CRT.diag |
|
||||
| **Minor** Val cfc_h_state attn_pool bootstrap unclear | Not addressed | §4.5 explicit: val_cfc_h_state bootstrapped via attn_pool on val-start, just as a file-boundary bootstrap |
|
||||
|
||||
---
|
||||
|
||||
## 1. Motivation (unchanged)
|
||||
|
||||
Intervention A's smoothness regularizer differentiated horizons at training time (jitter ratio 0.40) but failed inference WIN gate by ~25× (mean_run_len 3.9 vs target ≥100). Per `pearl_training_smoothness_does_not_transfer_to_inference`: position-wise penalty on training `[K, B, 5]` shapes weights only relative to the training geometry; inference's event-by-event `forward_step_into` doesn't inherit the constraint.
|
||||
|
||||
## 2. Hypothesis
|
||||
|
||||
The trunk already amplifies short-horizon evidence into multi-minute predictions. The bottleneck is the readout stage + the training/inference geometry mismatch. Three structural changes ship together:
|
||||
|
||||
1. **Per-horizon EMA readout (MTER)** between trunk and heads, with τ[h] = K_h
|
||||
2. **Stateful CfC trunk** — drop attn_pool reset at sequence boundaries; use only at file-boundary bootstrap
|
||||
3. **Stateful MTER** — h_view state persists across sequences AND epochs; reset only at trainer construction + file boundary
|
||||
|
||||
Training and inference exercise the SAME continuous trunk_h trajectory through the SAME EMA dynamics.
|
||||
|
||||
## 3. Training architecture
|
||||
|
||||
### 3.1 Forward (per-event/per-position)
|
||||
|
||||
```
|
||||
α[h] = 1 / τ[h] # frozen τ by default
|
||||
h_view[b, h, :] = (1 - α[h]) · h_view_prev[b, h, :] + α[h] · trunk_h[b, :]
|
||||
```
|
||||
|
||||
τ_init = K_h = `[30, 100, 300, 1000, 6000]`. Frozen by default; `--mter-tau-mode learnable` available as opt-in ablation.
|
||||
|
||||
### 3.2 Concatenated head input
|
||||
|
||||
Each head reads `[h_view[b, h, :128], trunk_h[b, :128]]` (256-d).
|
||||
|
||||
### 3.3 Stateful CfC trunk
|
||||
|
||||
CfC `h_old` carries across sequence boundaries within a file. attn_pool used only at file-boundary bootstrap. BPTT remains K-truncated.
|
||||
|
||||
### 3.4 Stateful MTER
|
||||
|
||||
`h_view` persists across K-loop iterations, sequence boundaries, AND epoch boundaries. Reset only at trainer construction + file boundary. Warmup cost: ~0.2% of total training (single τ_max=6000 bootstrap per session) + ~1.6% across 8 file boundaries = ~1.8% total.
|
||||
|
||||
### 3.5 Validation loop state isolation
|
||||
|
||||
Separate `val_h_view_state_d`, `val_cfc_h_state_d` buffers. At val-start:
|
||||
- `val_cfc_h_state := attn_pool(val_file_positions_0..K-1)` (treated as a file-boundary bootstrap)
|
||||
- `val_h_view_state := sentinel-bootstrap on first observation`
|
||||
|
||||
Val data walked sequentially. AUC accumulator only counts samples after `events_since_val_start ≥ τ_max` (= 6000 events ≈ 188 sequences). Below that threshold, samples contribute to the forward pass but NOT to AUC.
|
||||
|
||||
For val_n_seqs=1000: ~812 sequences contribute to AUC, ~188 are warmup. AUC reflects the equilibrium-distribution measurement.
|
||||
|
||||
### 3.6 B = 1 enforcement
|
||||
|
||||
Trainer construction errors if `cfg.n_batch != 1`. Multi-batch sequential training is out of scope for v5.
|
||||
|
||||
### 3.7 Loader mode
|
||||
|
||||
`LoaderMode::Sequential` for MTER training + validation. `Random` preserved for legacy use and the fallback path (per §11).
|
||||
|
||||
### 3.8 Multi-fold trainer state
|
||||
|
||||
New method `PerceptionTrainer::reset_for_new_fold()`:
|
||||
- Re-sentinel-bootstraps all stateful buffers (h_view, cfc_h_state, val_*)
|
||||
- Sets first_obs := 0 across the board
|
||||
- Re-initializes any cumulative counters
|
||||
|
||||
Called by external multi-fold orchestrators. v5 first cut assumes one trainer per Argo run (one fold per trainer); the method exists to support future multi-fold workflows without silent state-leakage bugs.
|
||||
|
||||
### 3.9 Heads kernel signature
|
||||
|
||||
`multi_horizon_heads_grn_fwd_batched` input:
|
||||
- Old: `h [B, HIDDEN_DIM=128]`
|
||||
- New: `h_view [B, N_HORIZONS, HIDDEN_DIM]` + `trunk_h [B, HIDDEN_DIM]`
|
||||
|
||||
Per-horizon weights `w1 [N_HORIZONS, HEAD_MID=64, HIDDEN_DIM*2=256]`, `w_skip [N_HORIZONS, HIDDEN_DIM*2]`. Other weights unchanged.
|
||||
|
||||
`multi_horizon_heads_grn_bwd_batched`:
|
||||
- Output old: `grad_h [B, HIDDEN_DIM]` (summed across horizons)
|
||||
- Output new: `grad_h_view [B, N_HORIZONS, HIDDEN_DIM]` (per-horizon) + `grad_trunk_h_from_skip [B, HIDDEN_DIM]` (summed across horizons from trunk_h half of concat input)
|
||||
|
||||
Substantial rewrite of forward Pass 1, Pass 3 + backward Pass 5, Pass 6. Bit-equiv goldens regenerate.
|
||||
|
||||
### 3.10 τ_raw update
|
||||
|
||||
Frozen mode (default): τ_raw set once at construction, never updated.
|
||||
|
||||
Learnable mode: dedicated `mter_tau_update.cu` kernel (5 floats, gradient descent with simple optimizer state) runs OUTSIDE captured graph after each training step. Not integrated into the existing AdamW kernel.
|
||||
|
||||
## 4. Inference path design (NEW IN V5 — addresses concern A)
|
||||
|
||||
The inference path (`forward_step_into` per-event) needs symmetric MTER + stateful CfC handling so the heads see the same input distribution they trained on.
|
||||
|
||||
### 4.1 `forward_step_into` MTER integration
|
||||
|
||||
Sequence of operations per event in the captured inference graph:
|
||||
```
|
||||
1. snap_feature_assemble_batched (existing)
|
||||
2. VSN (variable selection) (existing)
|
||||
3. Mamba2 stack 1 step_into (existing)
|
||||
4. LN_a step (existing)
|
||||
5. Mamba2 stack 2 step_into (existing)
|
||||
6. LN_b step (existing)
|
||||
7. CfC step (h_old from cfc_h_state_step_d) (existing — stateful by default in inference)
|
||||
8. **MTER forward (NEW)** — read trunk_h (= cfc h_new), update h_view_state_step_d,
|
||||
write h_view_step_out
|
||||
9. heads_grn_fwd_batched (NEW INPUT SHAPE) — reads
|
||||
concat(h_view_step_out[h], trunk_h) for each horizon
|
||||
```
|
||||
|
||||
### 4.2 Inference-side state buffers
|
||||
|
||||
Add to `PerceptionTrainer` (or whichever struct owns inference):
|
||||
- `h_view_state_step_d: CudaSlice<f32>` — `[1, N_HORIZONS, HIDDEN_DIM]` per-instance MTER state, persistent across events within a session
|
||||
- `h_view_first_obs_step_d: CudaSlice<i32>` — `[1, N_HORIZONS]` sentinel
|
||||
- New `forward_step_into_with_log_ring` variant unchanged from the existing `forward_step_into` except for the new MTER kernel invocation
|
||||
|
||||
State lifecycle:
|
||||
- Bootstrap: at trainer construction (or first call to forward_step_into after a session reset).
|
||||
- Updated each event by the MTER kernel inside the captured graph.
|
||||
- Reset method: `reset_inference_state()` re-sentinels h_view_first_obs_step + re-bootstraps cfc_h_state_step. Called by external orchestrators between independent sessions.
|
||||
|
||||
### 4.3 LOB backtest harness integration
|
||||
|
||||
The CRT.diag harness (`lob-backtest-sweep` workflow → `fxt-backtest` binary or similar) loads a `PerceptionTrainer` checkpoint and walks events one at a time via `forward_step_into`. Required harness changes:
|
||||
|
||||
- Allocate the new inference-side state buffers at startup (same lifecycle as `cfc_h_state_step_d`).
|
||||
- Bootstrap h_view sentinel on first event.
|
||||
- The h_view state is naturally per-instance (one ring per `LobSimCuda` instance), parallel to `cfc_h_state_step_d`.
|
||||
|
||||
The existing `LobSimCuda::alpha_probs_d_mut` mapped-pinned output buffer is unchanged. MTER adds upstream state but doesn't change the output contract.
|
||||
|
||||
If the harness has its own kernel-launch sequence (not via `PerceptionTrainer::forward_step_into`), it needs the MTER launch inserted between CfC and heads — search for "heads_grn_fwd" in the harness codepath and add MTER one launch upstream.
|
||||
|
||||
### 4.4 Symmetry verification
|
||||
|
||||
The training-time `step_batched` and inference-time `forward_step_into` must produce numerically identical h_view + heads-input given identical inputs. New `tests/forward_step_mter_symmetry.rs`:
|
||||
- Run K=32 training step with all-deterministic inputs → record h_view_state + heads_input at k=K-1
|
||||
- Run K=32 individual `forward_step_into` calls with identical state setup → record h_view + heads_input at each step
|
||||
- Verify step-by-step equivalence within 1e-5
|
||||
|
||||
This is the equivalent of the existing `forward_step_golden.rs` but for the new MTER+heads geometry. Catches integration regressions cheaply.
|
||||
|
||||
## 5. Mathematical contract
|
||||
|
||||
**Forward at (b, h), frozen mode:**
|
||||
```
|
||||
If first_obs[b, h] == 0:
|
||||
h_view[b, h, :] := trunk_h[b, :]
|
||||
first_obs[b, h] := 1
|
||||
Else:
|
||||
h_view[b, h, :] := (1 - 1/τ[h]) · h_view_prev[b, h, :] + (1/τ[h]) · trunk_h[b, :]
|
||||
```
|
||||
|
||||
**Heads input:** `concat(h_view[b, h, :], trunk_h[b, :])` ∈ ℝ²⁵⁶.
|
||||
|
||||
**Backward at (b, h):**
|
||||
```
|
||||
d_trunk_h[b, :] += (1/τ[h]) · d_h_view[b, h, :] # from MTER EMA path
|
||||
d_trunk_h[b, :] += d_trunk_h_from_skip[b, :] # from heads' skip-concat path (already summed)
|
||||
d_h_view_carry[b, h, :] = (1 - 1/τ[h]) · d_h_view[b, h, :] # to previous step's h_view
|
||||
```
|
||||
|
||||
**Theoretical bound (NOT empirical claim):** under white-noise trunk_h, h_view[h]'s variance scales as 1/(2τ[h]−1). Under colored trunk_h, actual reduction is looser. JSONL trace records actual.
|
||||
|
||||
## 6. Implementation surface
|
||||
|
||||
### 6.1 NEW files
|
||||
|
||||
| File | Description | LOC est. |
|
||||
|---|---|---|
|
||||
| `crates/ml-alpha/cuda/mter_readout.cu` | MTER forward + backward kernels | ~300 |
|
||||
| `crates/ml-alpha/cuda/mter_tau_update.cu` | Tiny optimizer kernel for τ_raw (learnable mode only) | ~80 |
|
||||
| `crates/ml-alpha/tests/mter_invariants.rs` | 7 invariant tests (see §8.1) | ~300 |
|
||||
| `crates/ml-alpha/tests/forward_step_mter_symmetry.rs` | Training-vs-inference numerical symmetry (§4.4) | ~200 |
|
||||
|
||||
### 6.2 MODIFIED files
|
||||
|
||||
| File | Action | LOC est. |
|
||||
|---|---|---|
|
||||
| `crates/ml-alpha/cuda/multi_horizon_heads.cu` | Heads kernels accept per-horizon concat input. Forward Pass 1/3 loop bounds, backward Pass 5/6 grad split + accumulation. Substantial rewrite. | **+350-400 / -100 (delta)** |
|
||||
| `crates/ml-alpha/cuda/gpu_log_ids.h` | `KID_SMOOTHNESS_CONTROLLER` removed; slot 0 reclaimed by `KID_MTER_READOUT`; other KIDs shift down by 1 (KID_OUTPUT_SMOOTHNESS unused → also removed) | ~10 |
|
||||
| `crates/ml-alpha/build.rs` | Register `mter_readout`, `mter_tau_update`; deregister `output_smoothness`, `smoothness_lambda_controller`; bump cache-bust v14 → v15 | ~10 |
|
||||
| `crates/ml-alpha/src/heads.rs` | `TAU_INIT_PER_HORIZON`, `HEAD_INPUT_DIM = 2 * HIDDEN_DIM`, `MTerTauMode` enum; remove `SMOOTHNESS_LAMBDA_RATIO` | ~30 / -15 |
|
||||
| `crates/ml-alpha/src/trainer/perception.rs` | (a) New fields: `h_view_state_d`, `h_view_first_obs_d`, `tau_raw_d`, `grad_h_view_carry_d`, `val_h_view_state_d`, `val_h_view_first_obs_d`, `val_cfc_h_state_d`. (b) Inference-side fields: `h_view_state_step_d`, `h_view_first_obs_step_d`. (c) `reset_for_new_fold()` method. (d) `reset_inference_state()` method. (e) Constructor allocates + bootstraps. (f) Stateful CfC h_old carry (drop attn_pool reset at sequence boundaries). (g) `step_batched`: MTER fwd between CfC and heads; MTER bwd between heads_grn_bwd and CfC bwd. h_view persists across sequences + epochs. Validation uses separate state with post-warmup AUC accumulator. B=1 enforcement. (h) `forward_step_into`: MTER kernel launch added between CfC step and heads launch; h_view_state_step buffers updated each event. (i) New kernel-step-trace records for `KID_MTER_READOUT`. (j) Remove smoothness wiring (kernel handles, launch calls, telemetry buffers, accessor) | +400 / -250 |
|
||||
| `crates/ml-alpha/src/data/loader.rs` | `LoaderMode::Sequential` variant; sequential cursor advancement; `is_file_boundary()` method; `next_val_sequence()` sequential variant | ~150 |
|
||||
| `crates/ml-alpha/src/gpu_log.rs` | Replace smoothness decoders with MTER decoders (RT_INPUT/STATE/OUTPUT) | ~30 / -50 |
|
||||
| `crates/ml-alpha/examples/alpha_train.rs` | New flags: `--loader-mode {random|sequential}` (default sequential), `--mter-tau-mode {frozen|learnable}` (default frozen), `--mter-tau-init <comma-floats>` (optional). Drop `--smoothness-base-lambda`. Drop `final_smooth_loss_per_horizon` summary; replace with `final_h_view_norm_per_horizon` (telemetry from JSONL → summary aggregate). | +40 / -15 |
|
||||
| `infra/k8s/argo/alpha-perception-template.yaml` | New workflow params: `loader-mode`, `mter-tau-mode`, `mter-tau-init`. Drop `smoothness-base-lambda`. | +20 / -10 |
|
||||
|
||||
### 6.3 DELETED files (in the same commit chain)
|
||||
|
||||
| File | Reason |
|
||||
|---|---|
|
||||
| `crates/ml-alpha/cuda/output_smoothness.cu` | Superseded by MTER |
|
||||
| `crates/ml-alpha/cuda/smoothness_lambda_controller.cu` | Superseded by MTER |
|
||||
| `crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs` | Tests for deleted kernel |
|
||||
| `crates/ml-alpha/tests/output_smoothness_grad_finite_diff.rs` | Tests for deleted kernel |
|
||||
|
||||
### 6.4 REGENERATED test goldens
|
||||
|
||||
All four golden-value tests need regeneration because both heads kernel signature AND inference path change:
|
||||
|
||||
| Test | What regenerates |
|
||||
|---|---|
|
||||
| `tests/heads_bit_equiv.rs` | Heads kernel output goldens (new input shape `[B, N_HORIZONS, HIDDEN_DIM]` + `trunk_h`) |
|
||||
| `tests/forward_step_golden.rs` | Inference path's per-event output (MTER kernel now part of the chain) |
|
||||
| `tests/perception_forward_golden.rs` | Full training forward goldens (MTER + new heads + stateful CfC) |
|
||||
| `tests/forward_graph_capture.rs` | CUDA Graph captured-inference goldens (same as forward_step_golden but graph-captured) |
|
||||
|
||||
The bit-equiv tests are REGRESSION GATES, not correctness gates. After kernel rewrites, capture new goldens from a reference run + commit them. Each golden file ~1-10 KB.
|
||||
|
||||
### 6.5 Total scope
|
||||
|
||||
- New: ~880 LOC (kernels + tests + symmetry test)
|
||||
- Modified: ~660 LOC delta
|
||||
- Deleted: ~150 LOC
|
||||
- Test goldens regenerated: 4 files
|
||||
- **Net: ~1390 LOC** (+200 over v4 due to inference path + symmetry test)
|
||||
|
||||
Estimated dev time: 2-3 days. One L40S retrain cycle for validation.
|
||||
|
||||
## 7. Commit boundaries (phased — addresses concern E)
|
||||
|
||||
8 commits, each buildable + testable:
|
||||
|
||||
1. **Loader**: `LoaderMode::Sequential` + file-boundary detection + sequential val support. **Required smoke after this commit**: run a no-MTER training with `--loader-mode sequential` and verify trunk converges to within `random_baseline_AUC[h] − 0.02` per horizon (12 epochs, ~20 min L40S). If regression observed: STOP — MTER plan blocks. (~150 LOC)
|
||||
2. **MTER forward kernel** + Rust wiring + invariant tests 1-3. (~600 LOC)
|
||||
3. **Heads kernel forward reshape** + heads_bit_equiv.rs golden regenerate. (~200 LOC + golden file)
|
||||
4. **MTER backward kernel** + Rust wiring + invariant tests 4-7. (~450 LOC)
|
||||
5. **Heads kernel backward reshape** + MTER backward integration. (~250 LOC)
|
||||
6. **Stateful CfC trunk** + stateful MTER state across sequences/epochs + val state isolation + B=1 enforcement + post-warmup AUC accumulator. (~300 LOC)
|
||||
7. **Inference path** (forward_step_into MTER integration) + LOB backtest harness integration + `forward_step_mter_symmetry.rs` + golden regenerations for `forward_step_golden.rs`, `perception_forward_golden.rs`, `forward_graph_capture.rs`. (~250 LOC + 3 golden files)
|
||||
8. **Smoothness machinery deletion** + CLI/Argo template plumbing + `gpu_log_ids.h` reclaim of slot 0 + decoder rename. (~100 LOC + ~150 LOC deletion)
|
||||
|
||||
Each commit ships independently testable code. Commit 8 finalizes the architecture.
|
||||
|
||||
## 8. Memory rules + horizon_lambda interaction
|
||||
|
||||
### 8.1 Memory rules (unchanged from v4)
|
||||
|
||||
All preserved. `feedback_no_atomicadd`, `feedback_no_htod_htoh_only_mapped_pinned`, `feedback_no_nvrtc`, `pearl_no_host_branches_in_captured_graph`, `pearl_cudarc_disable_event_tracking_for_graph_capture`, `pearl_first_observation_bootstrap`, `pearl_blend_formulas_must_have_permanent_floor`, `feedback_no_partial_refactor` (atomic per commit), `feedback_wire_everything_up`, `feedback_single_source_of_truth_no_duplicates`, `feedback_no_feature_flags` (no Cargo features; runtime CLI mode flags only), `feedback_no_quickfixes` (structural), `feedback_no_functionality_removal` (smoothness removed because superseded; rollback via git revert).
|
||||
|
||||
### 8.2 horizon_lambda interaction
|
||||
|
||||
λ[h] applied in heads_grn_bwd BEFORE the split into d_h_view and d_trunk_h_from_skip. Both downstream paths inherit λ[h] scaling. Controller's stability range [0.5, 2.0] preserved. Invariant test #5 verifies.
|
||||
|
||||
## 9. Validation gates
|
||||
|
||||
### 9.1 Pre-MTER trunk validation (embedded in commit 1)
|
||||
|
||||
After commit 1 (loader change), smoke a no-MTER 12-epoch run with `--loader-mode sequential`. **Gate**: per-horizon val AUC within `random_baseline_AUC[h] − 0.02`. If pass, MTER proceeds (commits 2-8). If fail, MTER blocked.
|
||||
|
||||
### 9.2 Local invariant tests (`mter_invariants.rs` + `forward_step_mter_symmetry.rs`)
|
||||
|
||||
1. **Sentinel bootstrap** — first_obs=0 + arbitrary trunk_h → h_view := trunk_h exactly.
|
||||
2. **EMA equilibrium** — constant trunk_h for 10τ steps → h_view ≈ trunk_h within 1e-3.
|
||||
3. **Per-horizon timescale separation** — step-change in trunk_h; settling-time ratio (h6000:h30) ≈ τ ratio within 20%.
|
||||
4. **Backward finite-diff** — analytic d_trunk_h, d_h_view_carry match FD within 5e-3.
|
||||
5. **horizon_lambda × MTER** — λ=[1,1,1,1,2]; d_trunk_h via h6000 path is exactly 2× the unweighted case.
|
||||
6. **Stateful training across sequences** — h_view at step 6000 ≈ true 6000-event EMA within 1e-2.
|
||||
7. **CfC stateful across sequences** — sequence-boundary h_old = previous h_new at k=K-1; attn_pool used only at file boundaries.
|
||||
8. **Training/inference symmetry** (separate file) — `step_batched` and `forward_step_into` produce numerically identical h_view + heads_input within 1e-5.
|
||||
|
||||
### 9.3 MTER training MUST gate
|
||||
|
||||
- Per-horizon val AUC within `lnfwd_baseline_AUC[h] − 0.02` (baseline h6000 = 0.7157).
|
||||
- No NaN/divergence within 12 epochs.
|
||||
|
||||
Implicitly tests stateful CfC training; if MUST passes, trunk training is acceptable.
|
||||
|
||||
### 9.4 MTER training WIN gate (intermediate)
|
||||
|
||||
Theoretical bound: under colored trunk noise, training-time h6000_jitter / h30_jitter ∈ [0.05, 0.4]. Real differentiation = inside this range.
|
||||
|
||||
### 9.5 Inference WIN gate (THE gate)
|
||||
|
||||
- **MUST**: h6000 mean_run_len ≥ 100 events.
|
||||
- **MUST**: h6000/h30 mean_run_len ratio ≥ 10×.
|
||||
- **STRETCH**: CRT.1 controller smoke shows mean PnL monotonically rises with conviction OR (win_rate rises AND mean_pnl at highest-conviction bucket > -10).
|
||||
|
||||
By construction (τ[h6000] = 6000-event integration + stateful training matches inference distribution), h6000's per-event output should change on the order of 100s of events. WIN target should be achievable.
|
||||
|
||||
## 10. Risks
|
||||
|
||||
| Risk | Severity | Mitigation |
|
||||
|---|---|---|
|
||||
| Commit-1 sequential-sampling smoke fails (trunk regresses) | **High** | MTER plan blocks. Diagnose: trunk's training dynamics specifically depend on random sampling. Alternative: per-horizon CfC branches with random sampling (Option 1 from earlier reviews). Documented as explicit fork. |
|
||||
| MTER kernel orchestration more complex than estimated | Medium | Phased 8-commit implementation in §7 limits any broken state to one commit. |
|
||||
| Heads kernel rewrite breaks bit-equiv goldens | High (planned) | Golden values regenerated at commit 3 + 7. Tests are regression gates, not correctness gates. |
|
||||
| Inference path integration with LOB backtest harness | Medium | `forward_step_mter_symmetry.rs` (test #8) catches numerical drift between training and inference. Harness integration tested in commit 7. |
|
||||
| File-boundary attn_pool bootstrap creates regime-mismatched h_view warmup | Low | ~1.6% of total training. h_view's slow horizons filter out the discontinuity. |
|
||||
| Smoothness deletion makes rollback non-trivial | Medium | `git revert <MTER-commit-chain>` is rollback. Diagnostic JSONL from prior runs preserved on MinIO. Phased commits enable partial rollback. |
|
||||
| Val state warmup bias on AUC measurement | Low | Post-warmup AUC accumulator (samples only count after `events_since_val_start ≥ τ_max`). Trivial change. |
|
||||
| Multi-fold trainer reuse causes state leakage | Low | `reset_for_new_fold()` method added explicitly. Current Argo workflow is one-fold-per-trainer; method is for future workflows. |
|
||||
| Frozen τ is suboptimal vs learnable | Low | Default frozen for safety. Learnable available via `--mter-tau-mode learnable` as opt-in post-WIN ablation. |
|
||||
| 12-epoch L40S retrain wall | Negligible | ~36 min total (5 build + 18 train + 13 CRT.diag) for the validation cycle. |
|
||||
|
||||
## 11. Open questions
|
||||
|
||||
1. Drop attn_pool entirely (use zero-init at file boundaries)? Preserve for now.
|
||||
2. ISV-driven τ controller for learnable mode. Defer until learnable τ proven insufficient.
|
||||
3. Per-dimension τ vectors. Defer.
|
||||
4. B > 1 sequential training. Out of scope.
|
||||
5. Val sequence count tuning. Default `n_val_seqs=1000` + post-warmup AUC accumulator. May need ↑ if h6000 AUC measurement is too noisy.
|
||||
|
||||
## 12. Out of scope
|
||||
|
||||
- Per-dimension τ vectors
|
||||
- Cross-horizon attention at readout
|
||||
- ISV-driven τ controller
|
||||
- B > 1 sequential training
|
||||
- Dropping attn_pool entirely
|
||||
|
||||
## 13. Predecessor and successor
|
||||
|
||||
**Predecessor:**
|
||||
- `pearl_training_smoothness_does_not_transfer_to_inference`
|
||||
- `pearl_state_amplifies_short_horizon_into_long_horizon`
|
||||
- `pearl_separate_aux_trunk_when_shared_starves`
|
||||
|
||||
**Successor (if commit-1 smoke fails):**
|
||||
- MTER shelved. Alternative path: per-horizon CfC branches with random sampling. Bigger architectural change. Write separate spec.
|
||||
|
||||
**Successor (if MTER training MUST + WIN pass):**
|
||||
- No follow-on; this IS the architecture.
|
||||
- Ablation candidates: learnable τ, ISV-driven τ.
|
||||
- Resume CRT.1 controller on differentiated h_view outputs.
|
||||
|
||||
**Successor (if training passes but inference WIN fails):**
|
||||
- JSONL trace localizes: τ collapse (learnable mode), h_view non-equilibration, model over-reliance on trunk_h skip-path, stateful CfC training degeneracy.
|
||||
- Likely path: zero-init the trunk_h half of heads' w1 (force heads to use h_view primarily).
|
||||
- Fallback: redefine inference WIN gate from mean_run_len to conviction × outcome.
|
||||
|
||||
---
|
||||
|
||||
**End of CRT.train Intervention B v5 design. Awaiting review.**
|
||||
@@ -0,0 +1,135 @@
|
||||
# CRT.train Output-Smoothness — ISV-Driven λ Controller (Spec Extension)
|
||||
|
||||
**Date:** 2026-05-21
|
||||
**Branch context:** `ml-alpha-phase-a` (HEAD `70ecfc0fe`; static-λ implementation shipped in 10 commits)
|
||||
**Predecessor spec:** `2026-05-20-crt-train-output-smoothness-design.md`
|
||||
**Status:** Design — extends the predecessor by replacing the static `SMOOTHNESS_LAMBDA_RATIO` with a signal-driven controller anchored on observed h30 jitter.
|
||||
|
||||
---
|
||||
|
||||
## 1. Motivation
|
||||
|
||||
The predecessor spec ships `λ[h] = smoothness_base_lambda × SMOOTHNESS_LAMBDA_RATIO[h]` where the ratio array `[1, 3.33, 10, 33.33, 200]` is a tuned constant. This violates two memory rules:
|
||||
|
||||
- [pearl_controller_anchors_isv_driven](memory/pearl_controller_anchors_isv_driven.md) — every controller anchor/target/cap must be ISV-driven, not hardcoded.
|
||||
- [feedback_isv_for_adaptive_bounds](memory/feedback_isv_for_adaptive_bounds.md) — adaptive bounds live in ISV, not constants.
|
||||
|
||||
The K-ratio assumption (h6000 should change 200× more slowly than h30) is *probably right* but the model's intrinsic noise (the observed h30 jitter) is what we should anchor on, not the abstract ratio of horizon lengths.
|
||||
|
||||
## 2. Hypothesis
|
||||
|
||||
Anchor target per-horizon on the OBSERVED h30 jitter EMA, scale down per horizon by `K_h30 / K_h`. Drive λ[h] to push observed jitter toward target. Use a permanent floor so the controller never fully disables.
|
||||
|
||||
## 3. Architectural Intervention
|
||||
|
||||
### 3.1 State (new buffers on `PerceptionTrainer`)
|
||||
|
||||
- `jitter_ema_d: CudaSlice<f32>` shape `[5]` — running EMA of `raw_per_h`.
|
||||
- `jitter_first_obs_d: CudaSlice<i32>` shape `[1]` — sentinel `0 → 1` on first observation (per [pearl_first_observation_bootstrap](memory/pearl_first_observation_bootstrap.md)).
|
||||
|
||||
`smoothness_lambda_d` keeps its name but its semantics change: it is no longer uploaded once at construction; it is updated each step inside the captured graph.
|
||||
|
||||
### 3.2 New CUDA kernel `cuda/smoothness_lambda_controller.cu`
|
||||
|
||||
```
|
||||
extern "C" __global__ void smoothness_lambda_controller(
|
||||
const float* raw_per_h, // [5] emitted by output_smoothness_loss_and_grad
|
||||
float* jitter_ema, // [5] in/out — EMA state
|
||||
int* first_obs, // [1] in/out — sentinel
|
||||
float base_lambda, // scalar — amplitude knob from cfg.smoothness_base_lambda
|
||||
float* lambda_out // [5] new λ values for the kernel's NEXT step
|
||||
);
|
||||
```
|
||||
|
||||
Math (per thread `h ∈ {0..4}`, single block of 5 threads):
|
||||
|
||||
```
|
||||
// 1. EMA update with sentinel bootstrap.
|
||||
if (first_obs[0] == 0):
|
||||
jitter_ema[h] = raw_per_h[h]
|
||||
if (h == 0): first_obs[0] = 1
|
||||
else:
|
||||
// Fixed-α floor at 0.5 (Wiener-α floor for non-stationary loops
|
||||
// per pearl_wiener_alpha_floor_for_nonstationary). The target
|
||||
// drifts as the policy co-adapts, so MSE-optimal α isn't right.
|
||||
jitter_ema[h] = 0.5 * jitter_ema[h] + 0.5 * raw_per_h[h]
|
||||
|
||||
__syncthreads() // all threads need jitter_ema[0] for target derivation
|
||||
|
||||
// 2. Target = ISV-anchored on observed h30 jitter.
|
||||
// target[0] = jitter_ema[0] ← h30 self-targets
|
||||
// target[h] = jitter_ema[0] * (HORIZONS[0] / HORIZONS[h]) for h ≥ 1
|
||||
float target_h
|
||||
if (h == 0):
|
||||
target_h = jitter_ema[0]
|
||||
else:
|
||||
target_h = jitter_ema[0] * (HORIZONS[0] / HORIZONS[h])
|
||||
|
||||
// 3. λ controller: push UP when observed > target; relax to floor otherwise.
|
||||
float excess_ratio = max(0.0, jitter_ema[h] / max(target_h, 1e-9) - 1.0)
|
||||
float lambda_new = base_lambda * (1.0 + excess_ratio)
|
||||
lambda_out[h] = max(LAMBDA_FLOOR, lambda_new)
|
||||
```
|
||||
|
||||
where `LAMBDA_FLOOR = 1e-4` — the permanent floor (per [pearl_blend_formulas_must_have_permanent_floor](memory/pearl_blend_formulas_must_have_permanent_floor.md): `max(real, floor)`, never blend).
|
||||
|
||||
### 3.3 Launch order in `step_batched`
|
||||
|
||||
Replace the current sequence:
|
||||
```
|
||||
... → BCE → output_smoothness → backward K-loop
|
||||
```
|
||||
with:
|
||||
```
|
||||
... → BCE → output_smoothness → smoothness_lambda_controller → backward K-loop
|
||||
```
|
||||
|
||||
The controller runs AFTER `output_smoothness` (which emits `raw_per_h`) and BEFORE the backward K-loop. The updated `smoothness_lambda_d` is used by the NEXT step's `output_smoothness` launch — one-step delay, standard closed-loop pattern.
|
||||
|
||||
### 3.4 Removals
|
||||
|
||||
- `pub const SMOOTHNESS_LAMBDA_RATIO: [f32; N_HORIZONS]` in `heads.rs` — DELETED (no longer load-bearing; ratio is now derived from `HORIZONS`).
|
||||
- The constructor's λ pre-multiplication block in `perception.rs` — REPLACED with `λ_init = [LAMBDA_FLOOR; 5]` upload.
|
||||
|
||||
The `cfg.smoothness_base_lambda` knob stays — it's the amplitude (controller equilibrium value when `excess_ratio == 0`).
|
||||
|
||||
## 4. Implementation Surface
|
||||
|
||||
| File | Change |
|
||||
|------|--------|
|
||||
| `crates/ml-alpha/cuda/smoothness_lambda_controller.cu` | NEW |
|
||||
| `crates/ml-alpha/build.rs` | Add `"smoothness_lambda_controller"` to KERNELS, bump cache-bust v12 → v13 |
|
||||
| `crates/ml-alpha/src/heads.rs` | DELETE `SMOOTHNESS_LAMBDA_RATIO` const |
|
||||
| `crates/ml-alpha/src/trainer/perception.rs` | Add `jitter_ema_d`, `jitter_first_obs_d`, controller fn handle + module; init λ to LAMBDA_FLOOR; launch controller in `step_batched` |
|
||||
| `crates/ml-alpha/tests/output_smoothness_grad_finite_diff.rs` | NO change — existing tests still valid (they test the smoothness kernel; the controller is separate) |
|
||||
| `crates/ml-alpha/tests/smoothness_lambda_controller_invariants.rs` | NEW — invariant tests for the controller |
|
||||
|
||||
The CLI flag `--smoothness-base-lambda` and the Argo plumbing stay unchanged.
|
||||
|
||||
## 5. Validation Gate Update
|
||||
|
||||
Spec §5 MUST/WIN/STRETCH gates unchanged. ADD a controller-specific invariant test:
|
||||
|
||||
**Controller MUST gates** (unit tests, no training run needed):
|
||||
1. First-observation sentinel: `first_obs == 0` → kernel sets `jitter_ema = raw_per_h`, `first_obs = 1`, `λ = [LAMBDA_FLOOR; 5]`.
|
||||
2. Steady-state symmetry: if `raw_per_h[h] / raw_per_h[0] == HORIZONS[0] / HORIZONS[h]` for all h, the controller emits `λ[h] = base_lambda` exactly (excess_ratio = 0 everywhere).
|
||||
3. Excess response: if `jitter_ema[h6000]` is 10× its target, `λ[h6000] = base_lambda * 10`.
|
||||
4. Floor invariant: when `base_lambda = 0`, `λ[h] = LAMBDA_FLOOR` for all h. The kernel does NOT collapse to producing zero loss + zero grad (because LAMBDA_FLOOR > 0); that's an intentional change vs the static-λ spec — the floor is a safety net.
|
||||
|
||||
If you want the "static λ=0 is a no-op" property back (kernel literally produces nothing), set `LAMBDA_FLOOR = 0.0` — but then the controller can self-disable and that violates the user's "with floor" requirement. Resolution: `LAMBDA_FLOOR = 1e-4` chosen so the residual smoothing gradient is empirically negligible (~`2 × 1e-4 × (Δp)² / N_pairs` ≈ `1e-7` per gradient entry — well below f32 grad noise).
|
||||
|
||||
## 6. Open Questions
|
||||
|
||||
1. **First-observation sentinel race?** With 5 threads writing `first_obs` after `__syncthreads()`, only thread `h == 0` writes the sentinel — single writer, no race. Verify in code.
|
||||
|
||||
2. **Should `target_h` for h=0 always equal `jitter_ema[0]` (self-target, controller passive)?** Yes per design; h30 has no slower horizon to anchor on, so it tracks the amplitude knob directly. `excess_ratio[0] = 0` always.
|
||||
|
||||
3. **`max(target_h, 1e-9)` epsilon — is this a tuned constant?** Technically yes, but it's a denominator-safety epsilon, not a behavioral knob. Per existing codebase convention (e.g., `REGIME_EPS = 1e-6` in `data/loader.rs`), epsilon constants are exempt from the no-tuned-constants rule.
|
||||
|
||||
## 7. Backward Compatibility
|
||||
|
||||
The CLI flag `--smoothness-base-lambda` stays — operator can still sweep it for amplitude. Default 0.0 used to be "kernel produces zero loss/grad"; now it means "controller maintains λ = LAMBDA_FLOOR for all h." Behaviorally near-equivalent (LAMBDA_FLOOR is tiny) but not literally zero. If literally-zero behavior is required, set `LAMBDA_FLOOR = 0.0` — but this conflicts with the user's "with floor" requirement.
|
||||
|
||||
---
|
||||
|
||||
**End of spec extension. Awaiting plan and execution.**
|
||||
282
docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md
Normal file
282
docs/superpowers/specs/2026-05-21-gpu-log-ring-design.md
Normal file
@@ -0,0 +1,282 @@
|
||||
# GPU Log Ring — Unified High-Performance In-Kernel Diagnostic Logging (Design)
|
||||
|
||||
**Date:** 2026-05-21
|
||||
**Branch context:** `ml-alpha-phase-a` (HEAD `0f664cf36`; ISV-driven λ controller shipped but per-step trajectory is invisible — `final_smooth_loss_per_horizon` is the only smoothness telemetry, emitted once at end-of-run)
|
||||
**Predecessor:** `2026-05-21-crt-train-isv-driven-lambda-controller.md` (the controller whose internal state we cannot inspect)
|
||||
**Status:** Design — pending plan + execution
|
||||
|
||||
---
|
||||
|
||||
## 1. Motivation
|
||||
|
||||
ml-alpha kernels run inside a captured CUDA graph at ~370 steps/sec on L40S. To diagnose why a controller is or is not converging, we need per-step visibility into its inputs, state, and outputs. Today the only available telemetry channels are:
|
||||
|
||||
- `tracing::info!` from Rust, which can only see what the Rust side already computes (per-epoch validation outputs, the existing `isv snapshot` log line driven by `horizon_lambda.cu`). Anything that lives only inside a kernel — like `smoothness_lambda_controller`'s `jitter_ema`, `target`, `excess_ratio`, `lambda_out` — is invisible without a DtoH copy after every step.
|
||||
- Per-step DtoH copy of every diagnostic buffer: forbidden by [feedback_no_htod_htoh_only_mapped_pinned](memory/feedback_no_htod_htoh_only_mapped_pinned.md), [feedback_cpu_is_read_only](memory/feedback_cpu_is_read_only.md), and adds a synchronisation barrier in the hot path.
|
||||
- Per-step `printf` from the kernel: violates [pearl_no_host_branches_in_captured_graph](memory/pearl_no_host_branches_in_captured_graph.md) and stalls every block on a serialised CPU queue.
|
||||
|
||||
The empirical motivator: at base_lambda=0.001 the smoothness regularizer produced +6.5pt mean AUC but `final_smooth_loss_per_horizon = [0.082, 0.079, 0.075, 0.082, 0.076]` — no horizon differentiation. We cannot tell whether (a) the controller never amplified λ at all, (b) λ amplified but oscillated, (c) λ amplified steadily but the gradient was overpowered by BCE, or (d) something else. The right tool is per-step λ + jitter_ema + target trajectory. The right channel is a unified GPU log ring readable by host without DtoH copy.
|
||||
|
||||
## 2. Hypothesis
|
||||
|
||||
A mapped-pinned ring buffer with deterministic slot reservation (no `atomicAdd`), header-last commit ordering, and step-counter-derived addressing yields:
|
||||
- One coalesced 64 B write per logging event from the kernel hot path
|
||||
- Zero DtoH copies, zero kernel-host synchronisation
|
||||
- Full per-step trajectory for every kernel that opts in
|
||||
- Drop-detection on the host side (step-counter gaps in successive valid records)
|
||||
|
||||
This solves the diagnostic gap without violating the GPU/CPU contract memory rules.
|
||||
|
||||
## 3. Architectural Intervention
|
||||
|
||||
### 3.1 Record format (64 B, cacheline-aligned)
|
||||
|
||||
```c
|
||||
struct LogHeader { // 16 B
|
||||
uint32_t magic; // 4 — 0xF0XHN710; commit sentinel, written LAST
|
||||
uint32_t step; // 4 — monotonic step counter
|
||||
uint8_t kernel_id; // 1 — kernel emitting (see gpu_log_ids.h)
|
||||
uint8_t record_type; // 1 — sub-type within the kernel
|
||||
uint8_t payload_words; // 1 — # of f32 in payload (≤ 12)
|
||||
uint8_t reserved; // 1
|
||||
uint32_t _padding; // 4
|
||||
};
|
||||
|
||||
struct LogRecord { // 64 B
|
||||
LogHeader header; // 16 B
|
||||
float payload[12]; // 48 B = 12 × f32
|
||||
};
|
||||
```
|
||||
|
||||
Total: 64 B = one cacheline; single coalesced store per emission.
|
||||
|
||||
The `magic` field is the **commit sentinel**: written last, after the payload and the other header fields. Host readers check `magic == LOG_MAGIC` and `step == expected_step_for_this_slot` to validate. A torn record (kernel mid-write while CPU reads) reads either the OLD magic (skip) or matches the new magic but not the expected step (skip).
|
||||
|
||||
### 3.2 Slot reservation without atomicAdd
|
||||
|
||||
A 1-thread tick kernel runs ONCE per training step (first in the captured graph, before any logging kernel) and increments a device counter:
|
||||
|
||||
```c
|
||||
extern "C" __global__ void gpu_log_tick(int* step_counter) {
|
||||
if (threadIdx.x == 0 && blockIdx.x == 0) {
|
||||
step_counter[0] += 1;
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
Single thread, single block, no atomics needed (sole writer per launch). Lives at the head of the captured graph; trainer never touches it from host.
|
||||
|
||||
Each logging kernel then derives its slot at emission time:
|
||||
```
|
||||
slot = (step * N_RECORDS_PER_STEP + kernel_id * N_RT_PER_KERNEL + record_type) & (N_SLOTS - 1)
|
||||
```
|
||||
|
||||
with N_KERNELS=8, N_RT_PER_KERNEL=4, N_RECORDS_PER_STEP=32, N_SLOTS=32768 (all power-of-2). Modulo becomes a bitwise AND. Each (step, kernel_id, record_type) maps to exactly ONE slot — no contention, no atomics, no allocation.
|
||||
|
||||
Step-wrap window: 32768 / 32 = 1024 steps ≈ 2.8 s at 370 steps/sec. Drain cadence 500 ms keeps ring at < 20% utilisation.
|
||||
|
||||
### 3.3 Single-writer commit (header-last)
|
||||
|
||||
```c
|
||||
__device__ __forceinline__ void log_record(
|
||||
LogRing* ring,
|
||||
const int* step_counter,
|
||||
uint8_t kernel_id,
|
||||
uint8_t record_type,
|
||||
const float* payload,
|
||||
int payload_words
|
||||
) {
|
||||
// Thread (0,0) of block 0 only — single writer per (step, kernel_id, record_type).
|
||||
if (blockIdx.x | threadIdx.x) return;
|
||||
const int step = step_counter[0];
|
||||
const int slot = (step * N_RECORDS_PER_STEP
|
||||
+ ((int)kernel_id) * N_RT_PER_KERNEL
|
||||
+ (int)record_type) & (N_SLOTS - 1);
|
||||
LogRecord* rec = &ring->records[slot];
|
||||
|
||||
// Payload + non-magic header fields first.
|
||||
#pragma unroll
|
||||
for (int i = 0; i < payload_words; ++i) rec->payload[i] = payload[i];
|
||||
rec->header.step = (uint32_t)step;
|
||||
rec->header.kernel_id = kernel_id;
|
||||
rec->header.record_type = record_type;
|
||||
rec->header.payload_words = (uint8_t)payload_words;
|
||||
rec->header.reserved = 0;
|
||||
rec->header._padding = 0;
|
||||
|
||||
__threadfence(); // make all the above globally visible …
|
||||
rec->header.magic = LOG_MAGIC; // … BEFORE the magic commit sentinel.
|
||||
}
|
||||
```
|
||||
|
||||
### 3.4 Mapped-pinned host visibility + drain
|
||||
|
||||
The ring is allocated via a `MappedF32Buffer`-style mapped-pinned host allocation (extended to support arbitrary record sizes; see §8). Device gets a device pointer to the same memory via `cudaHostGetDevicePointer`. Host sees writes via volatile reads — no DtoH memcpy, no synchronisation. Per [feedback_no_htod_htoh_only_mapped_pinned](memory/feedback_no_htod_htoh_only_mapped_pinned.md): only legitimate cross-direction channel.
|
||||
|
||||
CPU drainer is a tokio task in `alpha_train.rs` (or the trainer's `step_batched` driver, depending on lifetime). It:
|
||||
1. Reads the device-side `step_counter[0]` via mapped-pinned (one volatile u32 read).
|
||||
2. For each slot in `[last_drained_step .. step_counter] × N_RECORDS_PER_STEP`, reads `record.header.magic` — if not LOG_MAGIC, skip silently. If LOG_MAGIC and `header.step == expected_step_for_slot`, dispatch to the per-(kernel_id, record_type) decoder.
|
||||
3. Decoder emits a structured `tracing::info!` line with the payload fields named per the record type.
|
||||
|
||||
Drop detection: if the gap between successive valid records' `step` fields exceeds N_SLOTS/N_RECORDS_PER_STEP, the drainer fell behind. Emit a `tracing::warn!` with the gap size; downstream analysis sees the discontinuity.
|
||||
|
||||
### 3.5 Kernel-author API
|
||||
|
||||
A single header `cuda/gpu_log_ids.h` defines the kernel_id and record_type enums; both the kernels and the Rust decoder include/reference it.
|
||||
|
||||
```c
|
||||
// gpu_log_ids.h
|
||||
#define KID_SMOOTHNESS_CONTROLLER 0
|
||||
#define KID_OUTPUT_SMOOTHNESS 1
|
||||
#define KID_BCE 2
|
||||
#define KID_HORIZON_LAMBDA 3
|
||||
#define KID_HEADS_GRN_FWD 4
|
||||
#define KID_HEADS_GRN_BWD 5
|
||||
#define KID_CFC_STEP 6
|
||||
#define KID_RESERVED_7 7
|
||||
|
||||
#define RT_INPUT 0
|
||||
#define RT_STATE 1
|
||||
#define RT_OUTPUT 2
|
||||
#define RT_DEBUG 3
|
||||
```
|
||||
|
||||
In a kernel:
|
||||
```c
|
||||
#include "gpu_log_ids.h"
|
||||
|
||||
extern "C" __global__ void smoothness_lambda_controller(
|
||||
/* … existing args … */,
|
||||
LogRing* g_log_ring, // NEW: optional, NULL = no logging
|
||||
const int* g_step_counter // NEW
|
||||
) {
|
||||
/* … existing logic … */
|
||||
|
||||
if (g_log_ring != nullptr) {
|
||||
float payload[10] = { jitter_ema[0], jitter_ema[1], jitter_ema[2], jitter_ema[3], jitter_ema[4],
|
||||
lambda_out[0], lambda_out[1], lambda_out[2], lambda_out[3], lambda_out[4] };
|
||||
log_record(g_log_ring, g_step_counter, KID_SMOOTHNESS_CONTROLLER, RT_STATE, payload, 10);
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
The two new args (`g_log_ring`, `g_step_counter`) are device pointers, stable across captures. When logging is disabled at construction time, the trainer passes NULL for `g_log_ring` and the kernel skips the emission — no diagnostic overhead.
|
||||
|
||||
### 3.6 Decoder side (Rust)
|
||||
|
||||
A central `gpu_log.rs` defines:
|
||||
- `pub fn alloc_log_ring(n_slots: usize) -> Result<(MappedBuffer<LogRecord>, CudaSlice<i32>)>` — returns mapped-pinned ring + the device step_counter.
|
||||
- `pub fn drain_loop(ring: &MappedBuffer<LogRecord>, step_counter: &CudaSlice<i32>, ...) -> JoinHandle<()>` — tokio task that runs until cancelled, periodically draining new records and dispatching to decoders.
|
||||
- Decoders are registered via a const table indexed by `(kernel_id, record_type)`:
|
||||
```rust
|
||||
type DecoderFn = fn(step: u32, payload: &[f32]);
|
||||
static DECODERS: [[DecoderFn; N_RT_PER_KERNEL]; N_KERNELS] = [...];
|
||||
```
|
||||
- The default decoder (no entry registered) emits a generic `tracing::debug!("unknown log record kernel_id={} record_type={}", ...)`.
|
||||
|
||||
### 3.7 Wiring lifecycle
|
||||
|
||||
- Trainer construction: alloc log ring + step counter; pass device pointers into every kernel handle that opts in.
|
||||
- Trainer step: tick kernel runs first in the captured graph (incrementing step_counter); all logging-capable kernels then derive their slots from it.
|
||||
- Trainer drop: drain task cancelled, ring buffer dropped (pinned memory unmapped).
|
||||
|
||||
Compile-time gating: a Rust feature flag `cuda-diag-log` enables ring allocation + drain task. When disabled, the trainer passes NULL pointers and the drain task never starts — zero cost in production training runs.
|
||||
|
||||
## 4. Memory-Rule Compliance
|
||||
|
||||
| Rule | Compliance |
|
||||
|------|------------|
|
||||
| `feedback_no_atomicadd` | Slot reservation by `step × N_RPS + kernel_id × N_RT + record_type` modulo — single-writer per slot. No atomics. |
|
||||
| `feedback_no_htod_htoh_only_mapped_pinned` | Ring is mapped-pinned; host reads via volatile, no DtoH memcpy. |
|
||||
| `feedback_no_nvrtc` | Tick kernel + log helpers pre-compiled via build.rs. |
|
||||
| `pearl_no_host_branches_in_captured_graph` | Kernels' new args are device pointers, stable across captures. No host branching. |
|
||||
| `pearl_cudarc_disable_event_tracking_for_graph_capture` | No host malloc, no scalar arg changes between captures, no pointer swaps. |
|
||||
| `pearl_first_observation_bootstrap` | N/A (no EMA state in the ring) |
|
||||
| `pearl_one_unbounded_signal_per_reward` | N/A (diagnostic, not loss) |
|
||||
| `feedback_no_stubs`, `feedback_no_todo_fixme`, `feedback_no_hiding` | Implementation must complete every site or delete; no markers; no `_` hiding. |
|
||||
| `feedback_no_partial_refactor` | When a kernel signature gains the two new pointer args, every consumer must migrate atomically. |
|
||||
| `feedback_single_source_of_truth_no_duplicates` | One ring per training process. One header file (`gpu_log_ids.h`) is the source of truth for IDs. |
|
||||
| `feedback_wire_everything_up` | Ring is allocated, tick fires, decoders dispatch, drain task runs — no orphans. |
|
||||
|
||||
## 5. Validation Gate
|
||||
|
||||
Unit tests (no training run needed):
|
||||
|
||||
1. **Round-trip**: emit a known record from a test kernel; drain; verify decoder reconstructs payload exactly. Asserts no off-by-one in slot arithmetic, no payload truncation, no endianness surprise.
|
||||
2. **Magic validation**: write a deliberately torn record (payload only, no magic); verify drainer skips it without panic.
|
||||
3. **Step monotonicity**: emit records at steps 0, 1, 2, …; verify drainer sees monotonic step counter without gaps.
|
||||
4. **Ring wrap**: emit > N_SLOTS records; verify drain catches all (assuming fast enough) and reports wrap discontinuity correctly when slow.
|
||||
5. **NULL ring**: launch kernel with NULL ring pointer; verify no crash, no record emitted.
|
||||
6. **Multi-kernel coexistence**: emit from 3 different kernel_ids simultaneously in one step; verify all 3 records distinct and decodable.
|
||||
|
||||
Integration test:
|
||||
|
||||
7. **Live smoothness trajectory**: enable cuda-diag-log feature, run 100 trainer steps, drain, verify per-step records exist for smoothness_lambda_controller (RT_INPUT, RT_STATE, RT_OUTPUT) showing jitter_ema and lambda evolving monotonically toward equilibrium.
|
||||
|
||||
These are runnable as `cargo test -p ml-alpha --features cuda-diag-log --test gpu_log_ring_invariants`.
|
||||
|
||||
## 6. Risks
|
||||
|
||||
| Risk | Severity | Mitigation |
|
||||
|------|----------|------------|
|
||||
| Drainer falls behind, ring wraps, records lost | Low | Step-gap detection emits warn; bump N_SLOTS if observed in practice. |
|
||||
| Memory fence cost slows hot path | Low | Only thread (0,0) writes; fence cost amortised over kernel's natural completion fence. Benchmark before/after smoke. |
|
||||
| Decoder dispatch table drifts from kernel-side IDs | Medium | Single shared `gpu_log_ids.h` header included by both kernels and the Rust generated wrapper. Build.rs may verify counts match. |
|
||||
| Adding two pointer args to kernels triggers partial-refactor compile errors | Medium | Atomic refactor commits (per `feedback_no_partial_refactor`); feature flag gates the new args at the source level. |
|
||||
| Captured-graph re-record invalidates if pointers change | Low | Pointers are allocated once at trainer construction; never freed during training. |
|
||||
| Magic collision (random memory matches LOG_MAGIC) | Negligible | 32-bit sentinel + step-count match doubles up the validation. |
|
||||
| Performance regression on the captured-graph hot path | Low | The tick kernel is 1 thread, 1 block, 1 instruction. Logging kernels add 1 cacheline write. Microsecond-scale overhead. |
|
||||
|
||||
## 7. Open Questions
|
||||
|
||||
1. **Should the drainer run in `alpha_train` as a tokio task, or as a sidecar process attached to the same mapped-pinned buffer via shared-memory IPC?** Tokio task is simpler; sidecar would allow real-time monitoring tools to attach. Default: tokio task in-process; expose the buffer geometry so a sidecar can be added later without code changes.
|
||||
|
||||
2. **Variable-length payloads?** Current design fixes payload at ≤ 12 floats per record. Larger records (e.g., 256 B for matrix dumps) would require either a separate ring or a "continuation record" pattern. Defer until a use case appears.
|
||||
|
||||
3. **Drainer cadence — per-step vs per-second?** Per-second (500 ms tick) is sufficient for 2.8 s ring window. Per-step adds CPU contention for no diagnostic benefit. Default 500 ms.
|
||||
|
||||
4. **`gpu_log_ids.h` location.** `crates/ml-alpha/cuda/gpu_log_ids.h` is the kernel-side home. Rust side via `include_str!` + a small parser, OR a hand-maintained mirror in `gpu_log.rs`. Build.rs check would catch drift. Default: hand-mirror with build.rs verification.
|
||||
|
||||
5. **Should logging be on by default or opt-in?** Opt-in via cuda-diag-log feature flag. Production runs ship without it; investigation runs enable it.
|
||||
|
||||
6. **What happens if a logging-capable kernel is launched with `g_log_ring == NULL` but `g_step_counter != NULL` (or vice versa)?** Kernel branches on `g_log_ring != nullptr` only; step counter is unused without the ring. Safe under partial wiring during incremental rollout.
|
||||
|
||||
## 8. Implementation Surface
|
||||
|
||||
| File | Action |
|
||||
|------|--------|
|
||||
| `crates/ml-alpha/cuda/gpu_log_ring.cu` | NEW — defines `LogRecord`, `LogRing`, `log_record()`, the tick kernel |
|
||||
| `crates/ml-alpha/cuda/gpu_log_ids.h` | NEW — kernel_id and record_type enum macros (single source of truth) |
|
||||
| `crates/ml-alpha/build.rs` | Register `gpu_log_ring` in `KERNELS`; bump cache-bust v13 → v14 |
|
||||
| `crates/ml-alpha/src/gpu_log.rs` | NEW — Rust-side ring allocator, drain task, decoder registry |
|
||||
| `crates/ml-alpha/src/pinned_mem.rs` | Extend with a generic `MappedRecordBuffer<T>` (preserve `MappedF32Buffer` as type alias) |
|
||||
| `crates/ml-alpha/src/lib.rs` | Expose `pub mod gpu_log` behind `#[cfg(feature = "cuda-diag-log")]` |
|
||||
| `crates/ml-alpha/Cargo.toml` | Add `cuda-diag-log` feature |
|
||||
| `crates/ml-alpha/src/trainer/perception.rs` | Add `log_ring_d`, `log_step_counter_d` fields; allocate at construction; pass to logging-capable kernels |
|
||||
| `crates/ml-alpha/cuda/smoothness_lambda_controller.cu` | Add `g_log_ring`, `g_step_counter` args; emit RT_INPUT, RT_STATE, RT_OUTPUT records |
|
||||
| `crates/ml-alpha/tests/gpu_log_ring_invariants.rs` | NEW — 6 unit tests + 1 integration test |
|
||||
|
||||
Plus the wire-up commits for each subsequent kernel as it opts in (BCE, output_smoothness, horizon_lambda, heads_grn_fwd, heads_grn_bwd, cfc_step). Initial spec scope wires only smoothness_lambda_controller; remaining kernels are follow-on tasks once the ring is proven.
|
||||
|
||||
Estimated total LOC (initial scope): ~550 lines new code + ~50 lines of perception.rs wiring + per-kernel wiring (5 lines × 1 kernel initially).
|
||||
|
||||
## 9. Out of Scope (deferred)
|
||||
|
||||
- Sidecar IPC drain (a separate process attaching to the mapped-pinned buffer)
|
||||
- Variable-length payload records
|
||||
- Compression / quantisation of recorded floats
|
||||
- Auto-generation of the kernel_id/record_type enums via macros (manual maintenance + build.rs check)
|
||||
- Wire-up of kernels other than `smoothness_lambda_controller` in the first plan; each subsequent kernel is its own atomic refactor in a follow-on commit.
|
||||
- Visualisation tooling (plotly/grafana for the drained records) — out of scope; the tracing output is structured enough for offline analysis with jq/pandas.
|
||||
|
||||
## 10. Predecessor and Successor
|
||||
|
||||
**Predecessor:** `2026-05-21-crt-train-isv-driven-lambda-controller.md` — the controller whose per-step trajectory we cannot inspect.
|
||||
|
||||
**Successor (if WIN gate fails at base_lambda=0.001 and a second retrain is required):** The log ring enables per-step inspection of the controller's λ trajectory in the second retrain, telling us whether the issue is amplitude (mechanism #1 in the post-mortem) or trunk-sharing (mechanism #2 → intervention B).
|
||||
|
||||
**Successor (if WIN gate passes at base_lambda=0.001):** The ring is still valuable infrastructure for future controllers and kernel debugging, even if not immediately critical. Lower priority in that branch.
|
||||
|
||||
---
|
||||
|
||||
**End of GPU log ring design. Awaiting plan and execution.**
|
||||
@@ -0,0 +1,632 @@
|
||||
# Per-Horizon CfC Inference Architecture — Design
|
||||
|
||||
**Date**: 2026-05-21
|
||||
**Branch**: `ml-alpha-phase-a`
|
||||
**Author**: brainstorming session synthesis
|
||||
**Status**: design awaiting plan
|
||||
**Supersedes**: `2026-05-21-crt-train-intervention-b-multi-timescale-readout.md` (MTER, premise falsified)
|
||||
|
||||
---
|
||||
|
||||
## 1. Problem statement and verified mechanism
|
||||
|
||||
### 1.1 The bottleneck
|
||||
|
||||
`pearl_training_smoothness_does_not_transfer_to_inference` documented an empirical observation: the K-window output-smoothness regularizer (intervention A) produced real training-time horizon stratification (jitter ratio 0.40 between h30 and h6000 outputs) but ZERO inference-time `mean_run_len` differentiation (ratio 0.93). The original WIN gate (`mean_run_len[h6000] / mean_run_len[h30] ≥ 10×`, h6000 mean_run_len ≥ 100 events) was not approachable.
|
||||
|
||||
`pearl_random_sampling_sufficient_for_per_horizon_training_auc` (recorded 2026-05-21 during the Phase-1 diagnostic) further established that random batch sampling achieves identical peak per-horizon validation AUC to stateful sequential sampling (both ~0.745 at h6000). This refuted MTER's premise that training/inference EMA-distribution alignment was the bottleneck.
|
||||
|
||||
The remaining hypothesis: the bottleneck is on the inference-side architecture itself — specifically the CfC layer's per-channel τ dynamics under the K=32 training window.
|
||||
|
||||
### 1.2 The mechanism (math-derived, user-verified)
|
||||
|
||||
From `crates/ml-alpha/cuda/cfc_step.cu:57-58`:
|
||||
|
||||
```c
|
||||
const float decay = expf(-dt_s / fmaxf(tau[i], 1e-6f));
|
||||
h_new[i] = h_old[i] * decay + (1.0f - decay) * tanhf(pre[i]);
|
||||
```
|
||||
|
||||
CfC `τ` is per-channel (`Vec<f32>` of length `n_hid=HIDDEN_DIM`), initialized log-uniform over `[0.01, 1000]` per `crates/ml-alpha/src/cfc/trunk.rs:204-209`. `τ` is also trained.
|
||||
|
||||
For a channel with `τ_i = 6000` and `dt = 1` (event-rate inference, per `forward_step_into` perception.rs:3517):
|
||||
|
||||
| Quantity | Value |
|
||||
|---|---|
|
||||
| Per-step decay | `exp(-1/6000) ≈ 0.99983` |
|
||||
| Per-step input contribution `(1 − decay)` | `≈ 0.00017` |
|
||||
| Cumulative h after K=32 training steps from h=0 | `≈ 32 × 0.00017 × tanh(pre) ≈ 0.005 × tanh(pre)` |
|
||||
| Effective state magnitude during training | near zero |
|
||||
|
||||
For a channel with `τ_i = 1`:
|
||||
|
||||
| Quantity | Value |
|
||||
|---|---|
|
||||
| Per-step decay | `exp(-1) ≈ 0.37` |
|
||||
| Cumulative h after a few K=32 training steps | `≈ tanh(pre[last])` |
|
||||
| Effective state magnitude during training | full range |
|
||||
|
||||
Consequences:
|
||||
|
||||
1. Long-τ CfC channels produce vanishingly small `h` values during K=32 training.
|
||||
2. `heads_w_skip[head, i]` (head's per-channel projection) receives gradient `grad_logit × h[i] ≈ 0` for those channels. **The per-horizon head weights through long-τ channels are effectively frozen at initialization.**
|
||||
3. `τ[i]` itself receives gradient `d_decay × decay × dt / τ_eff²` where `d_decay = grad_h_new × (h_old − tanh(pre))`. With both h_old and the gradient signal near zero, **training does not push long-τ channels toward shorter τ either.**
|
||||
4. At inference (sequential `forward_step_into` over thousands of events), long-τ channels finally integrate and produce meaningful state values. But the heads never learned to use them, so those channels contribute essentially DC noise that is the same across all horizon heads.
|
||||
5. Per-horizon head outputs share dynamics dictated by the short-τ channels that DID train. Inference-time `mean_run_len` is therefore roughly the same across all horizons. **No differentiation.**
|
||||
|
||||
This mechanism is consistent with all observed phenomena: training-time AUC differentiation (heads' weights to short-τ channels do differ per horizon), training-time output stratification (intervention A nudges training-time outputs but the underlying state is shared), and the inference-time `mean_run_len` collapse.
|
||||
|
||||
### 1.3 The fix family
|
||||
|
||||
For inference-time `mean_run_len` to differ per horizon, each head's output must read from state with structurally different time constants. Three families of fix exist:
|
||||
|
||||
| Family | Mechanism | Verdict |
|
||||
|---|---|---|
|
||||
| A: shorten τ to fit K=32 | All CfC channels become fast; heads project from uniform short-τ state | **Fails the WIN gate** — does not produce multi-scale dynamics; only fixes dead-channel waste |
|
||||
| B: K-window inference reset | Training and inference geometries match | **Fails the WIN gate** — geometries match but all heads still share dynamics |
|
||||
| **C: route per-horizon heads via channels sorted by `CfC.tau`** | The existing per-channel CfC τ vector (`crates/ml-alpha/src/cfc/trunk.rs:138`, init log-uniform [0.01, 1000] per line 204) IS the per-channel multi-scale signal; sort by it, bucket, route heads | **Approach taken** |
|
||||
|
||||
**Bucketing source: CfC.tau, not Mamba2 A.** Verified against code (`feedback_trust_code_not_docs`):
|
||||
- `crates/ml-alpha/src/cfc/trunk.rs:138` — `tau_d: CudaSlice<f32>` is per-hidden-channel of size `HIDDEN_DIM=128`. **Per-channel, trained, multi-scale** (log-uniform init spans 5 decades).
|
||||
- `crates/ml-alpha/cuda/mamba2_alpha_kernel.cu:57-89` — Mamba2's `w_a` is shape `[hidden_dim, state_dim ≤ 32]` and the per-step gate is `sigmoid(a_proj[i,t,s])` with `a_proj = x @ W_a.T + b_a`. **The gate is input-dependent, recomputed every sample × every step. No per-channel intrinsic A_diag exists.**
|
||||
|
||||
CfC.tau is the natural bucketing source because the §1.2 decay math literally derives from its per-channel values. The mechanism that breaks differentiation (long-τ channels are dead in K=32 training) is the same mechanism that gives us the bucketing signal (τ-spectrum is multi-scale by construction).
|
||||
|
||||
Empirical anchor: `pearl_state_amplifies_short_horizon_into_long_horizon` showed Mamba2 SSM state lifts AUC 0.50→0.66 at K=6000. CfC's contribution (0.66 → 0.745) is a per-channel-τ-filtered smoother on top — confirming CfC.tau-based bucketing is reading the right signal.
|
||||
|
||||
---
|
||||
|
||||
## 2. Architecture
|
||||
|
||||
### 2.1 Forward flow
|
||||
|
||||
```
|
||||
snap_feature → vsn → mamba2_l1 → ln_a → mamba2_l2 → ln_b → attn_pool
|
||||
│
|
||||
▼ [B, HIDDEN_DIM]
|
||||
┌───────────────────────────────────┴───────────────────────────────────┐
|
||||
│ Phase 1: WARMUP (epochs 0..W, W decided by ISV A-stability EMA) │
|
||||
│ - All channels routed identically (no bucketing) │
|
||||
│ - Single CfC body trains W_in/W_rec/b/τ as before │
|
||||
│ - heads_w_skip/w1/w2/etc. train with full-dim reads │
|
||||
└───────────────────────────────────┬───────────────────────────────────┘
|
||||
▼
|
||||
┌───────────────────────────────────┴───────────────────────────────────┐
|
||||
│ Phase 1→2 TRANSITION (one batch, ISV-triggered, ALL ON DEVICE): │
|
||||
│ - GPU bitonic-sort cfc.tau_d → sorted_channel_indices_d │
|
||||
│ - GPU kernel: bucket_id_per_channel from static quintile bounds │
|
||||
│ - GPU kernel: tau_all_d reorder; bucket IQRs for Controller B │
|
||||
│ - GPU kernel: heads_w_skip_d ragged-compact reorder │
|
||||
│ - Freeze: bucket-channel-map; recapture CUDA graph │
|
||||
└───────────────────────────────────┬───────────────────────────────────┘
|
||||
▼
|
||||
┌───────────────────────────────────┴───────────────────────────────────┐
|
||||
│ Phase 2: ROUTED (epochs W..end): │
|
||||
│ │
|
||||
│ attn_pool out ────────► bucket_split (5 buckets by channel) │
|
||||
│ │ │
|
||||
│ ├─► CfC[shared W,b] τ=τ_0 ──► head_h30 ─► logit_h30
|
||||
│ ├─► CfC[shared W,b] τ=τ_1 ──► head_h100 ─► logit_h100
|
||||
│ ├─► CfC[shared W,b] τ=τ_2 ──► head_h300 ─► logit_h300
|
||||
│ ├─► CfC[shared W,b] τ=τ_3 ──► head_h1000 ─► logit_h1000
|
||||
│ └─► CfC[shared W,b] τ=τ_4 ──► head_h6000 ─► logit_h6000
|
||||
│ │
|
||||
│ Each branch carries its own h_state across forward_step_into calls │
|
||||
│ Heads are block-diagonal: head_h_k reads ONLY bucket_k channels │
|
||||
└────────────────────────────────────────────────────────────────────────┘
|
||||
```
|
||||
|
||||
### 2.2 Shapes and parameters
|
||||
|
||||
- `HIDDEN_DIM = 128` (current value, no widening). Confirmed at `crates/ml-alpha/src/heads.rs:20`.
|
||||
- `N_HORIZONS = 5` (matches existing per-horizon heads).
|
||||
- **Bucket-channel-count vector**: `bucket_dim_k = [25, 25, 25, 25, 28]` (sum = 128 = HIDDEN_DIM). Last bucket absorbs the remainder because HIDDEN_DIM is not divisible by 5.
|
||||
- **MAX_BUCKET_DIM = 28** (used for uniform CUDA block shape; idle threads in shorter buckets handled by uniform predicate; see §5.4 point 1).
|
||||
- CfC body weights (`W_in: [HIDDEN_DIM, HIDDEN_DIM]`, `W_rec: [HIDDEN_DIM, HIDDEN_DIM]`, `b: [HIDDEN_DIM]`) are **shared across all 5 branches**.
|
||||
- `tau_all_d: CudaSlice<f32>` total size `sum(bucket_dim_k) = 128`; addressed via `bucket_channel_offset[k]..bucket_channel_offset[k] + bucket_dim_k[k]`.
|
||||
- Per-branch h_state: `h_state_all_d: CudaSlice<f32>` of size `n_batch × sum(bucket_dim_k) = n_batch × 128`, same addressing pattern.
|
||||
- `heads_w_skip_d`: compact ragged storage of size `sum(bucket_dim_k) = 128` (each horizon reads only its bucket; total bytes = `128 × sizeof(f32) = 512` per ragged column, NOT `[5 × 128 = 640]`). Storage is **5× reduced** from the original block-diagonal-with-zeros layout.
|
||||
- `heads_w_skip_offset: [N_HORIZONS + 1]` — start+end indices of each horizon's compact weights.
|
||||
- `bucket_channel_offset: [N_HORIZONS + 1] -> u32` — start+end indices into HIDDEN_DIM-ordered channel space. Device-resident, frozen at transition. Last entry = HIDDEN_DIM for end-cap.
|
||||
- `bucket_id_per_channel: [HIDDEN_DIM] -> u8` — inverse map, device-resident, also frozen.
|
||||
|
||||
Total new device-resident metadata: `(N_HORIZONS + 1) * 4 + HIDDEN_DIM = 24 + 128 = 152` bytes — negligible.
|
||||
|
||||
### 2.3 Two-phase training
|
||||
|
||||
**Phase 1 (warmup, epochs 0..W):**
|
||||
|
||||
Identical to the current architecture: single CfC body operating on full HIDDEN_DIM state, full HIDDEN_DIM-wide heads, all backward gradients flow as today. Controller A (Section 3.1) monitors `CfC.tau` drift to detect when the trunk has stabilized enough to derive the bucketing.
|
||||
|
||||
**Phase 1→2 transition (one batch boundary, ISV-triggered, ALL ON DEVICE):**
|
||||
|
||||
Per `feedback_no_htod_htoh_only_mapped_pinned` and `feedback_cpu_is_read_only`, the entire transition happens on the GPU. No CPU sort, no DtoH bulk read of τ. The host only OBSERVES Controller A's scalar trigger signal (single float via mapped-pinned, per the standard host-decision pattern).
|
||||
|
||||
Quintile boundaries are KNOWN constants given `HIDDEN_DIM=128` and `bucket_dim_k=[25,25,25,25,28]`: cumulative-sum offsets are `[0, 25, 50, 75, 100, 128]`. Static device constants; no per-step or per-transition computation needed.
|
||||
|
||||
Device kernels launched at transition (between epochs, OUTSIDE captured graph):
|
||||
|
||||
1. **`tau_sort_kernel`** — sort `cfc.tau_d` ascending. Bitonic-merge sort, single block of 32 threads, 4 passes for HIDDEN_DIM=128 (uniform-block, no host branching). Output: `sorted_channel_indices_d: [HIDDEN_DIM]` device tensor (channel index in sorted-by-τ order).
|
||||
|
||||
2. **`bucket_assign_kernel`** — for each `sorted_position p ∈ [0, HIDDEN_DIM)`: lookup which bucket via static `[0, 25, 50, 75, 100, 128]` boundaries; write `bucket_id_per_channel[sorted_channel_indices_d[p]] = bucket_id`. Single block, HIDDEN_DIM threads.
|
||||
|
||||
3. **`bucket_iqr_kernel`** — for each bucket k, compute median + IQR of τ values within bucket. Block-per-bucket, BUCKET_SIZE threads, warp-shuffle reductions per `feedback_no_atomicadd`. Output: `bucket_tau_median_d: [N_HORIZONS]`, `bucket_tau_iqr_lo_d: [N_HORIZONS]`, `bucket_tau_iqr_hi_d: [N_HORIZONS]`. Used by Controller B as the IQR-widened center.
|
||||
|
||||
4. **`tau_reorder_kernel`** — copy `cfc.tau_d` values to `tau_all_d` in bucket-grouped order. Per the static `bucket_channel_offset = [0, 25, 50, 75, 100, 128]`, channel `c` in bucket `k` at within-bucket-position `p` writes `tau_all_d[bucket_channel_offset[k] + p] = cfc.tau_d[c]`. Preserves per-channel-within-bucket variation; median is metadata, not collapsed value.
|
||||
|
||||
5. **`heads_compact_kernel`** — reorder `heads_w_skip` into ragged-compact `heads_w_skip_d: [HIDDEN_DIM]`. For each bucket k, head_h_k's compact weights read from `existing_heads_w_skip[k * HIDDEN_DIM + channel_indices_in_bucket_k]`. Per the static `bucket_channel_offset`, output offsets are known constants.
|
||||
|
||||
6. **Constants populated on device** (no host-side bulk memcpy): `bucket_channel_offset_d: [N_HORIZONS+1] = [0, 25, 50, 75, 100, 128]`, `bucket_dim_k_d: [N_HORIZONS] = [25, 25, 25, 25, 28]`, `heads_w_skip_offset_d: [N_HORIZONS+1] = [0, 25, 50, 75, 100, 128]`. Two acceptable implementations:
|
||||
- (a) **Compile-time embedded**: kernels reference these as `__constant__` device-side arrays linked into the cubin; no runtime memcpy needed.
|
||||
- (b) **One-time init kernel**: a single device kernel writes the constants into device tensors at program start (before training). Single kernel launch, no host data transfer.
|
||||
Both implementations are equally valid; choice deferred to plan-writing (small detail).
|
||||
|
||||
7. **Seed-scoping** per `pearl_scoped_init_seed_for_reproducibility`: any reallocation of `tau_all_d` / `heads_w_skip_d` device buffers wraps the GPU-side init in `scoped_init_seed(cfg.seed)`. Per-allocation determinism preserved.
|
||||
|
||||
8. **Emit `kernel-step-trace` event** marking transition: the device-side `bucket_iqr_kernel` writes per-bucket median τ values into the existing GPU log ring buffer (mapped-pinned, per the GPU log ring spec). Host drains the ring asynchronously; no synchronous DtoH.
|
||||
|
||||
9. **Recapture the training CUDA Graph** per `pearl_no_host_branches_in_captured_graph` + `pearl_cudarc_disable_event_tracking_for_graph_capture` (bracket capture with disable/enable, no host-malloc, no scalar arg changes after capture).
|
||||
|
||||
The total device-only work at transition: 5 kernel launches + 1 graph recapture. Estimated wall time < 5ms.
|
||||
|
||||
**Phase 2 (routed, epochs W..end):**
|
||||
|
||||
Per-branch τ trains independently. Heads' `w_skip` operates on its bucket's BUCKET_DIM channels. Controllers B, C, D active. All training behavior otherwise identical (random sampling, K=32, intervention-A regularizer per-bucket).
|
||||
|
||||
### 2.4 Inference (`forward_step_into`)
|
||||
|
||||
Phase 2 only — production checkpoints have already-frozen routing.
|
||||
|
||||
```
|
||||
1. Existing pipeline up to attn_pool — unchanged.
|
||||
2. attn_pool output [B, HIDDEN_DIM] is conceptually viewed as N_HORIZONS contiguous buckets via bucket_channel_offset.
|
||||
3. Fused per-branch CfC step kernel (single launch, see Section 5.4 point 1):
|
||||
- Grid: (B, N_HORIZONS, 1)
|
||||
- Block: (MAX_BUCKET_DIM = 28, 1, 1) — uniform block shape; idle threads in shorter buckets (uniform predicate, no warp divergence)
|
||||
- Each block reads its bucket's start/size from `bucket_channel_offset[blockIdx.y]` and `bucket_dim_k[blockIdx.y]`
|
||||
4. Per-branch h_new values write back to h_state_all_d for next step.
|
||||
5. Heads forward (existing kernel, modified to read compact w_skip): produces [B, N_HORIZONS] logits.
|
||||
6. Sigmoid → probs → device-resident alpha_probs_d as today.
|
||||
```
|
||||
|
||||
### 2.5 Atomic refactor: MTER scaffolding removal
|
||||
|
||||
The following are deleted in the same commit chain as the per-horizon CfC implementation:
|
||||
|
||||
- `LoaderMode::Sequential` enum variant
|
||||
- `next_sequence_sequential` function
|
||||
- `sequential_file_idx`, `sequential_anchor_idx`, `last_call_was_file_boundary` cursor fields in MultiHorizonLoader
|
||||
- `notify_file_boundary` method on PerceptionTrainer
|
||||
- `last_seen_file_boundary` field
|
||||
- `train_graph_boundary_state` field and its recapture-on-boundary-change logic
|
||||
- `--loader-mode` CLI flag
|
||||
- `loader-mode` Argo template parameter
|
||||
|
||||
The transition-boundary CUDA Graph recapture infrastructure (`train_graph_boundary_state` mechanism) is **redirected** to the Phase 1→2 transition rather than file boundaries. Pattern preserved; trigger condition changed.
|
||||
|
||||
---
|
||||
|
||||
## 3. ISV-driven controllers
|
||||
|
||||
Per `pearl_controller_anchors_isv_driven` and `feedback_isv_for_adaptive_bounds`, every adaptive threshold or bound in the control loop derives from a signal. No hardcoded constants.
|
||||
|
||||
### 3.1 Controller A — Warmup-completion detector
|
||||
|
||||
Decides when CfC.tau has stabilized enough to derive the bucketing.
|
||||
|
||||
**Memory contract**: per `feedback_cpu_is_read_only` and `feedback_no_htod_htoh_only_mapped_pinned`, all signal computation happens on device. The host's role is to read a SINGLE FLOAT scalar via mapped-pinned (the `tau_change_ema` value at end of step) and compare against `noise_floor / 4` (a host-local scalar derived from the first-observation read). This is the standard host-decision pattern; no bulk data moves CPU↔GPU.
|
||||
|
||||
```
|
||||
Per-step signal (all on device):
|
||||
tau_change[t] = ‖cfc.tau_t − cfc.tau_{t-1}‖_F² // single GPU kernel: vector diff + Frobenius reduce → single float on device
|
||||
tau_change_ema = wiener_α_ema(tau_change[t]) // device-side EMA update; result is a single float on device
|
||||
α = diff_var / (diff_var + sample_var + ε) // Wiener-α, all device-side
|
||||
α = max(α, 0.4) // floor for non-stationary, per pearl_wiener_alpha_floor_for_nonstationary
|
||||
|
||||
Anchor (first-observation bootstrap, per pearl_first_observation_bootstrap):
|
||||
At t=0: device kernel reads tau_change[0] and writes noise_floor[k=0] to a mapped-pinned scalar slot.
|
||||
Host reads the mapped-pinned slot ONCE at t=0 to keep noise_floor as a host-local f32 (read-only).
|
||||
Subsequent steps NEVER re-read τ from device; only tau_change_ema scalar is read.
|
||||
|
||||
Trigger condition (host-side scalar compare):
|
||||
Per step, host reads tau_change_ema scalar from mapped-pinned device-visible slot (no DtoH; mapped-pinned IS the device-visible buffer).
|
||||
if tau_change_ema < max(noise_floor / 4, noise_floor / 16):
|
||||
trigger Phase 1→2 transition.
|
||||
// SNR margin derivation: 4 = standard SNR margin for "signal has decayed below detection" (4× below initial noise floor matches SP16 controller idiom for stabilized signals); 16 = 4² safety floor preventing degenerate near-zero noise_floor from collapsing the threshold to 0. Both factors anchored to the same first-obs noise_floor — they encode "4× and 16× below observed initial dispersion", not arbitrary constants.
|
||||
// max-with-floor per pearl_blend_formulas_must_have_permanent_floor
|
||||
|
||||
Safety cap (ISV-anchored, derived from observed dispersion):
|
||||
half_life = first step at which tau_change drops below noise_floor / 2 (observed online; first-obs bootstrap if no crossing within first 100 steps).
|
||||
hard_cap_warmup_steps = half_life × (1 + MAD(tau_change in first 100 steps) / median(tau_change in first 100 steps))
|
||||
// ISV-anchored cap: more dispersion in observed signal → larger cap multiplier
|
||||
// Replaces hardcoded 1.5× with a signal-derived dispersion factor
|
||||
Computed on device, tracked via device-side step counter; host reads single counter scalar via mapped-pinned for the safety check.
|
||||
```
|
||||
|
||||
When triggered → execute Phase 1→2 transition (Section 2.3, all-on-device).
|
||||
|
||||
**Controllers B, C, D, E follow the same memory contract**: signals computed on device, single-scalar mapped-pinned reads for host-side trigger decisions only. Controllers' state (EMAs, accumulators) is device-resident throughout training; no bulk DtoH ever occurs.
|
||||
|
||||
### 3.2 Controller B — Per-bucket τ bound (post-Adam projection)
|
||||
|
||||
Keeps each branch's τ within its bucket's empirical range during Phase 2.
|
||||
|
||||
**Why not loss-weight modulation**: per `pearl_adam_normalizes_loss_weights`, adding `loss += λ * τ_drift` would be a no-op — `tau_all_d` is its own parameter group and Adam normalizes m/sqrt(v). A K× loss-weight lift produces K× m, K²× v, identical step ratio.
|
||||
|
||||
**Mechanism**: hard projection AFTER each Adam step on the τ slice. This bypasses Adam's normalization entirely.
|
||||
|
||||
```
|
||||
At transition time, the all-on-device kernels (§2.3 step 3) produce:
|
||||
bucket_tau_q1_d[k], bucket_tau_q3_d[k] // per-bucket IQR, device-resident
|
||||
slack_factor_d[k] = derived per-bucket from observed Q3/Q1 ratio:
|
||||
slack_factor_k = sqrt(Q3_k / Q1_k)
|
||||
// ISV-anchored: wider IQR gets less additional slack; narrow IQR gets more
|
||||
// Avoids hardcoded 1.5; the slack is proportional to bucket's natural log-spread
|
||||
IQR_widened_lo_d[k] = bucket_tau_q1_d[k] / slack_factor_d[k]
|
||||
IQR_widened_hi_d[k] = bucket_tau_q3_d[k] * slack_factor_d[k]
|
||||
|
||||
After every Adam step that updates tau_all_d (single fused clamp kernel, device-resident bounds):
|
||||
for each channel c ∈ [0, HIDDEN_DIM):
|
||||
k = bucket_id_per_channel[c]
|
||||
tau_all_d[c] = clamp(tau_all_d[c], IQR_widened_lo_d[k], IQR_widened_hi_d[k])
|
||||
```
|
||||
|
||||
Hard projection on a single param group. Gradient signal can still drift τ within the widened IQR each step. Outside the IQR, Adam's m/v continue to accumulate (state preserved for when gradient pushes back inward) but the projection clips the visible parameter.
|
||||
|
||||
If a bucket's τ values consistently want to drift outside the IQR for >`tau_change_half_life` consecutive steps, that signals a Controller D dead-bucket condition or a Mamba2-state-shift; logged via kernel-step-trace, treated as a soft warning (not panic).
|
||||
|
||||
### 3.3 Controller C — Per-bucket smoothness (per-bucket LR multiplier)
|
||||
|
||||
Extends intervention A's existing single-λ ISV controller to 5 independent per-bucket regulators.
|
||||
|
||||
**Why not loss-weight modulation**: same Adam-normalization no-op critique as Controller B. The smoothness loss for bucket k backpropagates into `head_h_k` weight slice — its own parameter group. λ-weight on the loss term cancels under Adam's normalization.
|
||||
|
||||
**Mechanism**: per-bucket LR multiplier on the head_h_k parameter group's Adam step.
|
||||
|
||||
```
|
||||
Per bucket k, per training step:
|
||||
jitter_k = K-window position-variance of head_h_k output // existing intervention-A signal, scoped to bucket
|
||||
jitter_ema_k = wiener_α_ema(jitter_k)
|
||||
α = max(α, 0.4) // floor per pearl_wiener_alpha_floor_for_nonstationary (this IS a co-adapting closed loop)
|
||||
|
||||
target_jitter_k = sqrt(realised_label_variance_k) // ISV anchor from per-horizon label distribution; sentinel-bootstrap at step 0 per pearl_first_observation_bootstrap
|
||||
|
||||
smoothness_ratio_k = max(0, jitter_ema_k − target_jitter_k) / max(target_jitter_k, ε_floor)
|
||||
// signal-driven, max-with-floor, ε_floor = target_jitter_k_at_first_observation / 16
|
||||
|
||||
// LR multiplier per head_h_k parameter group:
|
||||
lr_multiplier_h_k = 1.0 / (1.0 + smoothness_ratio_k) // multiplicative attenuation when jitter exceeds target
|
||||
// Bounded in (0, 1]: when jitter is in spec, lr_mult=1; when far over, lr_mult→0
|
||||
|
||||
// Smoothness loss term WITHOUT λ-weight (just unit weight, because the LR multiplier provides the regulation):
|
||||
loss += smoothness_loss_k
|
||||
```
|
||||
|
||||
Where `smoothness_loss_k` is the existing K-window output-smoothness loss from intervention A, evaluated on horizon-k bucket's head output. The kernel is `output_smoothness.cu` called 5 times with different (bucket-id, head-output-slice) arguments. Each loss term contributes its gradient at unit weight; the per-bucket LR multiplier on the Adam step regulates the response amplitude — escaping Adam's m/sqrt(v) normalization.
|
||||
|
||||
Each loss term has exactly one unbounded multiplicand per `pearl_one_unbounded_signal_per_reward`.
|
||||
|
||||
### 3.4 Controller D — Dead-bucket detector (recovery-first, inheritance-last)
|
||||
|
||||
Post-warmup safety: detect if any bucket's CfC state never integrates. Per `feedback_no_functionality_removal`, recovery is attempted BEFORE the bucket is collapsed to a neighbor's dynamics.
|
||||
|
||||
```
|
||||
Per bucket k, every step in Phase 2:
|
||||
h_mag_k = mean(|h_state_branch_k|) across batch and bucket_dim_k channels
|
||||
h_mag_k_ema = wiener_α_ema(h_mag_k)
|
||||
α = max(α, 0.4) // co-adapting closed loop, floor per pearl_wiener_alpha_floor_for_nonstationary
|
||||
|
||||
First-observation bootstrap (NOT deferred):
|
||||
At Phase 2 first step: first_observation_floor[k] = h_mag_k[t=transition+1]
|
||||
Replace directly per pearl_first_observation_bootstrap; no waiting period.
|
||||
Note on bucket ordering: buckets are sorted ascending-by-τ (bucket 0 = shortest τ, bucket 4 = longest τ). Per §1.2 mechanism, short-τ channels integrate fast → larger h_mag at first observation; long-τ channels integrate slowly → smaller h_mag at first observation. So first_observation_floor is monotonically decreasing with bucket index k.
|
||||
|
||||
Threshold derivation (ISV-anchored, no inter-bucket coupling):
|
||||
For each bucket k INDEPENDENTLY:
|
||||
dead_threshold[k] = first_observation_floor[k] × decay_relative_floor
|
||||
decay_relative_floor = exp(-1) ≈ 0.368 // 1/e of first-observation magnitude is the standard half-life floor; derived from CfC's own decay formula, not a hardcoded fraction
|
||||
// No inter-bucket comparison; "dead" means bucket-k's own state has shrunk to 1/e of its initial value over the dead_window. This is the natural CfC time-constant criterion.
|
||||
|
||||
Detection window (ISV-anchored):
|
||||
dead_window_k = ceil(2 × half_life of h_mag_k_ema under current Wiener-α)
|
||||
(replaces the hardcoded "50 consecutive steps")
|
||||
|
||||
Dead-bucket condition:
|
||||
h_mag_k_ema < dead_threshold[k] for ≥ dead_window_k consecutive steps
|
||||
|
||||
On dead detection — RECOVERY-FIRST cascade:
|
||||
1. RECOVERY ATTEMPT 1: relax Controller B's slack factor for this bucket by ×2:
|
||||
slack_factor_d[k] := 2 × slack_factor_d[k] // doubles the τ envelope; if τ is being pinched, this releases it
|
||||
// The factor 2 is the natural log-doubling step; previously the slack was sqrt(Q3/Q1) — doubling captures one full log-spread step beyond the bucket's natural range
|
||||
Emit kernel-step-trace event: "bucket_k_recovery_widen"
|
||||
Re-check dead condition for next dead_window_k steps.
|
||||
|
||||
2. RECOVERY ATTEMPT 2 (if Attempt 1 fails): reassign N_swap channels from the most-active neighbor
|
||||
bucket into this dead bucket (channels with highest h_mag move).
|
||||
N_swap = ceil(bucket_dim_k * (1 - h_mag_k_ema / first_observation_floor[k])) // ISV-derived from how-far-below-floor the bucket has decayed
|
||||
// 0 channels swap if EMA still at floor; bucket_dim_k channels swap if EMA = 0
|
||||
// Replaces hardcoded "5 channels"
|
||||
N_swap = clamp(N_swap, 1, bucket_dim_k - 1) // ensure at least 1 swap, at most bucket_dim_k-1 (preserve target bucket structure)
|
||||
Recompute bucket_channel_offset and bucket_id_per_channel; recapture CUDA graph.
|
||||
Emit kernel-step-trace event: "bucket_k_recovery_channel_swap" with N_swap value
|
||||
Re-check dead condition for next dead_window_k steps.
|
||||
|
||||
3. INHERITANCE FALLBACK (only after both recoveries fail):
|
||||
Inherit τ from neighbor (k-1) or (k+1) — pick the neighbor with closer label_variance to bucket k.
|
||||
tau_all_d[k slice] := tau_all_d[neighbor slice]
|
||||
Emit kernel-step-trace event: "bucket_k_inheritance_fallback" (this is a degraded outcome, logged distinctly).
|
||||
Continue training.
|
||||
|
||||
After any of the 3 paths, reset Controller D state for bucket k (avoid repeated firing).
|
||||
```
|
||||
|
||||
Recovery-first satisfies `feedback_no_functionality_removal`: the bucket's architectural slot is preserved through two corrective attempts before degraded inheritance is permitted.
|
||||
|
||||
If inheritance is reached, the trading layer's `WeightedByRealizedSharpe` ensemble will detect correlated leaf Sharpe between the dead bucket and its neighbor and downweight the redundant one — graceful degradation, NOT silent collapse.
|
||||
|
||||
### 3.5 Controller E — Inference-time bucket validation (post-checkpoint)
|
||||
|
||||
Diagnostic only; runs in `crt-diag` tool, not in production hot path.
|
||||
|
||||
```
|
||||
At inference (offline diagnostic only), per bucket k:
|
||||
output_run_len_k_ema = wiener_α_ema(run length of sign(head_h_k_logit) being constant)
|
||||
// head_h_k emits a scalar logit per step; "run length" = number of consecutive steps with sign(logit_k - 0.5) unchanged
|
||||
// This is the metric CRT.diag already computes (harness.rs:374 `sum_run_length[b * n_horizons + h]`)
|
||||
τ_assigned_k = median tau across branch k's bucket_dim_k channels
|
||||
|
||||
Cross-check:
|
||||
Spearman rank correlation of (τ_assigned_k for k in 0..5) and (output_run_len_k_ema for k in 0..5)
|
||||
|
||||
Threshold (statistically anchored, not hardcoded):
|
||||
Use Spearman ρ, not Pearson — we only require monotonic ordering, not linear scaling.
|
||||
For N=5 ordered pairs, the threshold for "strictly increasing within sampling noise" is
|
||||
ρ ≥ (N−1)/N = 0.8 // exactly one out-of-order adjacent pair is the natural single-sample slack
|
||||
This derives from the discrete rank-correlation structure for N=5, not a hardcoded constant.
|
||||
|
||||
If ρ < 0.8:
|
||||
Warning logged; treated as a soft failure of the inference-differentiation gate.
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 4. Validation gates
|
||||
|
||||
The design is validated against three smokes, in order. Each smoke gates the next.
|
||||
|
||||
### 4.1 Smoke 1: training-time stability (~20 min L40S)
|
||||
|
||||
Argo run, random sampling, 20 epochs, no early-stop, default seed 16962.
|
||||
|
||||
| Quantity | Gate | Source |
|
||||
|---|---|---|
|
||||
| `val_loss` stays within 1.5× of Phase 1 peak across all 20 epochs | required | rules out catastrophic divergence like the sequential failure (val_loss 3 → 18) |
|
||||
| Best per-horizon `val_AUC` (h6000) | ≥ `chance + (random_baseline_h6000_auc − chance) × (1 − regression_tolerance)` where `regression_tolerance` derives from per-seed variance of random baseline (run a 3-seed baseline sweep BEFORE this smoke to anchor) | invariant-based per `pearl_tests_must_prove_not_lock_observations` — assert that the architecture didn't materially regress, not a fixed 0.72 threshold |
|
||||
| Controller A transition triggered | within ISV-derived hard cap (Section 3.1) | if cap hit before signal-driven trigger, surfaces as "soft" outcome; Smoke 2 may still pass with imperfect routing |
|
||||
| Controller D dead-bucket detector fires | 0 occurrences ideal; recovery-attempt-1 firing acceptable; recovery-attempt-2 firing soft warning; inheritance-fallback distinct outcome (logged, counted) | recovery cascade in §3.4 means each level is a different soft-outcome classification |
|
||||
| Controller B post-Adam projection clips | bounded; per-bucket fraction of params clipped per step < 50% sustained | confirms IQR widening matches training data — if half the params want to live outside the IQR sustained, that's a signal the bucketing source was wrong |
|
||||
|
||||
### 4.2 Smoke 2: inference-time differentiation (~10 min, CRT.diag on Smoke 1 checkpoint)
|
||||
|
||||
The actual WIN gate. Uses existing `crt-diag` tooling with the Phase 2 checkpoint.
|
||||
|
||||
| Quantity | Gate | Source |
|
||||
|---|---|---|
|
||||
| `mean_run_len[h6000] / mean_run_len[h30]` | **≥ 10×** | original WIN criterion |
|
||||
| `mean_run_len[h6000]` absolute | **≥ 100 events** | original WIN criterion |
|
||||
| Monotonic ordering: `mean_run_len[h30] < h100 < h300 < h1000 < h6000` | strictly increasing | confirms 5 distinct dynamics |
|
||||
| Controller E Spearman ρ: assigned τ ↔ observed run-length | ≥ 0.8 (Spearman, threshold derived from N=5 single-pair slack) | confirms multi-scale dynamics flow through to outputs |
|
||||
|
||||
### 4.3 Smoke 3: trading-sim end-to-end (~5-10 min)
|
||||
|
||||
Run `fxt-backtest` with the Phase 2 checkpoint against existing test scenarios.
|
||||
|
||||
| Quantity | Gate | Source |
|
||||
|---|---|---|
|
||||
| `fxt-backtest` completes without panic / NaN / inf-cascade | required | basic robustness check |
|
||||
| Per-leaf max position size | ≤ existing leaf bound | confirms Kelly sizing doesn't blow under new dynamics |
|
||||
| `WeightedByRealizedSharpe` ensemble weights at end | non-degenerate (no leaf at 100%) | confirms 5 useful signals, not 1+4 redundant |
|
||||
| Total PnL vs current-architecture baseline | ≥ 0.9× baseline (soft); stretch goal ≥ baseline | rules out regression in trading performance |
|
||||
| Per-leaf trade frequency: h30 leaf trades >> h6000 leaf trades | qualitative | confirms differentiation flows through to trading behavior |
|
||||
|
||||
### 4.4 Failure outcomes
|
||||
|
||||
Each row below corresponds to a distinct mechanism with a distinct next-step (per `feedback_kill_runs_on_anomaly_quickly`, gates must discriminate mechanisms).
|
||||
|
||||
| Outcome | Distinct mechanism | Next step (distinct) |
|
||||
|---|---|---|
|
||||
| Smoke 1 ✓ Smoke 2 ✓ Smoke 3 ✓ | DESIGN VALIDATED | Document pearl `pearl_per_horizon_channel_routing_via_cfc_tau`. Set up walk-forward and production training. |
|
||||
| Smoke 1 ✓ Smoke 2 fails ratio gate AND Controller A triggered cleanly | Bucketing source produces τ ordering BUT outputs don't differentiate | Dump per-bucket head_h_k output histograms. Likely cause: shared CfC body (per §5.8 out-of-scope item) is too restrictive — per-bucket independent W_in/W_rec/b becomes the next-spec scope. |
|
||||
| Smoke 1 ✓ Smoke 2 fails ratio gate AND Controller A hit hard cap | CfC.tau distribution hasn't separated into 5 distinct quintiles | Inspect CfC.tau histogram at hard cap. If most channels clustered in narrow band: longer Phase 1 needed (extend hard cap derivation) or τ-init was wrong (revisit log-uniform range). |
|
||||
| Smoke 1 ✓ Smoke 2 ✓ Smoke 3 fails PnL gate | Trading economic objective diverges from WIN criterion | SERIOUS finding requiring follow-up design. NOT code rollback; trade-off must be made explicit; user decision required. |
|
||||
| Smoke 1 fails val_AUC invariant gate | Bucket split + block-diagonal head mask costs AUC at training | Investigate: per-bucket head_h_k can't see enough channels to learn. Either widen bucket coverage (heads read OWN bucket + neighbors with soft mask) or revert to single-CfC. |
|
||||
| Smoke 1 fails val_loss stability (divergence cascade) | Controllers B/C/D are interacting destructively | Disable Controllers B and C (use plain Adam); re-run. If stable: re-design the per-bucket regulators. If not: deeper bug. |
|
||||
| Controller D recovery-attempt-1 fires (IQR widen) | Single bucket's τ values converging outside initial IQR | Likely benign: bucketing was conservative. Soft success if Smoke 2 + Smoke 3 still pass. |
|
||||
| Controller D recovery-attempt-2 fires (channel swap) | One bucket's CfC.tau-sorted channels weren't actually multi-scale | Bucketing source partially failed at the chosen bucket; investigate that quintile's channels post-hoc. |
|
||||
| Controller D inheritance-fallback fires | Both recoveries failed for a bucket | Logged as degraded outcome. Smoke 3 may still pass with ensemble rebalancing. If >1 bucket hits inheritance: bucketing source is fundamentally wrong — re-brainstorm.|
|
||||
|
||||
---
|
||||
|
||||
## 5. Implementation scope and performance
|
||||
|
||||
Per `feedback_no_partial_refactor` and `pearl_no_deferrals_for_complementary_fixes`, the per-horizon CfC architecture lands in one commit chain with MTER scaffolding atomically removed.
|
||||
|
||||
### 5.1 Files added (4)
|
||||
|
||||
- `crates/ml-alpha/src/cfc/bucket_routing.rs` — Controller A logic, freeze metadata, transition orchestration (kernel dispatch + graph recapture)
|
||||
- `crates/ml-alpha/cuda/cfc_step_per_branch.cu` — **single fused kernel** covering all 5 branches × n_batch in one launch
|
||||
- `crates/ml-alpha/cuda/heads_block_diagonal_fwd.cu` — heads forward with compact block-diagonal `w_skip` storage and bucket offset lookup
|
||||
- `crates/ml-alpha/cuda/bucket_transition_kernels.cu` — **5 device kernels for the Phase 1→2 transition** (all-on-device per `feedback_no_htod_htoh_only_mapped_pinned`): `tau_sort_kernel` (bitonic-merge, 1 block × 32 threads × 4 passes for HIDDEN_DIM=128), `bucket_assign_kernel` (HIDDEN_DIM threads, write `bucket_id_per_channel`), `bucket_iqr_kernel` (block-per-bucket median + IQR via warp-shuffle reduction), `tau_reorder_kernel` (channel-reorder tau into bucket-grouped `tau_all_d`), `heads_compact_kernel` (ragged-compact reorder of `heads_w_skip`)
|
||||
|
||||
### 5.2 Files modified (8)
|
||||
|
||||
- `crates/ml-alpha/src/cfc/trunk.rs` — `tau_d` replaced by `tau_all_d: CudaSlice<f32>` of size `sum(bucket_dim_k) = HIDDEN_DIM = 128`; `bucket_channel_offset: [N_HORIZONS+1]` (last entry = HIDDEN_DIM for end-cap); `bucket_dim_k: [N_HORIZONS]` (sizes per bucket — [25, 25, 25, 25, 28]). Checkpoint envelope adds these three vectors + per-channel τ values laid out compactly. Bump Checkpoint envelope (greenfield, no backward compat).
|
||||
- **Checkpoint envelope consumer audit (3 sites, all migrated atomically in this commit chain)**:
|
||||
1. `crates/ml-alpha/src/cfc/trunk.rs:30,334,398` — `Checkpoint` struct definition + save_checkpoint + load_checkpoint (the producer)
|
||||
2. `crates/ml-alpha/src/trainer/perception.rs:3826` — `CfcTrunk::load_checkpoint` consumer in `from_checkpoint`
|
||||
3. `crates/ml-backtesting/tests/checkpoint_smoke.rs:34-44` — smoke test that exercises load_checkpoint
|
||||
- No additional consumers found via `grep -rn "load_checkpoint\|CfcTrunk" --include="*.rs"`. Audit complete.
|
||||
- `crates/ml-alpha/src/cfc/step.rs` — binding for the per-branch fused kernel
|
||||
- `crates/ml-alpha/src/heads.rs` — block-diagonal mask on `heads_w_skip`; storage compacted from `[N_HORIZONS × HIDDEN_DIM]` to ragged `[sum(bucket_dim_k) = HIDDEN_DIM]` via `heads_w_skip_offset`; optimizer Adam moments shrink correspondingly
|
||||
- `crates/ml-alpha/src/trainer/perception.rs` — two-phase loop (warmup → transition → routed); Controllers A/B/C/D wired into `step_batched`; `forward_step_into` 5-branch fused CfC call + heads with compact w_skip
|
||||
- `crates/ml-alpha/src/data/loader.rs` — REMOVE `LoaderMode::Sequential`, `next_sequence_sequential`, `sequential_file_idx`, `sequential_anchor_idx`, `last_call_was_file_boundary`. `LoaderMode` reduces to a single-variant enum that may be deleted.
|
||||
- `crates/ml-alpha/examples/alpha_train.rs` — REMOVE `--loader-mode` CLI flag and `notify_file_boundary()` call. Add `--bucket-warmup-cap-steps` flag (default: ISV-derived from first 100 steps; CLI override for diagnostic purposes only).
|
||||
- `infra/k8s/argo/alpha-perception-template.yaml` — REMOVE `loader-mode` workflow param. Optionally add `bucket-warmup-cap-steps` param (default unset → ISV-derived).
|
||||
- `crates/ml-alpha/cuda/output_smoothness.cu` — intervention A regularizer scoped per-bucket. Same kernel, called 5× with different (bucket-id, head-output-slice) arguments, contributing to per-bucket jitter signal for Controller C.
|
||||
|
||||
### 5.3 Files deleted (0)
|
||||
|
||||
The existing MTER spec and plan documents (`docs/superpowers/specs/2026-05-21-crt-train-intervention-b-multi-timescale-readout.md` and corresponding plan) remain as historical record. They are not in the build path.
|
||||
|
||||
### 5.4 Performance optimizations (in scope from day one)
|
||||
|
||||
NVIDIA-grade kernel design is part of correctness here, not a separate optimization pass.
|
||||
|
||||
1. **Single fused CfC kernel for all 5 branches** (uniform-block-shape pattern)
|
||||
- `HIDDEN_DIM = 128` is not divisible by 5; bucket sizes are `[25, 25, 25, 25, 28]`. To keep a single fused launch with a uniform block shape:
|
||||
- Launch: `grid = (n_batch, N_HORIZONS, 1)`, `block = (MAX_BUCKET_DIM = 28, 1, 1)`
|
||||
- Each block reads its bucket's start offset from `bucket_channel_offset[blockIdx.y]` and its size from `bucket_dim_k[blockIdx.y]`
|
||||
- Threads with `threadIdx.x >= bucket_dim_k[blockIdx.y]` early-return (3 idle threads in 4 of 5 buckets — negligible). No branch divergence within a warp since the early-return is a uniform comparison against a per-block constant.
|
||||
- Each block addresses `tau_all_d[bucket_channel_offset[blockIdx.y] + threadIdx.x]`
|
||||
- Shared `W_in`/`W_rec`/`b` cooperative-staged into shared memory once per block, per `pearl_cooperative_staging_eliminates_redundant_reads`
|
||||
|
||||
2. **Compact block-diagonal heads storage** (ragged layout for uneven bucket sizes)
|
||||
- `heads_w_skip` was `[N_HORIZONS × HIDDEN_DIM]` = `[5 × 128]` = 640 floats with off-bucket entries forced to zero.
|
||||
- Becomes a ragged compact buffer of size `sum(bucket_dim_k) = 128` floats total — each head_h_k owns only its bucket's contiguous slice.
|
||||
- Storage **5× reduced** (640 → 128 floats); memory bandwidth proportionally reduced.
|
||||
- GRN heads kernel takes `heads_w_skip_offset[N_HORIZONS+1]` device array (start+end), reads from offset.
|
||||
- Zero data movement at hot path — bucket assignment is purely addressing metadata.
|
||||
|
||||
3. **No tensor reshuffling at bucket-split**
|
||||
- `bucket_id_per_channel: [HIDDEN_DIM] -> u8` and `bucket_channel_offset: [N_HORIZONS] -> u32` precomputed at transition, frozen, ~256 bytes device-resident
|
||||
- Kernels use the metadata for addressing — zero data shuffling on the hot path
|
||||
- Per-branch h_state laid out as `[n_batch × sum(bucket_dim_k) = n_batch × HIDDEN_DIM]` single contiguous allocation; per-branch slice via `bucket_channel_offset[k]..bucket_channel_offset[k+1]`, no allocator pressure
|
||||
|
||||
4. **CUDA Graph capture preserved**
|
||||
- Phase 1→2 transition happens OUTSIDE the captured graph (between epochs)
|
||||
- Phase 2 captured graph is monomorphic — no host branches, no scalar-arg changes per step
|
||||
- One-time graph recapture at transition — same boundary-aware pattern as the (now-removed) sequential mode's file boundary recapture, redirected to phase transition
|
||||
- Per `pearl_no_host_branches_in_captured_graph` and `pearl_cudarc_disable_event_tracking_for_graph_capture`
|
||||
|
||||
5. **Coalesced writes in heads forward**
|
||||
- Heads' per-horizon logit outputs already shape `[n_batch × N_HORIZONS]` from existing kernel
|
||||
- Confirm via inspection under new compact w_skip layout; if non-coalesced, swap thread role per `pearl_coalesce_via_thread_role_swap`
|
||||
|
||||
6. **Per-branch backward respects mask without divergence**
|
||||
- Backward CfC kernel: same per-branch fused pattern as forward, single launch
|
||||
- Backward heads: gradient on `heads_w_skip` only fills compact `[sum(bucket_dim_k) = 128]` positions; off-bucket positions DO NOT EXIST in the storage (it's compact, not block-zero-padded). Adam moments live only over the compact storage.
|
||||
- No branch divergence: bucket addressing is uniform within a block
|
||||
|
||||
7. **Adam step kernel reuses existing path + per-bucket LR multipliers (Controllers B and C)**
|
||||
- Per-branch τ values are contiguous slices of single `tau_all_d` device tensor (size HIDDEN_DIM = 128, unchanged from pre-bucketing layout) — same Adam-update kernel as today
|
||||
- Adam moment buffers (`m`, `v`) for `tau_all_d` are single contiguous allocations matched to the param shape
|
||||
- **Per-bucket LR multiplier** (Controller C mechanism): the Adam update kernel takes a per-parameter-group LR multiplier vector. For `head_h_k` weight slices, the multiplier is `lr_multiplier_h_k` from Controller C. For `tau_all_d`, multiplier stays 1.0 (Controller B operates via post-Adam projection, not LR).
|
||||
- **Post-Adam projection** (Controller B mechanism): after `tau_all_d` is updated by Adam, a single clamp kernel applies `tau_all_d[c] = clamp(tau_all_d[c], IQR_widened[bucket_id_per_channel[c]].lo, IQR_widened[bucket_id_per_channel[c]].hi)`. Single launch, threadIdx covers all 128 channels.
|
||||
- **Seed-scoping** (per `pearl_scoped_init_seed_for_reproducibility`): the Phase 1→2 transition's `tau_all_d` and `heads_w_skip_d` reallocations MUST be wrapped in `scoped_init_seed(cfg.seed)`. Any xavier-init or random-init helpers within the transition use this scope.
|
||||
|
||||
8. **`fxt-backtest` inference reuses the same kernels**
|
||||
- `forward_step_into` at deployment uses the SAME fused per-branch step kernel as training
|
||||
- 5 independent h_state buffers (sum = HIDDEN_DIM floats per batch slot) — identical total to the current single-h_state memory footprint
|
||||
- Layout: `h_state_all_d: [n_batch × HIDDEN_DIM]` contiguous; per-branch slices addressed via `bucket_channel_offset`
|
||||
|
||||
9. **Launch-order assertion (per `pearl_canary_input_freshness_launch_order`)**
|
||||
- During Phase 1, NO kernel reads `bucket_id_per_channel`, `bucket_channel_offset`, or any Phase-2-only tensor. Phase 1 captured graph and Phase 2 captured graph are DISJOINT — different kernel pointer sets, different metadata tensors. The transition (between epochs) recaptures from scratch.
|
||||
- During Phase 2, the bucket metadata tensors are read by every per-branch CfC kernel; they're written ONLY at the transition by the device kernels in §2.3 (no HtoD; static constants like `bucket_channel_offset=[0,25,50,75,100,128]` can be either compile-time embedded or initialized once at program start). `bucket_id_per_channel` is the only metadata that depends on actual τ values; it's computed by `bucket_assign_kernel` from `sorted_channel_indices_d` (also device-side).
|
||||
|
||||
### 5.5 Trading-layer integration (no code changes, validation only)
|
||||
|
||||
The trading layer is already structured around per-horizon probs:
|
||||
|
||||
- `ml-backtesting/policy/mod.rs`: 5 leaf strategies indexed by `horizon_idx: u8`; each runs per-horizon ISV-Kelly
|
||||
- `ml-backtesting/sim/mod.rs`: `alpha_probs_d` is `[N_HORIZONS]` device slice — interface preserved
|
||||
- `ml-backtesting/harness.rs`: hardcoded horizon labels `[30, 100, 300, 1000, 6000]` — unchanged
|
||||
- `harness.rs` CRT.diag: already computes `flip_count`, `mean_run_len`, `run_length_hist` per (batch, horizon)
|
||||
- Policy's `WeightedByRealizedSharpe`: adaptive ensemble across the 5 leaves — auto-rebalances when leaf signals genuinely differentiate (which the new architecture provides)
|
||||
|
||||
**Integration risk** is operational, not interface-level: NaN propagation, exposure bound behavior under suddenly-differentiated leaf signals, capital allocation responsiveness. Smoke 3 covers these.
|
||||
|
||||
### 5.6 Memory updates (4 entries in same commit chain)
|
||||
|
||||
- New: `pearl_per_horizon_channel_routing_via_cfc_tau.md`
|
||||
> "Horizon-conditioned channel routing where per-horizon heads read state channels sorted by `CfC.tau` (the existing per-channel CfC time-constant vector). The multi-tau state already exists in CfC; this design structurally binds each horizon's head to a quintile of the tau spectrum, producing inference-time temporal differentiation. Verified: the bucketing source must be a real per-channel signal in the codebase — Mamba2's `A` is input-dependent, not a per-channel rate, so it's NOT the bucketing source despite being a tempting candidate."
|
||||
- Update: `pearl_training_smoothness_does_not_transfer_to_inference.md` — mark mechanism understood, link to this design
|
||||
- Update: `project_mter_commit1_smoke_outcome.md` — link to successor design; mark verdict closed
|
||||
- New: `pearl_trading_layer_designed_for_differentiated_per_horizon_dynamics.md`
|
||||
> "5-horizon trading ensembles with WeightedByRealizedSharpe can absorb genuinely-differentiated per-horizon signals without code change. The bottleneck for inference-time horizon differentiation was always upstream in the alpha generation, not in the trading layer."
|
||||
|
||||
### 5.7 Verification matrix (commit chain complete only when ALL pass)
|
||||
|
||||
- `cargo check --workspace --all-targets` clean
|
||||
- `SQLX_OFFLINE=true cargo test -p ml-alpha --lib` passes
|
||||
- `SQLX_OFFLINE=true cargo test -p ml-backtesting --lib` passes (downstream consumers still green)
|
||||
- Local CUDA smoke on RTX 3050 (required, per `feedback_nvidia_grade_perf_for_kernels` — local sm_86 smoke catches UB before cluster deploy). With bucket sizes [25,25,25,25,28] and HIDDEN_DIM=128 total, memory footprint is trivially below 4GB.
|
||||
- Argo Smoke 1 (training-time stability)
|
||||
- Argo Smoke 2 (CRT.diag inference-time differentiation)
|
||||
- Argo Smoke 3 (fxt-backtest trading sim)
|
||||
|
||||
### 5.8 Explicitly out of scope
|
||||
|
||||
These were considered during brainstorming and deferred. Each would require its own design and is NOT to be included in the per-horizon CfC commit chain.
|
||||
|
||||
- Per-bucket independent W_in/W_rec/b CfC bodies (shared body validated first; future spec if shared body underperforms)
|
||||
- Mamba2 A initialization constraints (let Mamba2 learn whatever it learns; if Phase 1→2 transition reveals A clustering, future spec)
|
||||
- Production walk-forward training of the new architecture (after design validation, separate spec)
|
||||
- H100 training (smokes are on L40S per `feedback_default_to_l40s_pool`; if production needs H100, separate spec)
|
||||
|
||||
---
|
||||
|
||||
## 6. References
|
||||
|
||||
### Memory entries this design depends on
|
||||
|
||||
- `pearl_training_smoothness_does_not_transfer_to_inference` — predecessor pearl; documents the bottleneck
|
||||
- `pearl_random_sampling_sufficient_for_per_horizon_training_auc` — refutes MTER's premise
|
||||
- `pearl_state_amplifies_short_horizon_into_long_horizon` — empirical anchor for Mamba2 multi-scale capacity
|
||||
- `pearl_controller_anchors_isv_driven` — every controller anchor ISV-driven
|
||||
- `pearl_first_observation_bootstrap` — sentinel-bootstrap pattern
|
||||
- `pearl_wiener_optimal_adaptive_alpha` — α derivation
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` — α floor 0.4 in control loops where target drifts
|
||||
- `pearl_blend_formulas_must_have_permanent_floor` — max-with-floor not blend
|
||||
- `pearl_one_unbounded_signal_per_reward` — controller term hygiene
|
||||
- `pearl_no_host_branches_in_captured_graph` — CUDA Graph constraint
|
||||
- `pearl_cooperative_staging_eliminates_redundant_reads` — shared body weights staging
|
||||
- `pearl_coalesce_via_thread_role_swap` — heads kernel output coalescing
|
||||
- `feedback_no_partial_refactor` — atomic migration of consumers
|
||||
- `feedback_isv_for_adaptive_bounds` — bounds derive from ISV
|
||||
- `feedback_default_to_l40s_pool` — smoke on L40S
|
||||
- `project_mter_commit1_smoke_outcome` — context for why MTER is abandoned
|
||||
|
||||
### Code references (verified during spec authorship)
|
||||
|
||||
- `crates/ml-alpha/cuda/cfc_step.cu:57-58` — CfC math `decay = exp(-dt/τ); h_new = h_old·decay + (1-decay)·tanh(pre)`
|
||||
- `crates/ml-alpha/src/cfc/trunk.rs:138` — `tau_d: CudaSlice<f32>` of size HIDDEN_DIM=128 — **the bucketing source**
|
||||
- `crates/ml-alpha/src/cfc/trunk.rs:204-209` — current log-uniform τ init over [10ms, 1000s]
|
||||
- `crates/ml-alpha/src/cfc/trunk.rs:30,334,398` — Checkpoint struct, save_checkpoint, load_checkpoint (envelope producer + consumer)
|
||||
- `crates/ml-alpha/src/trainer/perception.rs:3826` — load_checkpoint consumer in `from_checkpoint`
|
||||
- `crates/ml-backtesting/tests/checkpoint_smoke.rs:34-44` — smoke test exercising load_checkpoint (3rd consumer)
|
||||
- `crates/ml-alpha/cuda/mamba2_alpha_kernel.cu:57-89` — verified Mamba2's input-dependent selective scan gate (NOT a per-channel A_diag)
|
||||
- `crates/ml-alpha/src/mamba2_block.rs:551-554,631-633` — verified `w_a` shape `[hidden_dim, state_dim≤32]`
|
||||
- `crates/ml-alpha/src/trainer/perception.rs:3319` — `forward_step_into` inference path
|
||||
- `crates/ml-alpha/src/trainer/perception.rs:3517` — `dt_s = 1.0` event-rate per CfC step
|
||||
- `crates/ml-alpha/src/heads.rs:20` — `HIDDEN_DIM = 128`
|
||||
- `crates/ml-backtesting/src/policy/mod.rs:23` — per-horizon Kelly leaf structure
|
||||
- `crates/ml-backtesting/src/harness.rs:147` — horizons `[30, 100, 300, 1000, 6000]`
|
||||
|
||||
### Diagnostic Argo workflows
|
||||
|
||||
- `alpha-perception-8gz72` — original MTER commit-1 smoke (sequential, patience=5, false-fail)
|
||||
- `alpha-perception-diag-seq-120533` — 20-epoch sequential no-early-stop diagnostic (revealed catastrophic divergence at epoch 7+, peak h6000=0.745 at epoch 6)
|
||||
- `alpha-perception-diag-rnd-120534` — 20-epoch random control no-early-stop (stable throughout, peak h6000=0.745 at epoch 5)
|
||||
|
||||
---
|
||||
|
||||
## 7. Open questions deferred to plan-writing
|
||||
|
||||
These are detail-level choices that the implementation plan (writing-plans skill) should pin down before Task 1. Q1 (originally about Mamba2's A_diag) is RESOLVED — bucketing source is `CfC.tau`, not Mamba2 A. The remaining questions:
|
||||
|
||||
- **Q-label-var**: Exact form of label-variance signal for Controller C's `target_jitter_k` — likely existing per-horizon BCE label distribution; need to verify the per-bucket scoping is sound and that `sentinel-bootstrap at step 0` lands on a non-degenerate value
|
||||
- **Q-ckpt-tests**: Existing `checkpoint_smoke_tests::save_load_roundtrip_preserves_all_weights` (`cfc/trunk.rs:485`) currently asserts bit-equality of `tau_d`; needs migration to assert bit-equality of `tau_all_d` + `bucket_channel_offset` + `bucket_dim_k` + `bucket_id_per_channel`. Plan must include this test migration.
|
||||
- **Q-smoothness-api**: The existing `output_smoothness.cu` kernel API may or may not support per-bucket invocation natively — need to inspect the kernel signature at plan-writing time and decide between (a) per-bucket wrapper, (b) extending the kernel with a `bucket_offset` argument
|
||||
- **Q-3-seed-baseline**: Before Smoke 1, run a 3-seed baseline sweep on the CURRENT architecture (HEAD f22f3f948) to establish `random_baseline_h6000_auc` mean ± stddev. This anchors the Smoke 1 invariant gate's `regression_tolerance`. ~30 min total wall.
|
||||
- **(RESOLVED in §3.4)**: Controller D Recovery-Attempt-2's channel swap stride is now ISV-derived from `bucket_dim_k * (1 - h_mag_k_ema / first_observation_floor)`. No longer an open question.
|
||||
- **Q-recapture-cost**: CUDA graph recapture cost at Phase 1→2 transition — empirically measure during plan execution. If > 100ms, consider keeping bucket metadata as a graph-input rather than recompiling.
|
||||
Reference in New Issue
Block a user