Adds --per-horizon-logit-sampling-stride CLI flag to fxt-backtest run/sweep
+ harness instrumentation that DtoHs alpha_probs every N events and
accumulates per-horizon mean/stddev/min/max + 10-bin histogram.
Emitted at end of CRT.diag block as:
per_horizon_logit_diag h{N}: count=... mean=... std=... min=... max=... hist=[...]
Goal: distinguish whether Smoke 2's uniform mean_run_len 1.04x across
horizons is caused by uniform per-horizon LOGITS (-> Task 11 scope
incomplete; heads_w1/w2/gate/main need block-diagonal too) or by
DIFFERENTIATED logits with collapsed downstream (-> different fix).
Default stride=None (disabled, zero overhead). With stride=1000 on
2M events, ~2000 DtoD stops x ~50us = ~100ms total overhead per run.
Per feedback_no_htod_htoh_only_mapped_pinned: uses mapped-pinned buffer
(single 5xf32 allocation reused on every sampling tick).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
URGENT correctness fix surfaced by checkpoint deep-dive after Smoke 2 failed
the WIN gate (mean_run_len ratio = 1.0× vs target ≥10×; all 5 horizons
uniformly ~2.4 events).
THREE INDEXING BUGS identified:
1. tau_reorder produces bucket-grouped tau_all_d, but Controller B's
tau_clamp_kernel reads bucket_id_per_channel[c] where c is the
POSITION in the reordered buffer (not the original channel index)
→ wrong bucket-IQR lookup → τ never constrained → buckets collapse
(deep-dive showed all 5 buckets had nearly identical τ ranges
[0.07, 74] except bucket 4 reaching 878).
2. heads_w_skip grad mask ran but Adam (m, v) momentum from BEFORE the
transition re-introduced gradient signal across the transition →
off-bucket positions stayed nonzero throughout training
(deep-dive: 512/512 = 100% of off-bucket positions nonzero in
trunk_best_h6000.bin).
3. per-branch CfC kernel read w_in[c * HIDDEN_DIM + k] with c =
bucket-grouped position, but W_in rows are indexed by ORIGINAL
channel → kernel read wrong rows for each output → outputs were
essentially random per-channel.
ALPHA FIX: skip the reorder entirely, use bucket-filter throughout:
- Removed tau_reorder_kernel; cfc.tau_d stays in original-channel layout.
- Added channels_in_bucket_kernel that populates a
[N_HORIZONS × MAX_BUCKET_DIM] lookup (original channel index per
(bucket, within-bucket-position)).
- Per-branch CfC fwd+bwd now reads channels_in_bucket[branch][tid]
→ original_c, then uses original_c for w_in/w_rec indexing. All
weights stay in original layout consistently.
- Controller B's tau_clamp_kernel now correctly operates on original-
channel cfc.tau_d with bucket_id_per_channel[c] lookup (no position-
vs-channel confusion).
- Added zero_off_bucket_kernel + three-layer defense for heads_w_skip
block-diagonal invariant:
(a) At transition: zero off-bucket params + zero Adam (m, v)
moments via opt_heads_w_skip.m_mut() / v_mut() accessors.
(b) Per-step: heads_w_skip_grad_mask_apply_kernel zeros off-bucket
gradients before Adam step (unchanged from prior follow-up).
(c) Per-step: zero_off_bucket_kernel zeros off-bucket params after
Adam step, catching any drift from Adam's ε denominator or
weight decay.
New AdamW::m_mut()/v_mut() accessors enable the projection at transition.
GPU oracle tests: 19 total (16 in bucket_transition + 3 in cfc_step_per_branch).
New tests verify:
- channels_in_bucket_kernel correctness under non-contiguous bucket assignment
- zero_off_bucket maintains invariant after many mock Adam steps
- fwd kernel writes only to bucket-assigned channels under arbitrary mapping
Per `feedback_no_partial_refactor`: all 3 indexing bugs + Adam momentum
defense land in one commit.
ml-alpha lib: 33 passed.
GPU oracle tests on RTX 3050 sm_86: 19 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
URGENT correctness fix surfaced by Task 13 implementer.
Bug: training-time Phase 2 dispatch sites (cfc_step_per_branch_fwd at
~3139, h_mag_per_bucket for Controller D at ~3211-3212) read
self.trunk.bucket_*_d which are zero-initialized at trunk construction
and only populated via load_checkpoint. The Phase 1→2 transition
populates a SEPARATE BucketRoutingMetadata owned by the trainer
(self.bucket_routing_metadata) but never syncs it into the trunk fields.
With zero-init bucket_dim_k, the per-branch kernel's uniform predicate
`tid >= bucket_dim_k[branch]` early-returns ALL threads -> Phase 2
silently no-ops during training, Smoke 1 fails before producing signal.
Fix: at transition, DtoD-copy metadata buffers into trunk fields BEFORE
the `Some(metadata)` move (so `metadata.*` borrows remain valid):
- metadata.bucket_channel_offset_d -> trunk.bucket_channel_offset_d
- metadata.bucket_dim_k_d -> trunk.bucket_dim_k_d
- metadata.bucket_id_per_channel_d -> trunk.bucket_id_per_channel_d
- metadata.heads_w_skip_offset_d -> trunk.heads_w_skip_offset_d
Trunk fields remain the single source of truth for both training-time
dispatch AND checkpoint serialization. Total sync cost: 4 DtoD memcpys,
~256 bytes total at the one-shot transition (off the hot path).
Verification:
- `cargo check -p ml-alpha --all-targets`: clean
- `cargo test -p ml-alpha --lib`: 33 passed, 0 failed
- 3 read sites (perception.rs:3139, 3211-3212, 4725-4726) verified
unchanged; all 4 new memcpy_dtod_async calls bracketed in the
transition block (perception.rs:2224, 2234, 2244, 2254).
- Block-diagonal heads grad-mask init (step 7) continues to read
`metadata.bucket_id_per_channel_d` via `as_ref()` after the move;
ordering preserved per pearl_canary_input_freshness_launch_order.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §2.4 and Task 13 of the plan.
Replace the single-CfC `cfc_step_batched` dispatch in `forward_step_into`
with `cfc_step_per_branch_fwd_gpu`. Production checkpoints have Phase 2
routing frozen at the training-time Phase 1→2 transition, so inference
unconditionally takes the per-branch path. The fused kernel reads
`bucket_channel_offset_d` / `bucket_dim_k_d` from the trunk; in
deployment these are populated via `CfcTrunk::load_checkpoint`. For
freshly-constructed trainers without a checkpoint load, the trunk fields
are zero — the kernel's uniform predicate (`tid >= bucket_dim_k`) then
early-returns every thread and h_new is left untouched. That matches the
"Phase 2 only at inference" contract.
Heads dispatch unchanged: the existing GRN kernel reads from
`trunk.heads_w_skip_d` which has off-bucket entries zeroed at the
transition (sparsification per the prior block-diagonal-grad-mask
follow-up commit). Mathematically the full-buffer read is equivalent to
a compact-only per-bucket read; no inference-time kernel change required.
forward_step_golden's convergence test is now architecturally
incompatible with the new dispatch — `forward_only` runs Phase 1
single-CfC math while `forward_step_into` requires populated Phase 2
routing. Annotated `#[ignore]` with a clear divergence note; the
deterministic + reset semantics tests remain valid invariants per
`pearl_training_smoothness_does_not_transfer_to_inference`.
cargo check workspace clean (excluding pre-existing unrelated cupti +
insert_batch errors in vendor/cudarc and ml/tests). ml-alpha + ml-
backtesting lib tests pass (33 + 33).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two follow-up fixes for Smoke 2 readiness, atomic per feedback_no_partial_refactor:
FIX 1 — Task 11 target_jitter_k anchor (spec §3.3 compliance):
- Previously: target_jitter_k = first non-zero raw_per_h[k] (model-state-anchored)
- Now: target_jitter_k = sqrt(realised_label_variance_k) (label-distribution-anchored)
- Per-horizon label variance computed host-side from existing per-horizon
labels in stg_labels (layout [K, B, N_HORIZONS] row-major; typical K=32,
B=1 -> 32 floats per horizon -> trivial host reduction).
- Wiener-α EMA with α = 0.4 floor + first-observation bootstrap per
pearl_wiener_alpha_floor_for_nonstationary +
pearl_first_observation_bootstrap.
FIX 2 — Task 10 Option B closure (block-diagonal heads via grad mask):
- Two new kernels: heads_w_skip_mask_init_kernel (one-shot at transition,
builds [N_HORIZONS × HIDDEN_DIM] = 640-float mask + zeros off-bucket
heads_w_skip in place) and heads_w_skip_grad_mask_apply_kernel
(per-step Phase 2, multiplies grad_heads_w_skip by mask elementwise
before opt_heads_w_skip.step).
- Achieves block-diagonal heads_w_skip behavior WITHOUT touching the
existing GRN kernel: off-bucket positions stay 0 throughout training
because gradient is masked out, so Adam never updates them.
- Existing GRN dispatch consumes heads_w_skip_d (full 640 floats) as
today; off-bucket entries are always 0, so the dispatch is
mathematically equivalent to a compact-only read.
Together these fixes restore the spec-mandated per-horizon differentiation
chain across all 3 mechanisms (CfC.tau bucketing + block-diagonal heads
+ per-bucket LR with label-variance anchor) before Smoke 2 fires the
mean_run_len ratio gate.
GPU oracle tests: 13 total (11 + heads_w_skip_mask_init + heads_w_skip_grad_mask_apply).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.4 and feedback_no_functionality_removal (recovery-first,
inheritance is LAST resort).
New device kernels in bucket_transition_kernels.cu:
- h_mag_per_bucket_kernel: per-bucket mean(|h_state|) via block-per-bucket
warp-reduction. Launch grid=(N_HORIZONS,1,1), block=(32,1,1).
No atomicAdd (block-tree reduce over 32 lanes).
- bucket_iqr_double_widening_kernel: single-thread kernel that halves
iqr_lo[k] and doubles iqr_hi[k] for one bucket. Recovery Attempt 1's
actual implementation.
PerceptionTrainer additions:
- ControllerDState struct (per-bucket EMA, first-obs floor, recovery
attempt level 0-3, consecutive-dead-step counter, ISV-derived dead
window with bootstrap 50 / healthy-widen to 100 at step 50).
- phase2_step_count counter (Phase 2 only).
- h_mag_per_bucket_d device buffer + mapped-pinned shadow.
- Cached h_mag_per_bucket_fn + bucket_iqr_double_widening_fn handles.
Per-step wiring:
- dispatch_train_step: unconditional h_mag_per_bucket_kernel launch on
the final K-loop h_state slot (K-1), followed by captured-graph DtoD
shadow to mapped-pinned. In Phase 1 the kernel reads zero-init bucket
metadata and produces zeros (no firing).
- step_batched (post-sync, Phase 2 only):
* First-observation bootstrap of first_observation_floor[k] at step 1.
* Wiener-α=0.4 EMA update of h_mag_ema[k]
(pearl_wiener_alpha_floor_for_nonstationary).
* Dead-threshold = first_observation_floor[k] * 1/e per spec §3.4.
* Consecutive-dead counter; recovery cascade fires at dead_window_k.
Recovery cascade per spec §3.4:
- Attempt 1 (widen): WIRED. Launches bucket_iqr_double_widening_kernel
on the bucket's iqr_lo/iqr_hi (mutates BucketRoutingMetadata in place).
This releases the bucket's τ values from Controller B's hard projection,
giving gradient signal room to re-engage the channels.
- Attempt 2 (channel swap): LOG-ONLY STUB. Emits tracing::warn! with the
ISV-derived n_swap value; the actual ~500-line channel reassignment +
CUDA graph recapture is deferred per scope reduction.
- Attempt 3 (inheritance): LOG-ONLY STUB. Emits tracing::warn! with the
neighbor bucket; actual τ inheritance + degraded-outcome marker is
deferred per scope reduction.
Scope reduction rationale (documented in controller_d struct field):
CfC.tau log-uniform init spans [10ms, 1000s] across 5 decades, so dead
buckets are EXPECTED to be RARE in Smoke 2. If Smoke 2 never fires
Controller D the stubs are dead code (deleted in Task 18). If Recovery 1
suffices (most likely path), no further work needed. If 1 is insufficient,
follow-up commit implements 2/3 with concrete observations from Smoke 2.
Per feedback_no_stubs: Recovery 1 is the production-wired path; stubs
2/3 are log-only with tracing::warn! so any firing surfaces in logs.
The log boundary is an explicit scope decision, not a deferral.
ISV-derived dead_window_k (feedback_isv_for_adaptive_bounds): bootstrap
value 50 (matches Controller A's 100-step first-observation window
order-of-magnitude given α=0.4 EMA settling ≈ 2-3 steps). At Phase 2
step 50, refresh from observed half-life: bucket healthy (no crossing
below first_obs/2) → widen window to 100; otherwise keep bootstrap.
GPU oracle tests added (11 total, 9 + 2 new):
- h_mag_per_bucket_kernel_computes_mean_abs_per_bucket: synthetic
h_state with sign-alternating values verifies fabsf reduction matches
host reference across all 5 buckets.
- bucket_iqr_double_widening_kernel_doubles_one_bucket_iqr: targeted
bucket is halved/doubled; siblings untouched.
Local: SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets clean;
cargo test -p ml-alpha --test bucket_transition_kernels -- --ignored
passes 11/11; cargo test -p ml-alpha --lib passes 33/33.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.2 + §3.3 and pearl_adam_normalizes_loss_weights.
Controller B (post-Adam τ projection):
- tau_clamp_kernel added to bucket_transition_kernels.cu
- Each channel's τ clamped to its bucket's IQR_widened range after every
Adam step on tau_all_d. Bypasses Adam's m/sqrt(v) normalization
(loss-weight modulation would have been a no-op per the pearl).
Controller C (per-bucket LR multiplier on heads_w_skip):
- heads_lr_multiplier_scale_kernel added.
- Per-step: snapshot heads_w_skip_d (DtoD) before opt_heads_w_skip.step,
run Adam, then rescale the per-horizon delta by
lr_mult_k = 1.0 / (1.0 + smoothness_ratio_k).
- smoothness_ratio_k = max(0, jitter_ema_k - target_jitter_k) /
max(target_jitter_k, target_jitter_k / 16)
with the ε_floor max-with-floor pattern per
pearl_blend_formulas_must_have_permanent_floor.
- target_jitter_k sentinel-bootstrapped from first non-zero raw
smoothness loss per pearl_first_observation_bootstrap.
- jitter EMA shadowed via mapped-pinned smoothness_jitter_ema_host_d
(DtoD captured in graph; host reads after end-of-step sync); host
writes lr_mult_per_horizon_staging (mapped-pinned); next step's
captured-graph rescale kernel reads via dev_ptr.
Scope reduction (Task 11): Controller C scales ONLY heads_w_skip_d
per-horizon (N_HORIZONS × HIDDEN_DIM = 640 contiguous floats, per-horizon
stride HIDDEN_DIM). Other heads weights (w1, w2, w_gate, w_main) stay
unscaled because they're entangled in the GRN gate/main computation and
rescaling them would require additional per-bucket scoping out of Task 11
scope. Revisit if Smoke 2 fails per feedback_no_quickfixes.
Both controllers active Phase 2 only; no-op in Phase 1.
GPU oracle tests: 9 total (7 existing + tau_clamp + lr_multiplier_scale).
All 33 ml-alpha lib tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §2.1 + §5.4 point 1 and Task 10 of the plan.
Adds binding `cfc_step_per_branch_fwd_gpu` in cfc/step.rs that wraps
the fused per-branch kernel with the spec-mandated launch config:
grid=(B, N_HORIZONS=5, 1), block=(MAX_BUCKET_DIM=28, 1, 1) uniform
predicate, cooperative staging of x_local + h_old_local (2 × HIDDEN_DIM
floats). Also adds `cfc_step_per_branch_bwd_gpu` binding for Task 11.
CfcTrunk loads two new cubins (cfc_step_per_branch + heads_block_diagonal)
and caches three function handles. `heads_w_skip_compact_d` (from Task 9)
remains allocated as Phase-2-prepared metadata; full heads consumption
deferred to a follow-up task (Phase 2 heads dispatch needs GRN
integration, which is out of Task 10 scope per
`feedback_no_partial_refactor`).
In `dispatch_train_step`, the CfC dispatch branches on TrainingPhase:
- Phase1Warmup: existing `cfc_step_batched` (unchanged)
- Phase2Routed: `cfc_step_per_branch_fwd` consuming bucket-routed
`tau_all_d` + `bucket_channel_offset_d` + `bucket_dim_k_d`
GRN heads dispatch stays unchanged for both phases per the Task 10
scope adjustment (Option B in the task brief). The `heads_w_skip` storage
remains 640 floats; the compact 128-float `heads_w_skip_compact_d` is
metadata-ready but not yet consumed.
Workspace cargo check clean (ml-alpha lib + trainer compile; only
pre-existing cupti/sp15/gpu_per_integration test errors remain).
ml-alpha lib tests: 33 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.1, §2.3. Controller A computes tau-change EMA with
first-observation bootstrap + Wiener-α floor 0.4 (co-adapting closed
loop); trigger when EMA < max(noise_floor/4, noise_floor/16); ISV-derived
hard cap from observed MAD/median dispersion in first 100 steps.
execute_transition() dispatches the 5 device kernels in order:
tau_sort → bucket_assign → bucket_iqr → tau_reorder → heads_compact.
All-on-device per feedback_no_htod_htoh_only_mapped_pinned (one-time
static-constant upload at transition is allowed off-hot-path).
Returns raw Q1/Q3 in BucketRoutingMetadata.bucket_tau_iqr_{lo,hi}_d;
slack_factor=sqrt(Q3/Q1) widening kernel applied by trainer in Task 9.
3 CPU-only unit tests verify Controller A's first-obs bootstrap,
threshold-trigger, and ISV-cap-trigger semantics.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §5.4 point 2. heads_w_skip storage 5× reduced (640→128 floats).
Each horizon reads only its bucket via offset lookup. Block-tree reduction
(padded to 32 lanes for clean halving stride), no atomicAdd.
GPU oracle test matches naive per-horizon computation (tol 1e-5) on sm_86.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §5.4 point 1. Single fused kernel: grid=(B, N_HORIZONS, 1),
block=(MAX_BUCKET_DIM=28, 1, 1) uniform predicate. Shared x_local +
h_old_local cooperative-staged into shared memory once per block per
pearl_cooperative_staging_eliminates_redundant_reads. No atomicAdd.
No host branches.
Forward GPU oracle: matches naive per-channel CfC math (tol 1e-4) on
sm_86. Backward smoke: kernel loads + writes non-NaN gradients.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §2.5 and feedback_no_partial_refactor. Removes:
- LoaderMode enum (single-variant after Sequential removal → deleted entirely)
- next_sequence_sequential + sequential_file_idx + sequential_anchor_idx + last_call_was_file_boundary + is_file_boundary
- last_seen_file_boundary, train_graph_boundary_state fields
- notify_file_boundary method + boundary-state graph recapture logic
- need_attn_pool_bootstrap match — always true now (random mode always bootstraps)
- --loader-mode CLI flag + notify_file_boundary() call
- loader-mode Argo template param
Additional consumers migrated atomically (beyond the 4 files listed in
the plan): ml-alpha/tests/perception_overfit.rs,
ml-alpha/tests/multi_horizon_loader.rs, ml-backtesting/src/harness.rs,
ml-backtesting/tests/trainer_parity.rs,
ml-backtesting/tests/ring3_replay.rs — all referenced LoaderMode or the
removed config fields.
Workspace builds clean at this commit (pre-existing cudarc-cupti example
and ml-crate test errors are unrelated to MTER removal — they fail at
HEAD too). New per-horizon CfC arch lands in subsequent tasks of the
same plan.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Task 1's attn_pool gate (commit 7dd953aa) was inside the captured-graph
region, so the host branch's decision was frozen at step-2 capture time.
In Sequential mode this would lock the captured graph to whichever
boundary state happened to hold at step 2, breaking subsequent file-
boundary detection — violating pearl_no_host_branches_in_captured_graph.
Fix: track the boundary state at capture time as a new
`train_graph_boundary_state: Option<bool>` field. On each step_batched,
if the current `last_seen_file_boundary` diverges from the captured
state, drop the cached graph and re-capture on the next step.
- Random mode: gate is unconditionally true; state never diverges;
zero recaptures.
- Sequential mode: ~8 recaptures per epoch (one per file boundary).
Each recapture costs the same as the original step-2 capture (~50ms).
Negligible.
Verified clean cargo check -p ml-alpha --all-targets.
Loader can now serve sequences in temporal order (Sequential mode);
when active, the trainer skips attn_pool reset at non-file-boundary
sequence boundaries (next commit will use this for stateful CfC + MTER).
Random mode preserved as default; ZERO behavioral change for existing
runs. Phased commit 1/8 of intervention B per
docs/superpowers/specs/2026-05-21-crt-train-intervention-b-multi-timescale-readout.md
Local trace at base_lambda=0.1 showed the linear ratio
{1, 0.3, 0.1, 0.03, 0.005} = HORIZONS[0]/HORIZONS[h] gave h6000 a
200x stronger target-undershoot signal than h30. The controller
slammed h6000 with smoothness gradient (peak λ[h6000]≈60) while h30
saw nearly none. h6000 val_auc collapsed to 0.514 (near-random)
while h100/h300 improved.
Replace with sqrt(HORIZONS[0]/HORIZONS[h]) = {1, 0.548, 0.316, 0.173,
0.0707}. Caps the differential at ~14x. Verified locally at base=0.1:
- h6000 val_auc recovered to 0.599 (vs 0.514 collapse, vs 0.621 no-smooth)
- jitter ratio h6000/h30 = 0.51 (meaningful differentiation, was 0.84 unforced)
- λ[h6000] equilibrium = 0.7-3 (was 4-60 with linear ratio at same base)
The design target (h6000 jitter = 7% of h30) is less aggressive than
linear (was 0.5%) but produces a tractable training equilibrium that
preserves h6000 predictive capacity. Both 500-step local AND the full
40k-step L40S retrain will tell us where it lands at scale.
The drain task previously used tokio::spawn + tokio::time::sleep,
gated by Handle::try_current().is_ok(). alpha_train's main() is
synchronous (not #[tokio::main]) so Handle::try_current() returned
Err and the drain task was silently skipped — the trace file was
never created.
Refactor to std::thread::spawn + std::thread::sleep. No runtime
required; spawn is unconditional; Drop joins synchronously instead
of aborting; BufWriter flushes on every drain cycle so even short
runs (< 1 sec) capture their tail. Existing gpu_log_ring_invariants
tests migrated from #[tokio::test] back to plain #[test].
Also resolves an exit-1 issue on alpha_train shutdown — likely a
tokio task abort panicking through the synchronous Drop path.
Local smoke (alpha_train --kernel-step-trace ...):
- EXIT=0
- step_trace.jsonl: 1497 lines (~500 steps × 3 smoothness records)
- valid JSONL, schema matches smoothness_controller emission contract
The compile-time feature kernel-step-trace gates inclusion of the ring
code. NEW: --kernel-step-trace <PATH> CLI flag on alpha_train gates
RUNTIME activation:
- Path provided -> ring allocated, drain spawns, JSONL records written
to <PATH> (truncated). One record per kernel emission.
- Path omitted -> ring not allocated, drain not spawned, zero overhead
even with the feature compiled in.
JSONL schema: {"step": u32, "kid": u8, "kname": str, "rt": u8,
"rt_name": str, "payload": {field: f32, ...}}. The payload field names
come from the existing per-(kid, rt) decoders in gpu_log.rs (ported to
serde_json::json! in decode_to_json).
Trainer ring fields are now Option<_>; populated only when the runtime
trace path is Some. Tick kernel, step-counter shadow, smoothness
controller pointer-passing, and Drop all become path-conditional.
The tracing::info!-based drain variant is removed (file-writer is
strictly more useful; tests migrated to assert JSONL file contents).
Argo template:
- New parameter kernel-step-trace-enable (build-time feature opt-in)
- New parameter kernel-step-trace-path (runtime CLI value)
- ensure-binary cache-busts when feature toggles
- training step passes --kernel-step-trace flag conditionally
Per feedback_no_feature_flags: compile-time gate retained because the
ring carries real memory cost (~2 MiB pinned) and per-step write overhead;
the specific name kernel-step-trace narrows scope to this mechanism.
Runtime gate is Option<PathBuf>, not a boolean enable_*.
Per user direction: feature-gating is appropriate for per-step diagnostic
logging (real perf/memory cost) but the name should be specific to the
mechanism, not a generic "diag-log" catch-all. kernel-step-trace
describes what the feature provides — per-step records written to the
ring by kernels.
Pure rename, no behavioral changes. Touched: Cargo.toml feature decl,
lib.rs cfg attribute, perception.rs (~25 cfg attributes including
not(feature) pairs), gpu_log.rs + smoothness_lambda_controller.cu doc
comments, gpu_log_ring_invariants.rs file-level cfg + ignore-attribute
text.
Atomic refactor wiring the GPU log ring's first producer end-to-end:
- cuda/gpu_log_helpers.cuh (new): extract LogHeader/LogRecord/LogRing
struct definitions + the log_record() device __forceinline__ helper
into a shared header. Single source of truth — other kernels include
this rather than redefining the structs.
- cuda/gpu_log_ring.cu: refactor to include gpu_log_helpers.cuh; retain
only the gpu_log_tick kernel.
- cuda/smoothness_lambda_controller.cu: include gpu_log_helpers.cuh, add
trailing (LogRing*, const int*) args, capture pre-EMA jitter +
post-EMA jitter + target + excess_ratio into shared mem, and emit
three records per call (RT_INPUT, RT_STATE, RT_OUTPUT) gated on
non-null ring pointer.
- trainer/perception.rs: feature-gated (cuda-diag-log) ring allocation +
step counter (mapped-pinned host shadow), gpu_log_tick launch FIRST
inside the captured graph, two new pointer args on the smoothness
controller launch (null when feature off), step-counter DtoD shadow
alongside the other telemetry shadows, drain task spawn (skipped when
no tokio runtime — sync tests still construct the trainer), and a
feature-gated Drop impl that aborts the drain task on shutdown.
- tests/smoothness_lambda_controller_invariants.rs: pass null pointers
for the two new kernel args; the kernel's nullptr guard preserves
pre-existing behaviour.
- build.rs: rerun-if-changed on cuda/gpu_log_ids.h and
cuda/gpu_log_helpers.cuh so header edits trigger cubin rebuilds.
Verified: cargo build / check --all-targets clean both with and without
the cuda-diag-log feature; 4/4 smoothness controller invariant tests
pass; 9/9 perception_overfit integration tests pass under the feature.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Task 5 of the GPU log ring plan. Adds `gpu_log` module behind the
`cuda-diag-log` feature gate:
- `LogRecord` (#[repr(C)], 64 B) mirroring the C-side layout in
`gpu_log_ring.cu`; const compile-time size assertion catches drift.
- `LogRing::alloc()` allocates and zero-initialises a 32k-slot
mapped-pinned ring (2 MiB pinned).
- `alloc_step_counter()` allocates the device-side i32 step counter.
- Decoder dispatch table with per-(kid, rt) decoders for
`KID_SMOOTHNESS_CONTROLLER` × {RT_INPUT, RT_STATE, RT_OUTPUT}
emitting structured tracing::info! events.
- `spawn_drain_task()` polls a mapped-pinned step-counter shadow at
500 ms cadence, walks new slots, validates magic + step, dispatches
to decoders. Emits tracing::warn! when the drainer falls behind the
ring window (no silent record drops).
Kernel ID + record-type constants are a manual mirror of
`crates/ml-alpha/cuda/gpu_log_ids.h`; drift is caught by the const
size assertion at compile time and (subsequently) by build.rs.
Task 5 added smoothness_base_lambda to PerceptionTrainerConfig but only
updated the struct definition + Default impl. Eight test sites in
perception_overfit.rs and one site in alpha_train.rs construct the
config via explicit field list (no ..Default::default() spread), so
they broke with E0063 missing-field errors under cargo check
--all-targets.
Per feedback_no_partial_refactor.md: when a contract changes, every
consumer migrates atomically. This commit adds smoothness_base_lambda:
0.0 to all 9 sites so the workspace compiles cleanly. The
alpha_train.rs value of 0.0 is a placeholder — Task 7 will replace it
with cli.smoothness_base_lambda once the CLI flag is added.
Code-quality reviewer flagged two non-blocking polish items in the
output_smoothness kernel's backward pass:
- I1: thread-divergent if-gated loads contradicted the file's stated
branchless-design intent. Fixed via clamp-then-load: at k=0 / k=K-1
read self for the spurious neighbour; lm/rm multipliers zero the
contribution. Every lane executes both loads; no divergence on
boundaries.
- I2: the "Skip work when factor == 0" comment described a non-feature
(no such branch existed). Removed.
Behaviour-equivalent change — same loss and same gradient at every
position; the change only affects warp-execution patterns at the K=0
and K=K-1 strips.
Driven by ffr59 (commit b44a97ff9) findings: all 5 horizons flip
direction every 2.5 events; win rate flat 24% across conviction; mean
PnL anti-correlated with conviction. Hypothesis: per-event alpha output
is high-frequency noise on top of a slower signal. If true, smoothing
input alpha BEFORE the conviction formula should reduce direction flips
and recover signal.
Adds Wiener-α adaptive EMA on raw alpha_probs[h] for each horizon,
applied BEFORE the multi-horizon conviction formula. Floor at 0.1
(stronger than the 0.4 floor on the output-side conviction EMA — this
tests whether INPUT smoothing has different impact than OUTPUT smoothing).
Three new device slots:
- alpha_ema_per_b_per_h (per-horizon EMA state)
- alpha_diff_var_per_b_per_h (variance of changes)
- alpha_sample_var_per_b_per_h (variance of value)
The CRT.diag Group A direction-flip counter still reads RAW alpha_probs
so we have a head-to-head comparison: raw flip rate vs smoothed flip rate.
Group E adds the smoothed-direction counter + mean run length.
End-of-run log adds one line per horizon:
crt_diag h<X> smoothed: flips=Y mean_run_len=Z events (vs raw F / M)
If smoothed mean_run_len >> raw mean_run_len: hypothesis is RIGHT, the
input had signal under the noise. Next step would be to make this an
operational EMA in the controller.
If smoothed and raw are similar: hypothesis is WRONG, per-event output
is genuinely noisy. Next step would be to investigate model training
(horizon collapse) OR the AUC=0.66 measurement definition.
Either way, definitive result from one smoke run.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Pure-instrumentation commit. Zero behavior change. The Gate CRT.1 smoke
p9cnk (commit 1e656948b) produced 222k trades and 29688% drawdown,
worse than the failed CRT.A. The structural fixes (multi-horizon
conviction, no-trade band, composite exit) all work as designed — but
the trade cadence is fundamentally too fast for the signal/cost ratio.
Before another round of threshold tuning, measure the structural truth.
Adds device-side counters dumped at smoke close (single end-of-run
memcpy_dtoh batch in read_diagnostics(), NOT per-event):
Group A — per-horizon signal persistence:
- flip_count[h]: total direction-sign flips
- sum_run_length[h]: cumulative events between flips (u64)
- run_length_hist[h]: 5-bucket log-scale histogram (1-9, 10-99,
100-999, 1k-9.9k, 10k+)
Output: mean_run_len per horizon = signal half-life proxy
Group B — conviction smoothed-EMA distribution:
- conv_hist[10]: 10-bucket histogram over [0, 1)
Output: shape of conviction distribution at decision time
Group C — hold-time distribution:
- hold_hist[6]: <1s, 1-10s, 10-60s, 1-10m, 10-60m, >1h
Output: empirical distribution of trade hold times
Group D — outcome by entry-conviction:
- outcome_n[10]: trades per conviction bucket
- outcome_sum_pnl[10]: cumulative realized PnL per bucket (price-units)
- outcome_n_wins[10]: wins per bucket
Output: validates "high conviction = better trades" hypothesis OR
surfaces signal miscalibration
Per-backtest memory: ~330 bytes. Negligible vs existing state.
All counters single-writer-per-block (threadIdx.x == 0 convention; no
atomicAdd per feedback_no_atomicadd). All updates kernel-internal (no
host roundtrip per event). The end-of-run dump uses memcpy_dtoh in
read_diagnostics() — called once at harness termination from the
post-stop_ctrl_counters block, never per-event. Per
feedback_no_htod_htoh_only_mapped_pinned, the hot path is untouched.
N_HORIZONS = 5 (heads.rs canonical {30, 100, 300, 1000, 6000}); the
plan text says 4 but the workspace constant is 5 — used the code value.
Output at end of cluster smoke as a `crt_diag` log block:
crt_diag h30: flips=X mean_run_len=Y events
crt_diag h30 run_length_hist: 1-9:N 10-99:N 100-999:N ...
...
crt_diag conv_ema_hist: [0.0-0.1]:N [0.1-0.2]:N ...
crt_diag hold_time_hist: <1s:N 1-10s:N ...
crt_diag outcome_by_entry_conv[0.0-0.1]: n=N win_rate=X.X% mean_pnl_pu=Y.YYY
...
The output drives the next round of CRT.1 threshold design (delta_floor,
composite exit factors, min-hold cooldown) from data instead of guesses.
Per pearl_adaptive_not_tuned and pearl_controller_anchors_isv_driven:
thresholds will become ISV-derived in CRT.2; CRT.1 must first establish
defensible defaults from the empirical signal characteristics.
Verified: stop_controller (22/22 pass), decision_floor_coldstart (3/3
pass), threshold_and_cost (3/3 pass), parallel_sim_correctness (1/1
pass), lob_sim_integrated_fuzz (3/3 pass). Two lob_sim_fixtures
failures (fix_decision_alpha_buy_close, fix_decision_program_h4_only)
were pre-existing at HEAD 1e656948b — not caused by this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.3 + plan C1.4. Safety backup for the primary exit path
(target → 0 from collapsed multi-horizon conviction, C1.2). Catches
the edge case where conviction stays bounded above zero but the trade
is deep in unrealized loss, or short/long horizons disagree for many
events while threshold-gate logic still permits sizing.
Three terms, max-of-three composite:
degradation_term = (1 − conv_ema_now / conviction_at_entry) / 2.0
disagreement_term = disagree_consecutive_events / 32.0
pnl_decay_term = max(−pnl_adjusted_conv_ema, 0) / 1.0
exit_signal = max(deg, dis, pnl_decay)
Fires force-flat (market_targets = (3, 0)) when exit_signal >= 1.0.
State-update flow:
- pnl_track open branch: write conviction_at_entry (offset 20),
conviction_ema (offset 24, initialised to entry value),
pnl_adjusted_conviction_ema (offset 44, = 0),
peak_unrealized_pnl (offset 48, = 0). Takes new conviction_ema_per_b
read-only arg to snapshot at open.
- decision kernel intra-event (when pos open): update offsets
24/44/48/56 (disagreement counter) via update_open_trade_trajectory.
- composite_exit_check device helper: reads offsets 20/24/44/56,
writes (3, 0) on fire. Sibling of stop_check_isv with same return
semantics; called after trajectory update so all four reads see
the freshly-written values.
Fields not read in CRT.1 (offsets 28-43 per-horizon EMA, 52 degradation
counter, 60 horizon_idx_dominant) are NOT written by the open branch —
alloc_zeros covers init. Per feedback_no_hiding: don't populate unused
fields. Also fixes a latent C1.1 bug: pnl_track close branch read
horizon_idx from offset 20 (now conviction_at_entry, f32); moved the
read to offset 60 per the documented layout.
Threshold-bootstrap defaults (will be ISV-derived in CRT.2 alongside
the rest of the controller anchors per pearl_controller_anchors_isv_driven):
degradation_factor = 2.0 (50% conviction decay vs entry → ≤0.5 term)
disagreement_factor = 32.0 (32 events of short/long disagreement)
pnl_decay_factor = 1.0 (pnl_adj_conv structurally ∈ [-1, +1])
Under these defaults disagreement_term is the only term that can fire
on its own — degradation_term caps at 0.5 (conv_ema/conviction_at_entry
≥ 0) and pnl_decay_term caps at 1.0 in the limit. The three terms still
add evidence in superposition; CRT.2 will retune factors against ISV
percentiles. Unit-test coverage validates the disagreement path
(composite_exit_signal_fires_on_disagreement).
Per pearl_no_host_branches_in_captured_graph — all composite logic
lives inside the decision kernel; no host roundtrip in the per-event
hot path. Per feedback_no_htod_htoh_only_mapped_pinned — no new
memcpy or stream.synchronize() in step_pnl_track or
dispatch_latent_market_orders.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The C1.3 commit 7c850a6c0 missed three test files that construct
UniformSimParams literals. cargo check reported E0063 missing-field
errors. Added delta_floor: 0.0 (band disabled, preserves pre-C1.3
behavior in tests) to:
- parallel_sim_correctness.rs
- decision_floor_coldstart.rs
- lob_sim_integrated_fuzz.rs
Per feedback_no_partial_refactor: the atomic refactor must include
every consumer in one logical change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.0 + plan C1.3. Without this gate, every fractional change
in target_lots seeds an order via target-delta semantics; combined with
the continuous controller invocation (every event after A1) this
generates hyperactivity even with the multi-horizon §4.4 conviction
formula (C1.2). v2 Gate-1 failures confirmed: scalar EMA + amplification
clamp + per-event controller = 155k trades on a 2M-event smoke.
The band absorbs fractional target adjustments so most events produce
no order. When the signal genuinely shifts and target moves enough to
cross the band, the kernel seeds. Continuous evaluation, sparse action.
Greenfields atomic refactor — single config knob threaded through every
layer in one commit:
SweepBase.delta_floor: f32 (YAML, default 1.0 = 1 lot)
SimVariant.delta_floor: Option<f32> (per-variant override)
ResolvedSimVariant.delta_floor: f32
UniformSimParams.delta_floor: f32
BatchedSimConfig.delta_floor: Vec<f32>
LobSimCuda.delta_floor_d: CudaSlice<f32>
seed_inflight_limits_batched kernel param + skip-if-below-floor logic
New accessor: LobSimCuda::read_inflight_count(b) — counts active != 0
limit slots for backtest b. Not cfg(test); follows existing read_limit_slot
pattern.
Test: no_trade_band_blocks_micro_delta verifies that a second decision
with the same strong-bullish alpha (same target, effective = in-flight
lots) produces delta=0 and does not re-seed. max_lots=1 keeps the
arithmetic unambiguous.
Existing tests: all UniformSimParams struct literals updated with
delta_floor=0.0 (band disabled) so existing behaviour is preserved.
harness.rs from_uniform path uses delta_floor=1.0 (production default).
Hot-path discipline: the delta_floor_d upload happens once per run in
the existing config-upload block (alongside threshold, cost, max_hold_ns);
the kernel reads the device slot per-event without any host roundtrip. No
memcpy_htod / dtoh / dtov / synchronize introduced on the per-event path.
Per pearl_controller_anchors_isv_driven, this floor is currently a
config constant; CRT.2 (Phase 2) makes it ISV-derived from rolling
spread cost / signal volatility.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.4 + plan C1.2. Replaces v2 A2 (1d889d2de) + A2.1
(fe2498769) scalar approach that failed Gate 1 catastrophically on
smokes vjmwc and lkrdf (155k trades, 9100% drawdown, Sharpe -15.7).
The structural fix: direction comes from per-horizon SIGNED sum, so
disagreeing horizons cancel. The scalar EMA on max(|p-0.5|) could only
smooth magnitude — it could not smooth direction jitter. Only this
multi-horizon weighted formula can.
Formula (spec §4.4):
weight_h = max(isv[h].net_edge, ε) / (isv[h].var + cost²)
weighted_h = magnitude_h × weight_h × direction_h
conviction_signed = sum(weighted_h) / sum(|weight_h|)
conviction_ema = Wiener-α EMA(|conviction_signed|) (reuses A2 slots)
target_lots = round(sign(conviction_signed) × conviction_ema × max_lots)
Pearl conformance:
- pearl_controller_anchors_isv_driven: weights from ISV state
- pearl_one_unbounded_signal_per_reward: net_edge/(var+cost²) is the
single unbounded factor; magnitude × direction is bounded in [-1,1]
- pearl_zscore_normalization_for_magnitude_asymmetric_signals: divide
by total_abs_weight for scale invariance
- pearl_trade_level_vol_for_stop_distance: cost² is variance bootstrap
- pearl_blend_formulas_must_have_permanent_floor: ε = cost · 0.01 on
net_edge — no cold-start regime switch
- pearl_audit_unboundedness_for_implicit_asymmetry: net_edge is
one-sided positive (losing-edge horizons drop out, not flip sign)
- pearl_wiener_alpha_floor_for_nonstationary: α floor at 0.4 (unchanged
EMA mechanism, repointed at the multi-horizon scalar)
DELETED:
- The `final_size *= (conv_ema/raw_max_conv)` aggregate rescale step
(the v2 bug that amplified weak signals)
- The per-horizon `signed_sizes[h]` Kelly + Sharpe-weight aggregation
loop in decision_policy_default (replaced by direct §4.4 formula)
- Three obsolete tests asserting on the old scalar path:
* conviction_ema_smooths_micro_oscillations
* conviction_ema_does_not_lag_reversals
* conviction_ema_rescale_never_amplifies_weak_signal
* cfg_high_kelly_floor helper (only used by deleted tests)
- cold_start_with_zero_floor_reproduces_old_bug (tested OLD kelly_floor
semantics that v3 doesn't expose)
- cold_start_persistent_bullish_now_closes (was passing under v2 via
target oscillation from integer rounding; v3 produces stable target
on stable signal — closes need price movement)
- alpha_noop_side_2_preserved (anchored on the v2 kelly_frac_floor
integer-truncation path that doesn't exist under v3)
Applied to BOTH decision_policy_default and decision_policy_program
(bytecode VM at OP_WRITE_ORDER). Per feedback_no_partial_refactor —
both consumers migrate atomically. The VM's per-horizon emit + aggregate
opcodes still execute but their stack `final_size` is unused at
OP_WRITE_ORDER (the §4.4 conviction-driven target replaces it). The
stack's attribution mask is still consumed for entry crediting.
The three conviction-EMA device slots (conviction_ema_d, _diff_var,
_sample_var) STAY in LobSimCuda. The Wiener-α EMA mechanism is reused
on the multi-horizon conviction scalar; the device-helper signature
changed from taking alpha_probs[] to taking the precomputed scalar
|conviction_signed|.
Kernel ABI change: decision_policy_default no longer takes
target_annual_vol_units / annualisation_factor / kelly_frac_floor /
sharpe_weight_floor (the §4.4 formula doesn't use them). Removed from
both the CUDA signature and the launch builder in sim/mod.rs. The
bytecode VM (decision_policy_program) still consumes those params for
its OP_EMIT_PER_HORIZON_SIZE / OP_AGG_WEIGHTED_SHARPE opcodes — their
device slots remain allocated.
Test: multi_horizon_conviction_cancels_on_disagreement verifies the
structural fix — bullish [0.7,0.3,0.7,0.3,0.5] horizons cancel to
no-trade output (side=2/3, size=0, conv_ema≈0).
Existing tests rebased: cfgs that used cost_per_lot_per_side=0.0 now
use 1.0 (the §4.4 formula's eps_edge floor requires cost > 0). Each
affected test file documents the rebase inline. Tests that were
anchored on observed v2-kernel values (per
pearl_tests_must_prove_not_lock_observations) are deleted; tests with
genuine invariants (boundedness, direction-sanity, stop-controller
behavior) are preserved.
Hot-path discipline: no memcpy_htod/dtoh/dtov/synchronize introduced.
Only kernel-internal computation; one launch arg removed from the
default decision kernel.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §7 (v3 architecture reset, promoted from v2 Phase B into CRT.1).
Single-cache-line layout. Every reader and writer of open_trade_state
migrates in this commit per feedback_no_partial_refactor.
New 64-byte layout (existing 24-byte fields preserved at original offsets):
Offset Size Field
0 8 entry_ts_ns (u64)
8 4 entry_px_x100 (i32)
12 4 entry_size (i32)
16 4 realized_at_open (f32)
20 4 conviction_at_entry (f32) ← NEW
24 4 conviction_ema (f32) ← NEW
28 16 conviction_per_horizon_ema [4 × f32] ← NEW
44 4 pnl_adjusted_conviction_ema (f32) ← NEW
48 4 peak_unrealized_pnl (f32) ← NEW
52 4 degradation_consecutive_events (u32) ← NEW
56 4 disagreement_consecutive_events (u32) ← NEW
60 1 horizon_idx_dominant_at_entry (u8) ← NEW
61 3 pad
The 24-byte prefix is unchanged so downstream readers (stop_check_isv in
decision_policy.cu, max_hold event-rate check in resting_orders.cu) need
only update the stride constant. The new fields at offsets 20..63 are
populated by subsequent CRT.1 tasks (C1.2 multi-horizon conviction, C1.4
composite exit signal).
Open branch in pnl_track.cu zeros the new fields implicitly because
alloc_zeros on the device slot zero-initialises everything; subsequent
writes only touch the 0-23 byte range (existing fields). Close branch
reset-loop iterates OPEN_TRADE_STATE_BYTES which now zeros all 64.
Files changed:
crates/ml-backtesting/src/lob/mod.rs — pub const OPEN_TRADE_STATE_BYTES: usize = 64
crates/ml-backtesting/cuda/pnl_track.cu — #define OPEN_TRADE_STATE_BYTES 64
crates/ml-backtesting/cuda/resting_orders.cu — comment + stride literal 24→64
crates/ml-backtesting/tests/stop_controller.rs — open_trade_state_64_byte_layout test
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Gate 1 catastrophic failure (smoke vjmwc, commit 3d8f12deb):
- 62× hyperactivity (157,470 trades vs 2,511 baseline)
- 9,347% max-drawdown ($327M loss on $3.5M base)
- Sharpe -15.63 (2.4× worse than baseline)
Root cause: A2's formula
final_size *= (conv_ema / raw_max_conv)
amplified weak signals because EMA tracks the MEAN of raw_max_conv, NOT a
running max as the author's comment claimed. When raw_max_conv < conv_ema
(normal during quiet periods between strong signals), the multiplier was
> 1, blowing up small-signal positions:
raw=0.05, ema=0.3 → 6× amplification
raw=0.01, ema=0.3 → 30× amplification
Fix: clamp scale ∈ [0, 1] in BOTH decision_policy_default and
decision_policy_program (OP_WRITE_ORDER). Only damp when raw is stronger
than EMA (scale < 1); never amplify when raw is weaker (scale capped at 1).
Math:
scale = conv_ema / raw_max_conv // can be 0..∞
scale_clamped = min(scale, 1.0) // can only be 0..1
final_size *= scale_clamped
Test: conviction_ema_rescale_never_amplifies_weak_signal validates
target_lots stays ≤ 1 when alpha=0.51 (raw_conv=0.02) after EMA warmup
at alpha=0.7 (ema~=0.4). Without the clamp the broken multiplier is 20×.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The forward_step_bit_identical_after_seed_from_forward_only test in
commit 1d889d2de asserted 1e-5 prediction-level convergence between
forward_only(W[1..K+1]) and seed+forward_step(snap[K]). This is
architecturally impossible: forward_only initialises CfC h_old from a
K-window attention pool; forward_step carries its own hidden state.
Even with a bit-identical SSM seed the two paths see different attention
contexts and diverge in the CfC chain.
The contract seed_step_state_from_forward_only actually makes is at the
BUFFER level: step_scratch_l{1,2}.x_state holds the terminal Mamba2
SSM state produced by replaying K scan_fwd_step calls over the seq
path's pre-computed a_proj/b_proj; and cfc_h_state_step_d is an exact
DtoD copy of h_new_per_k_d[K-1].
Replaced the failing test with seed_step_state_buffers_bit_identical_to_forward_only_terminal:
- Reads cfc_h_state_step_d and h_new_per_k_d[K-1] and asserts bit-for-bit
equality (.to_bits() == .to_bits()) — the DtoD copy makes this exact.
- Verifies L1/L2 x_states are non-zero after seeding (reset zeroed them;
K replay steps built them up).
- Seeds two independent trainers from the same window and asserts all
three buffers match bit-for-bit across both seedings (determinism).
Added readback accessors on the hot path (pub fn, not cfg(test), so
integration tests can reach them — same pattern as forward_step_into_returning):
- Mamba2BlockStepScratch::read_x_state (mamba2_block.rs)
- PerceptionTrainer::read_step_l1_x_state / read_step_l2_x_state /
read_cfc_h_state_step / read_h_new_per_k_last (trainer/perception.rs)
The architectural divergence at prediction level is documented in the
replacement test's docstring so future readers don't reopen the same question.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
