The liquid_tau_rk4_step kernel modelled f(x) = (1/tau)*(1-x), which has
fixed point x=1.0 for any tau. liquid_mod_buf was initialised to [1.0;4]
and the dynamics could never move it away from that fixed point since every
kernel call reduces to f(1.0)=0 (no perturbation path exists). The
velocity_mod multiplier in c51_grad_kernel was therefore always 1.0 — a
mathematical identity with zero measurable effect on spread_scale.
Additionally, no ISV slot existed for liquid_mod (violates §4.C.6
GPU-drives spec), and the associated LiquidTrainableAdapter supervised
path was never wired into the DQN backward pass.
Decision C (delete): remove liquid_tau_rk4_step kernel (experience_kernels.cu),
liquid_mod param + velocity_mod line (c51_grad_kernel.cu), liquid_mod_buf
allocation, liquid_tau_rk4_kernel field, update_liquid_tau() method and its
call site in update_eval_v_range() (gpu_dqn_trainer.rs). per_branch_q_gap_ema_buf
is retained — it is used independently for trajectory backtracking state
snapshot/restore. Supervised LiquidTrainableAdapter unchanged.
cargo check -p ml: 8 warnings (baseline), 0 errors.
Audit docs updated: dqn-wire-up-audit.md + ml-supervised-to-dqn-concept-audit.md.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
SoftReset registry category (isv_grad_balance_targets, isv_grad_scale_limit)
implemented via bootstrap-write + existing kernel's EMA. CPU writes bootstrap
values (1.0 for targets, 2.0 for limit) at fold boundary; grad_balance_isv_update
kernel's adaptive-rate EMA (alpha in [0.01, 0.30]) converges from bootstrap
toward observed values over subsequent epochs (implicit decay_bars via EMA time
constant).
Option A chosen (minimal): the existing grad_balance_isv_update kernel already
uses adaptive-rate EMA — not a hard-write — so bootstrap-at-fold-boundary is
sufficient. The EMA time constant ~1/alpha provides the decay, satisfying the
decay_bars=500 intent without a new ISV slot (Option B rejected as over-
engineered: new ISV slot + kernel change for no material behavioral improvement).
Changes:
- training_loop.rs::reset_named_state: add dispatch arms for both SoftReset
entries, writing bootstrap constants (CPU-born input, not adaptive output;
complies with spec §4.C.6 GPU-drives-CPU-reads)
- trainer/mod.rs: fold-boundary reset loop now calls both fold_reset_entries()
and soft_reset_entries(); stale "handled separately (Plan 2)" comment removed
- smoke_tests/soft_reset.rs: 4 CPU-only tests verify registry classification,
entry count, mutual exclusivity from fold_reset, and bootstrap constant invariants
- docs/dqn-wire-up-audit.md: D.5 SoftReset dispatch row added
No CPU-side adaptive computation. No new ISV slot. No behavioral change to
training beyond writing known-good bootstrap values at fold boundaries.
Plan 2 Task 4. Spec §4.D.5.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan 1 A.5 audit found mamba2_scan_projected_bwd kernel + host call at
gpu_dqn_trainer.rs::mamba2_backward are ALREADY fully wired in the
adam_grad CUDA-graph child. Plan 2 Task 2 narrows from "implement" to
"validate".
Two smoke tests confirm correctness:
- mamba2_backward_gradients_propagate: grad Frobenius norm 0.25 after 3
epochs (>> 1e-8 threshold), ruling out silent no-op like compute_iqr
had pre-Task-A.6.
- mamba2_backward_grad_check: kernel-level reference check (B=2 K=2
SH2=4 STATE_D=4); max rel_err d_gate=2.6e-7, d_x_out=2.6e-7,
d_context=6.7e-8 — all well within 15% threshold (near machine
epsilon, confirming bit-identical host/GPU results).
No production code change — test-only accessors exposed via #[cfg(test)]
impl blocks on GpuDqnTrainer and DQNTrainer. Audit doc updated.
Plan 2 Task 2. Spec §4.D.1.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Revises Plan 2 to match the post-Plan-1 reality:
1. AdaptiveController → AdaptiveMonitor throughout (read-only observers
per spec §4.C.6 2026-04-24 revision; GPU kernels compute, CPU reads).
Applies to Task 3 (per-branch gamma) and Task 5 (liquid_mod audit).
2. Task 2 (D.1 Mamba2 backward) scope narrowed from "implement" to
"validate existing". Plan 1 A.5 audit confirmed
mamba2_scan_projected_bwd kernel + mamba2_backward host call are
already fully wired. Task 2 now adds a finite-difference grad-check
smoke + non-zero grad-propagation smoke; escalates if either fails.
3. Task 3 (D.2 per-branch gamma) rewritten for GPU-drives compliance:
- Replace scalar GAMMA_EFF_INDEX=43 with 4 per-branch slots at
43-46 (DIR/MAG/ORD/URG).
- New per_branch_gamma_update_kernel.cu reads v-range + health,
writes 4 slots. Deletes the Plan 1 gamma_update_kernel.cu.
- 4 read-only monitors (or one consolidated PerBranchGammaMonitor).
- ISV slots 44-48 shift downstream; fingerprint re-tails at 50-51;
ISV_TOTAL_DIM grows 49 → 52.
- c51_loss_kernel and iql_value_kernel read per-branch γ from ISV.
- No CPU-side γ computation anywhere.
4. Header architecture block updated: new ISV slot count target,
GPU-drives principle explicit, current Plan-1 layout table inline
for reference.
5. Pre-plan verification updated: expects AdaptiveMonitor trait (not
AdaptiveController); checks for 6 Plan-1 monitor files.
6. Removed "schema version bump 1 → 2" language — layout fingerprint
auto-recomputes from seed bytes; no integer version space exists.
Task 4 (D.5 soft fold transitions), Task 5 content (D.7 liquid audit),
Task 6 (D.3+D.6+D.8 coordinated state-layout migration), Task 7
(validation run) structure preserved. Internal slot numbers in those
tasks will reconcile to the new layout when they land (no pre-emptive
edits since those indices depend on whether Task 1 quantile slots or
Task 3 per-branch gamma slots land first — re-compute at exec time).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Local smoke-test run after Plan 1 C.6 completion surfaced three issues:
1. StateResetRegistry missing dispatch arms for the 8 ISV slots
pre-allocated in ac9bcab94 (isv_epoch_idx, isv_epsilon_eff, isv_tau_eff,
isv_gamma_eff, isv_kelly_cap_eff, + 3 SchemaContract which already no-op).
Fold boundary would error "unknown name 'isv_epoch_idx'". Added dispatch
arms in training_loop.rs::reset_named_state — reset each to 0.0; GPU
kernels repopulate on next epoch.
2. controller_activity smoke test's single 50% threshold was designed for
reactive CPU-compute controllers. Under GPU-drives-CPU-reads, tau is a
Polyak-EMA cosine schedule that fires every epoch by design (95%),
gamma is health-coupled monotonic (may fire every epoch as health
drifts). Split threshold per-controller: reactive (anti_lr, grad_clip,
cql_alpha, cost_anneal) = 0.50; schedule-based (tau, gamma) = 1.00.
3. examples/train_baseline_rl was broken since the f64→f32 ABI refactor
(d64adc14f) — hp_f64 returning f64 assigned to f32 fields. Added
hp_f32 helper that narrows JSON-born f64→f32 at ingest boundary. Use
hp_f32 for f32 fields, hp_f64 for f64 fields (learning_rate,
entropy_coefficient, weight_decay). No more "as f32" casts at call
sites. Also fixed replay_buffer_vram_fraction + bars_per_day f64→f32.
Local smoke tests now pass:
- controller_activity: ok (1 passed, 29.5s)
- multi_fold_convergence: ok (1 passed, 3 folds x 20 epochs, 534.9s)
- 24 new monitor + registry unit tests: all passing
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The test observed `mon.observe(0.0)` first but the monitor's initial
last_value is also 0.0, so the first observation didn't fire. Test
expected 2/3 but got 1/3.
Changed observe sequence to (1.0, 1.0, 1.5) mirroring grad_balancer's
correct pattern: first differs from initial 0.0 → fires, second matches
→ no fire, third differs → fires = 2/3.
No production code change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
atoms_update_kernel.cu: 4-block kernel (grid=(4,1,1), block=(256,1,1)) replaces
the CPU loop that launched adaptive_atom_positions 4× per branch. Each block reads
ISV v-range slots [23..31) and spacing_raw params[52..56); computes softmax +
cumsum; writes atom_positions_buf[branch * num_atoms]. Same formula as the
existing adaptive_atom_positions kernel in experience_kernels.cu.
Deleted recompute_atom_positions and warm_start_atom_positions from GpuDqnTrainer
and their fused_training.rs delegates. step_atom_positions now calls
launch_atoms_update. training_loop.rs: recompute_atom_positions → launch_atoms_update;
warm_start_atom_positions + compute_reward_quantiles block removed.
AtomsMonitor reads ISV[V_CENTER_DIR_INDEX=23] as representative scalar;
diagnose() exposes all 8 v-range slots [23..31). state_reset_registry.rs
"follow-up task" descriptions updated for all 4 adaptive slots.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
kelly_cap_update_kernel.cu: single-thread cold-path kernel aggregates Kelly
win/loss stats across all n_envs from portfolio_states[n_envs, PS_STRIDE].
Computes mean half-Kelly × health-coupled safety multiplier (0.5+0.5*health)
and writes ISV[KELLY_CAP_EFF_INDEX=44]. Matching Bayesian priors from
trade_physics.cuh (prior_wins=2, prior_losses=2) prevent cold-start collapse.
GpuExperienceCollector::portfolio_states_dev_ptr() added — exposes device
pointer for the portfolio_states buffer without a GPU→CPU roundtrip.
GpuDqnTrainer::launch_kelly_cap_update takes ps_dev_ptr + n_envs.
FusedTrainingCtx delegates to trainer. training_loop.rs launches at epoch
boundary after gamma_update (requires both fused_ctx and gpu_experience_collector).
KellyCapMonitor reads ISV[44] for HEALTH_DIAG via AdaptiveMonitor trait.
state_reset_registry.rs description updated; audit docs updated (Invariant 7).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
tau_update kernel computes Polyak-EMA tau from ISV[EPOCH_IDX=39,
TOTAL_EPOCHS=40, LEARNING_HEALTH=12] with cosine schedule + health-coupled
floor (rises to 0.01 at full collapse). Single-thread cold-path kernel
writes ISV[TAU_EFF_INDEX=42].
TauMonitor is a read-only observer exposing tau_eff, epoch_idx, total_ep,
health, progress, fire_rate in DiagSnapshot.
Wires EPOCH_IDX_INDEX CPU-born input write at epoch start via existing
write_isv_signal_at accessor (pinned-memory zero-copy). Also launches
tau_update kernel at epoch boundary before any tau consumer.
Adds read_isv_signal_at() to GpuDqnTrainer (symmetric counterpart to
write_isv_signal_at).
Deleted: compute_cosine_annealed_tau free function (fused_training.rs),
apply_health_coupled_tau_floor method (GpuDqnTrainer), last_tau_eff
cached field + last_tau_eff() delegate. 4 consumer sites in
fused_training.rs now read ISV[TAU_EFF_INDEX] directly.
Tests: 3 monitor unit tests pass. cargo check -p ml at 8-warning baseline.
Plan 1 Task 13. Spec §4.C.6 (2026-04-24 revision).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Lands the static-configuration branch of Plan 1 C.6: cql_alpha,
conviction_floor (no-op if IQL_BRANCH_SCALE_FLOOR already serves),
plan_threshold written to dedicated ISV slots at constructor. Consumer
kernels (CQL, backtest_plan, experience) read from ISV instead of
config fields / hardcoded literals.
Also pre-allocates ISV slots for the 6 upcoming GPU-kernel tasks
(atoms, gamma, kelly_cap, tau, epsilon outputs + CPU-born epoch inputs).
Those slots start at 0; GPU kernels in follow-up commits fill them.
New ISV slots:
- EPOCH_IDX_INDEX=39, TOTAL_EPOCHS_INDEX=40 (CPU-born inputs)
- EPSILON_EFF_INDEX=41, TAU_EFF_INDEX=42, GAMMA_EFF_INDEX=43,
KELLY_CAP_EFF_INDEX=44 (GPU-written in follow-up tasks)
- CQL_ALPHA_INDEX=45, PLAN_THRESHOLD_INDEX=46 (static config; this commit)
- Task 15 confirmed no-op: IQL_BRANCH_SCALE_FLOOR_INDEX=36 already
serves conviction-floor role (constructor + ISV read in iql kernel).
Layout fingerprint auto-updated via seed-byte edits; fingerprint
re-tail at [47..49). ISV_TOTAL_DIM 39 -> 49.
GpuDqnTrainConfig gains total_epochs field; fused_training.rs passes
hyperparams.epochs at construction; default 0 for smoke tests.
write_isv_signal_at bound extended from ISV_DIM(23) to ISV_TOTAL_DIM(49)
so tail slots are writable by CPU.
cql_alpha consumer: compute_cql_logit_gradients reads base from
ISV[CQL_ALPHA_INDEX] instead of config.cql_alpha; adaptive formula
(base x health x (1-regime_stability)) unchanged.
plan_threshold consumers: experience_kernels.cu (experience_state_gather,
experience_action_select, experience_env_step) and backtest_plan_kernel.cu
(backtest_plan_state_isv) read plan_thr from ISV[ISV_PLAN_THRESHOLD_IDX=46]
via isv_signals pointer already present in both kernels; null-guard defaults
to 0.5f for smoke-scale runs without ISV warmup.
ISV_PLAN_THRESHOLD_IDX=46 defined in state_layout.cuh (included by both
plan kernel files); value must match PLAN_THRESHOLD_INDEX in gpu_dqn_trainer.rs.
StateResetRegistry entries added for all new slots: SchemaContract for
TOTAL_EPOCHS/CQL_ALPHA/PLAN_THRESHOLD; FoldReset for EPOCH_IDX and
GPU-written per-fold slots.
No behavioural change: all threshold/base values remain at their prior
defaults; consumers now read adaptive values from ISV instead of baked-in
config or hardcoded literals.
Plan 1 Tasks 12, 15, 16 + pre-allocation for 9, 10, 11, 13, 14.
Spec §4.C.6 (2026-04-24 GPU-drives-CPU-reads revision).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Reverts commits d76849f31 (Batch A: atoms, gamma, kelly_cap, cql_alpha)
and 4189da563 (Batch B: tau, epsilon, conviction_floor, plan_threshold).
The reverted commits implemented 8 controllers under the old
AdaptiveController trait which had CPU-side update() that computed
adaptive values and write_output() that pushed them to ISV. This
violates the architectural principle codified in the §4.C.6 revision
(commit cf36091ee): GPU kernels compute all adaptive decisions; CPU is
pure observation.
The 8 migrations will be re-implemented under the new AdaptiveMonitor
pattern:
- 6 reactive mechanisms (atoms, gamma, kelly_cap, tau, epsilon,
grad_balancer) get new GPU kernels + read-only CPU monitors.
- 3 static mechanisms (cql_alpha, conviction_floor, plan_threshold)
get ISV constructor-writes (no kernel, no monitor).
After this revert:
- AdaptiveController trait is back on main (from Task 8's 419c24b4f).
It will be replaced with AdaptiveMonitor in the next commit per the
revised Plan 1 Task 8.
- StateResetRegistry (from Task 2's b688827d6) stays intact.
- Tasks 1-7 completed work unchanged.
Tests: cargo check -p ml at 8-warning baseline after revert.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
User raised that the CPU-controller pattern (where CPU computes adaptive
values from ISV and writes back to ISV) violates the architectural
principle: GPU kernels should compute adaptive decisions; CPU-side code
is pure observation.
Replaces the AdaptiveController trait with read-only AdaptiveMonitor:
- No update() — CPU doesn't compute adaptive values.
- No write_output() — CPU doesn't write adaptive values to ISV.
- read() returns the GPU-computed value from ISV.
- diagnose() emits HEALTH_DIAG snapshot.
- observe() + fire_rate() track how often the GPU-computed value changes.
Reclassifies the 9 adaptive mechanisms:
- 6 reactive get GPU kernel + CPU monitor: atoms, gamma, kelly_cap, tau
(Polyak EMA), epsilon, grad_balancer (last is already GPU-driven).
- 3 static get ISV constructor-write, no monitor: cql_alpha,
conviction_floor, plan_threshold.
New ISV slots for GPU-written adaptive outputs (EPSILON_EFF, TAU_EFF,
GAMMA_EFF, KELLY_CAP_EFF) and CPU-born inputs (EPOCH_IDX, TOTAL_EPOCHS),
plus slots for the 3 static configs.
Rationale:
- Unified adaptive machinery, no CPU-side special cases.
- Per-sample granularity available (Expected SARSA τ in c51_loss_kernel
is already the exemplar — reads ISV q_gap + health per sample).
- Zero CPU→GPU config transfer in any path.
- ISV is single source of truth for every adaptive value.
Plan 1 Tasks 8-17 restructured:
- Task 8 creates AdaptiveMonitor trait (not AdaptiveController).
- Tasks 9, 10, 11, 13, 14, 17: GPU kernel + monitor pairs.
- Tasks 12, 15, 16: static ISV writes only.
Follow-up commits will revert Batch A (d76849f31) and Batch B (4189da563)
which implemented the old CPU-compute pattern, then re-implement under
this design.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Unified protocol for every adaptive controller in the DQN trainer.
FireRateStats for controller_activity smoke. DiagSnapshot for
standardised HEALTH_DIAG emission. IsvBus abstraction over pinned
device-mapped memory (&mut [f32] wrapper).
No concrete controller migration yet — Tasks 9-17 migrate existing
controllers (atoms → gamma → Kelly → cql_alpha → tau → epsilon →
conviction_floor → plan_threshold → balancer) one per sub-commit.
Old scaffolding for each controller is removed in the SAME commit
that migrates its last consumer to the new protocol, per
feedback_no_partial_refactor.md.
Tests: 2 unit tests on FireRateStats and DiagSnapshot.
Plan 1 Task 8. Spec §4.C.6.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Enumerated every DtoH/HtoD/memcpy call in the DQN hot path
(run_full_step → child graphs). Classified 55 call sites:
31 OK-pinned, 20 COLD-PATH, 4 MIGRATED.
Fix 1 (gpu_dqn_trainer.rs): removed dead cuMemcpyDtoHAsync_v2 in
run_causal_intervention_unconditional — result (causal_mean_scratch)
was never consumed; now stays on device.
Fix 2 (fused_training.rs + training_loop.rs): compute_iqr() was
called inside submit_aux_ops (captured aux_child graph). Its sync
cuMemcpyDtoH_v2 cannot be graph-captured — silently ran only during
capture, then HtoD replayed stale IQR data on every step. Removed
from submit_aux_ops; added refresh_iqn_iqr() called once per epoch
in process_epoch_boundary.
Fix 3 (gpu_iqn_head.rs): tau_buf (CudaSlice<f32>) + tau_host (f32) +
cuMemcpyHtoDAsync_v2 each step replaced by tau_pinned (*mut f32) +
tau_dev_ptr (u64) via cuMemAllocHost_v2(DEVICEMAP) +
cuMemHostGetDevicePointer_v2. CPU writes *tau_pinned = tau; EMA
kernel reads via tau_dev_ptr — zero PCIe overhead. Drop updated.
Audit table populated in docs/dqn-gpu-hot-path-audit.md.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Populated docs/dqn-wire-up-audit.md with every pub module and CUDA
kernel in the DQN path. Each entry classified Wired / Partial / Orphan /
Ghost / OUT-of-DQN-scope with the action plan linking to the plan+task
that resolves any non-Wired status.
No Orphan left unclassified. Orphans fall into three buckets:
1. Scheduled for wiring by a later Plan (gpu_statistics → Plan 2 D.2;
tlob_loader → Plan 2 D.8).
2. OUT-of-DQN-scope because supervised consumers exist (PPO kernels,
xLSTM, KAN trainable adapter, flash_attention, benchmarks).
3. Genuinely unused — escalated to user review in the task output,
not deleted autonomously (streaming_dbn_loader, unified_data_loader,
training/orchestrator, training_pipeline, inference_validator,
model_loader_integration, paper_trading/mod.rs,
portfolio_transformer, regime_detection/mod.rs).
Summary: 109 total modules/kernels, 74 wired, 7 partial, 11 orphan,
0 ghost, 17 OUT-of-DQN-scope.
Plan 1 Task 6. Spec §4.A.5, Invariant 2.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Implements spec §4.A.2 structural layout fingerprint with tail placement
rather than head placement (spec alternative: §4.A.2 Step 5.3 alt).
Head placement (ISV[0..2)) was rejected because isv_signals[0] and [1]
are actively written by the isv_signal_update kernel (Q-drift EMA and
gradient-norm EMA). Shifting those would require updating every literal
reference in experience_kernels.cu — a larger change than warranted for
pure contract enforcement. Tail placement leaves all existing indices
intact, touches zero kernel .cu files, and fulfils the same design contract.
Key changes:
- ISV_LAYOUT_FINGERPRINT_LO_INDEX = 37, HI_INDEX = 38 (u64 across 2×f32).
- LAYOUT_FINGERPRINT_CURRENT: u64 = FNV-1a of slot-list seed bytes.
Value: 0x85d4d76b578a7c17. Any slot change updates seed bytes,
which updates the hash automatically.
- Constructor writes fingerprint after zero-init; calls
check_layout_fingerprint() to self-verify before returning.
- check_layout_fingerprint(): reads pinned slots [37..39), recomposes u64,
fails-fast on mismatch with "retrain required" message.
- Error message does NOT mention migration as an option.
- Pre-commit hook rejects `fn migrate_isv|upgrade_isv` names — makes the
no-migration rule structurally enforced (check_no_isv_migrations).
- ISV_TOTAL_DIM: 37 → 39.
- Zero existing index shifts (no kernel literal sites affected).
- StateResetRegistry entry renamed ISV_SCHEMA_VERSION → ISV_LAYOUT_FINGERPRINT.
- ResetCategory::SchemaContract docstring updated to remove "migration" framing.
- docs/isv-slots.md: updated table + design note for tail placement.
Tests: state_reset_registry 3 unit tests pass with renamed entry.
cargo check -p ml clean (pre-existing warnings only).
Plan 1 Task 5. Spec §4.A.2 (tail-placement alternative).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
User raised a backward-compat concern: integer "schema version"
semantically implies a family of coexisting versions with upgrade
paths between them, inviting the forbidden pattern
(feedback_no_legacy_aliases.md, feedback_no_partial_refactor.md).
Replace with a compile-time structural fingerprint:
- ISV[0..2) stores a u64 FNV-1a hash of the slot layout (split across
two f32 lanes preserving raw bits).
- The hash is computed by a const fn over the slot list; any slot
change automatically updates the fingerprint. No human decides
"what version is this now".
- Checkpoint load is fail-fast only. Error message does NOT mention
migration as an option.
- Pre-commit hook rejects any `fn migrate_isv|upgrade_isv` to make
the no-migration rule structurally enforced (landed as part of
Plan 1 Task 5 hook extension).
- StateResetRegistry entry renamed ISV_SCHEMA_VERSION →
ISV_LAYOUT_FINGERPRINT (Task 2's landed code touched in Task 5's
commit per no-partial-refactor).
Updated:
- spec §4.A.2 (fingerprint design + rationale)
- spec §5 landing-order note (fingerprint auto-updates on layout change)
- Plan 1 architecture line
- Plan 1 Task 2 StateResetRegistry test + entry naming
- Plan 1 Task 5 — full rewrite of implementation steps
- Plan 1 exit criteria #6
- Plan 2 pre-plan verification (grep fingerprint constants, not version==0)
- Plan 2 Task 6D.1 (update fingerprint seed, not bump version)
- Plan 2 Task 6D exit criteria #7
- Plan 3 pre-plan verification
- docs/isv-slots.md ISV[0..2) row
No code changed. Plan 1 Task 5 is not yet implemented — this commit
realigns the spec + plans so the implementer subagent works from the
corrected design.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Define MAG_QUARTER=0, MAG_HALF=1, MAG_FULL=2, NUM_MAGNITUDES=3 in
state_layout.cuh. Migrate the one unambiguous semantic comparison in
trade_physics.cuh (compute_target_position_4branch magnitude scaling).
Arithmetic operations on mag_idx (mag_idx±1, 2-mag_idx, mag_idx<2) are
runtime values, not semantic comparisons, and are not migrated.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Define DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3, NUM_DIRECTIONS=4 in
state_layout.cuh. Migrate all literal dir_idx/raw_dir comparisons to named
constants across experience_kernels.cu (15 sites), trade_physics.cuh (3 sites),
and backtest_metrics_kernel.cu (2 sites). Raw integer comparisons (== 0/1/2/3)
eliminated from all direction-branch consumers.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Add BRANCH_DIR/BRANCH_MAG/BRANCH_ORD/BRANCH_URG/NUM_BRANCHES to
state_layout.cuh. Migrate the 4 clear literal branch-index accesses
in branch_grad_balance_kernel.cu (branch_norms_dev[0..3] in
grad_balance_isv_update).
Other branch-keyed arrays (branch_starts, branch_lens, branch_norms,
branch_scales) use the runtime `branch = blockIdx.y` variable or loop
counter — no raw literal index, so no migration needed. Rust call sites
use struct fields (branch_0_size etc.) or loop vars — not literals.
No behavioural change; pure refactor.
Plan 1 Task 4D. Spec §3 Invariant 8.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Add PLAN_PARAM_* constants to state_layout.cuh and replace raw
plan_params/pp slot accesses in experience_kernels.cu and
backtest_plan_kernel.cu.
Sites migrated: pp[0..5] in both files (12 sites), plan_params_ptr
indexed accesses at slots 0, 1, 4 in experience_kernels.cu (4 sites),
plan_params[w*6+4] in backtest_plan_kernel.cu (1 site). Loop variable
pp[p] in noisy-plan path left as-is (p is a runtime loop counter, not
a literal slot index).
No behavioural change; pure refactor.
Plan 1 Task 4C. Spec §3 Invariant 8.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Add PLAN_ISV_* constants to state_layout.cuh and replace raw
plan_isv[N] and pisv[N] accesses in experience_kernels.cu and
backtest_plan_kernel.cu. backtest_plan_kernel.cu gains an
#include "state_layout.cuh" so the constants are visible.
6 code sites migrated (plan_isv[0..5] in experience_kernels.cu),
plus 6 pisv[N] local-array accesses in backtest_plan_kernel.cu
(same semantic slots).
No behavioural change; pure refactor.
Plan 1 Task 4B. Spec §3 Invariant 8.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replace raw ps[N] access across all CUDA kernels with PS_* named
constants defined in state_layout.cuh. 32 constants total covering
the full 38-slot portfolio state: position/cash/value, DSR stats,
Kelly stats, plan fields, and OFI scratch range.
Notable deviations from pre-written plan mapping: the plan's
PS_VALUE/PS_POSITION/PS_CASH values (0,1,2) had the wrong semantics.
Code wins (feedback_trust_code_not_docs): ps[0]=position, ps[1]=cash,
ps[2]=portfolio_value. All 148 raw ps[digit] accesses in
experience_kernels.cu and trade_stats_kernel.cu migrated. In-comment
range references (e.g. "ps[30..37]") left as documentation.
trade_stats_kernel.cu: migrated base+14 offset to PS_KELLY_WIN_COUNT.
docs/dqn-named-dims.md: PS_* table corrected to match actual code layout.
No behavioural change; pure refactor.
Plan 1 Task 4A. Spec §3 Invariant 8.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces the ad-hoc scattered fold-boundary reset calls with a single
registry-driven iteration. Adding new fold-reset state now requires
adding a registry entry AND a dispatch arm in reset_named_state in the
same commit (Invariant 2 Wire-It-Up).
Step 3.4 correction: plan_state entry removed from the registry —
plan_state_buf exists only in GpuBacktestEvaluator (val path), not in
the training-path fused ctx. No training-side fold reset is applicable.
New behaviour: isv_learning_health, isv_sharpe_ema, isv_q_means are
now properly reset to baseline at each fold boundary (previously unset,
which allowed signals from fold N to bias fold N+1 initialisation).
Tests: 3 registry unit tests pass; cargo check -p ml clean (8 pre-existing
warnings only, no new).
Authority: spec §4.A.1. Plan 1 Task 3.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Typed classification of every piece of training state by reset
lifecycle: FoldReset, WindowReset, SoftReset(decay_bars),
TrainingPersist, SchemaContract.
Replaces the previous vibes-driven fold-boundary reset logic with an
explicit registry. Adding new state requires adding an entry in the
same commit (Invariant 2 Wire-It-Up).
Consumer wiring (reset_fold_state, reset_window_state helpers that
iterate the registry) follows in Task 3 of Plan 1.
Tests: 3 unit tests covering category lookup, fold-reset filtering,
unknown-name handling.
Authority: docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §4.A.1
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan 1 Task 1. Creates the five audit docs plus config/metric-bands.toml
that track Invariants 2, 7, 8 per the DQN v2 spec, and extends the
pre-commit hook with two checks:
- component-adding commits must touch an audit doc (Invariant 7)
- added code may not contain TODO/FIXME/XXX/HACK/TBD/unimplemented!/
todo! markers (Invariant 9)
Tests: manually verified by staging a TODO-marked file; commit
rejected with the correct error message.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Self-review caught that the CqlAlphaSeedCoupledController read_signals
example used unimplemented!("plumbed by caller"), which violates the very
Invariant 9 the plan enforces. Replaced with a concrete implementation
that reads SEED_FRACTION_EMA_INDEX from the ISV bus (a Plan 1 reserved
controller-signal slot).
No other plan had placeholder violations. The remaining TBD/TODO/FIXME
grep hits across Plans 1-5 are either the pre-commit hook's rejection
patterns (correctly quoted) or enforcement scans that DETECT the markers
in audit docs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Fifth and final plan decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §2, §4.A.3,
§4.A.4, §4.A.4.1). No new policy mechanisms — orchestration, automation,
regression detection, and the final sign-off run.
Covers spec sections:
- §4.A.3 Multi-seed × multi-fold Argo harness (N=5, K=6 defaults via
--multi-seed / --folds flags on argo-train.sh)
- §4.A.4 Regression detection hard-stop (2N consecutive error-band
metrics → self-SIGTERM with TerminationReason)
- §4.A.4.1 nsys profile harness with --profile flag and
compare-nsys-profiles.py for 20% per-epoch regression detection
- §2 Tier 1 + Tier 2 + Tier 3 final exit criteria via scripts/validation/
suite (tiered tier1/tier2/tier3 check scripts + wrapper)
- §2 DQN v2 rollout close-out (spec marked LANDED, retrospective,
plan-level validation docs archived)
Plan structure: 6 tasks — 4 infrastructure (harness, regression-detect,
nsys, validation-scripts) + final-run task + rollout-closeout task. The
final-run task is the sole exit-gate for the DQN v2 rollout: ALL THREE
tiers must PASS on the first final-validation run (per stop-on-anomaly:
retrying with different seeds hoping for a different outcome is anti-
pattern).
Preserves all 9 invariants. Every mechanism wired, no stubs, no TODO/FIXME.
Retrospective written from actual implementation experience after Plans
1-5 all land.
This is the terminal plan. Plans 1-5 together constitute the complete
DQN v2 rollout. After Plan 5's final-run exit passes, the spec is marked
LANDED and the rollout is declared complete.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Fourth of five sequential plans decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §4.E).
Covers the six Part E IN decisions:
- §4.E.1 TFT Variable Selection Network across 6 feature groups
(market/OFI/TLOB/MTF/portfolio/plan_isv) with vsn_feature_selection kernel
- §4.E.2 Gated Residual Network — CONDITIONAL branch (ADOPT vs CANONICALISE)
based on docs/ml-supervised-to-dqn-concept-audit.md row
- §4.E.3 Multi-quantile IQN heads (5/25/50/75/95) replacing single CVaR output
- §4.E.4 Encoder-Decoder separation — explicit StateEncoder + per-branch
ValueDecoder with D.3 horizon-decomposed V_short/V_long sub-heads
- §4.E.5 Attention-weight interpretability — 7 new ISV slots [65..72) for
per-group attention focus EMAs + Mamba2 retention proxy
- §4.E.6 Multi-task auxiliary heads — next-bar return MSE + 5-bar regime CE,
ISV-coupled aux-weight schedule (sharpe-reactive)
ISV_TOTAL_DIM seals at 72 with this plan's allocations. Part E audit doc
closes out all rows (zero TBD/evaluate remain per Invariant 9).
Plan structure: 8 tasks (6 feature tasks + audit close-out + validation).
TDD-disciplined steps. Task 2 documents the CONDITIONAL branch decision
pathway explicitly (ADOPT vs CANONICALISE Branch A/B structure).
Plan 4 exit gate: Tier 1 convergence RETAINED (no regression vs Plan 3
baseline) + aux heads produce measurable signal. Tier 2 + Tier 3 checked
by Plan 5.
Preserves all 9 invariants. Zero stubs. Zero TODO/FIXME. Every Part E item
either lands fully or is explicit OUT (xLSTM, KAN); no third path.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Second of five sequential plans decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md).
Covers spec sections:
- §4.C.1 Quantile-based atom support (8 new ISV slots, q_quantile_reduce kernel)
- §4.D.1 Mamba2 backward pipeline completion (no atomicAdd — per-sample arrays + host reduce)
- §4.D.2 Per-branch gamma via AdaptiveController (4 new ISV slots)
- §4.D.5 Soft fold-boundary transitions (extends StateResetRegistry with SoftReset category)
- §4.D.7 Liquid Time-constant audit (trace fire rate, decide wire-or-delete)
- §4.D.3 + §4.D.6 + §4.D.8 coordinated state-layout migration (horizon-decomposed V +
plan_isv[6] + TLOB integration with atomic ISV schema version bump 1→2)
Plan structure: 7 tasks + pre-plan verification, all TDD-disciplined with
bite-sized steps (test → fail → implement → pass → commit). Task 6 is
explicitly atomic to preserve Invariant 5 (state-layout consistency).
Plan 2 prerequisites: Plan 1 landed (StateResetRegistry, AdaptiveController
trait, ISV schema version, audit docs). Plan 2 exit: all 9 invariants
preserved; convergence-scaffolding run passes with 16 new ISV slots
populated; ISV schema version == 2.
Preserves all 9 invariants. No stubs. No TODO/FIXME. Every new module
wired to production path in the same task it lands in.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
First of 5 implementation plans for the DQN v2 unified integrated
policy system spec (d13b53586, 336ee40b9). Plan 1 covers:
Task 1: Audit doc scaffolding + pre-commit hook (Invariant 7, 9)
Task 2: StateResetRegistry definition (A.1)
Task 3: Wire StateResetRegistry into fold-boundary reset (A.1)
Task 4: Named-dimension refactor — ps[], plan_isv[], plan_params[],
branch indices, dir/mag sub-indices (Invariant 8)
Task 5: ISV schema version at ISV[0], fail-fast on checkpoint
mismatch (A.2)
Task 6: Orphan audit populated (A.5)
Task 7: Hot-path purity audit populated; MIGRATE calls fixed (A.6)
Task 8: AdaptiveController trait + harness (C.6)
Tasks 9–17: Migrate each adaptive controller to the trait in spec-
specified order (atoms → gamma → Kelly → cql_alpha → tau
→ epsilon → conviction_floor → plan_threshold → balancer)
Task 18: Plan 1 validation run (smoke + 3-epoch L40S)
Plan 1 landing criteria: all 9 invariants preserved across every
commit, all smoke tests pass, audit docs fully populated. Plan 2
(temporal core) starts only after Plan 1 passes exit criteria.
Plans 2–5 cover Parts B, C (minus C.6 already in Plan 1), D, E, and
final validation. Planned as separate documents in docs/superpowers/
plans/2026-04-24-dqn-v2-plan-{2..5}-*.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
User-mandated addition:
Invariant 9 — No deferred work, no stubs, no TODO/FIXME
Every commit lands complete. No stubs (placeholder return values), no
TODO/FIXME/XXX/HACK/TBD markers, no half-finished implementations. If
it can't finish in this commit, it doesn't start. Stubs and deferrals
train the network on semantic emptiness — invisible in convergence
metrics but burn GPU time optimising against partly-fake signal.
Authority: feedback_no_stubs.md, feedback_no_todo_fixme.md,
feedback_no_quickfixes.md elevated to first-class spec enforcement.
Pre-commit hook greps for forbidden markers.
Part E decisions made concrete (per Invariant 9):
- E.1 TFT VSN: IN (full VSN extension)
- E.2 GRN: IN if A.5 audit finds absent (otherwise mark existing as canonical)
- E.3 Multi-quantile heads: IN (5/25/50/75/95 quantile decomposition)
- E.4 Encoder-decoder: IN (extends D.3 to full trunk/head separation)
- E.5 Interpretability: IN (attention-weight ISV diagnostics)
- E.6 Auxiliary heads: IN (next-return + regime-classification)
- xLSTM: OUT (redundant with Mamba2+TLOB; YAGNI)
- KAN: OUT (function approx not a bottleneck)
No "evaluate later" items. Every concept either lands or explicitly
doesn't, with rationale.
§8 decisions tightened:
- D.8 TLOB mode: decision criterion = per-step cost benchmark ≤ 5ms
- C.6 controller order: fixed in spec (atoms→gamma→Kelly→cql_alpha→
tau→epsilon→conviction_floor→plan_threshold→balancer)
- A.3 resource plan: auto-switch L40S→H100 at 12 GPU-hour threshold
- A.4 metric band init: [mean − 3σ, mean + 3σ] from last 3 good runs
No discretionary "we'll see" remain. All decisions data/order/budget/
statistic-driven.
Counter updates: "seven invariants" → "nine" in §3 header and §5
cleanup contract.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Self-review pass on the DQN v2 unified design spec (d13b53586). Fixes
applied inline:
Ambiguity fixes:
- §3 Invariant 3: explicit definition of "training step loop" (per-step
code in fused_training::step_fused and captured graph children; NOT
per-epoch / per-fold / per-checkpoint code).
- §2 Tier 1: "stable epochs" defined as epochs [warmup_end..run_end)
with warmup_end from B.3's seed-phase decay or explicit flag.
- §4 A.2: migration mechanism clarified as fail-fast initially; helpers
added only when a specific migration need arises (no speculative
scaffolding).
- §4 A.3: N=5, K=6 stated as DEFAULTS with ≥ MINIMUMS per Invariant 6;
resource budget flagged (~8-10 GPU-hours per validation pass).
- §4 B.3: "≥ 100K experiences" (minimum), warm-restart clarified as
runtime condition not feature flag.
- §4 C.3: adaptive KL threshold mechanism made explicit (second ISV
slot for threshold EMA, third for amplification multiplier).
- §4 C.6: cleanup timing explicit — old scaffolding removed in the SAME
commit that migrates its last consumer (no deferred pass).
- §4 D.1: DQN-path-specific scope clarified; validation via grad-norm
smoke + integration test.
- §4 D.6: plan_isv index naming made consistent with existing layout;
coordinated state-layout migration commit called out.
New Invariant 8 — Named dimensions, not indices:
Every dimension, slot, offset, or semantic position in a multi-field
buffer has a named constant. Raw numeric indices (`[0]`, `[6]`, `[23]`)
appear ONLY in the definition site. Named constants specified for ISV
slots (existing), portfolio state ps[0..30), plan_isv[0..7), plan_params
[0..6), state-vector offsets, branch indices (BRANCH_DIR/MAG/ORD/URG),
action sub-indices (DIR_SHORT/HOLD/LONG/FLAT, MAG_QUARTER/HALF/FULL).
Enforcement: audit pass during A.1/A.2, lint via grep for raw-index
access outside definition modules.
Landing order revision:
- Dependency graph explicit (A.2 before new ISV slots, D.1 before
D.2-4-8, state-layout changes in ONE coordinated commit).
- Spec decomposition note added — writing-plans may produce multiple
sequential plans rather than one monolithic plan.
New doc tracked by pre-commit: docs/dqn-named-dims.md (Invariant 8).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Eliminates the f64→f32 cudarc ABI trap (feedback_cudarc_f64_f32_abi.md,
task #82) at the type level: hyperparameters consumed by CUDA kernels
now live as f32 in Rust, cast once at the TOML/PSO ingest boundary
instead of at every kernel call site.
Structs changed:
- DQNHyperparameters (crates/ml/src/trainers/dqn/config.rs) —
~85 scalar fields migrated from f64 → f32. Covers all
kernel-facing scalars: reward weights (w_pnl/w_dd/w_idle,
dd_threshold, cea_weight, micro_reward_*, price_confirm_weight,
book_aggression_weight, hold_quality_weight), exploration
(epsilon_* and the 4 branch mults, noisy_sigma_*, count_bonus,
noise_sigma, q_gap_threshold), distributional RL (v_min, v_max,
reward_scale, iqn_lambda, qr_kappa, spectral_*,
gradient_collapse_multiplier), fill simulation (5 fill_*
fields), risk/Kelly (kelly_fractional, kelly_max_fraction,
max_leverage, max_position_absolute, minimum_profit_factor),
ensemble/curiosity (curiosity_weight,
curiosity_q_penalty_lambda, ensemble_*, beta_*, variance_cap),
anti-LR (anti_lr_*, adversarial_dd_threshold,
beta_penalty_strength), walk-forward (wf_*),
experience (avg_spread, transaction_cost_multiplier,
holding_cost_rate, churn_penalty_scale, contract_multiplier,
margin_pct, tick_size, bars_per_day, cash_reserve_percent),
misc kernel scalars (gamma, tau, huber_delta, q_clip_*,
shrink_perturb_*, regime_replay_decay, per_alpha,
per_beta_start, dt_target_return, etc.). Also
`noisy_epsilon_floor: Option<f32>` and
`count_bonus_coefficient: Option<f32>`.
- `computed_v_min` / `computed_v_max` now return f32.
- `compute_max_position` returns f32 (f64 internally for the
notional division).
Fields preserved as f64 (precision-sensitive, NOT kernel-facing
scalars — per task spec and feedback_cudarc_f64_f32_abi.md):
- `learning_rate` — tested at 1e-10 tolerance; f32 rounds to
2e-12 for a 1e-5 LR.
- `entropy_coefficient` — tested at 1e-9 tolerance.
- `weight_decay` — tiny 1e-5..1e-3 range.
- `adam_epsilon` — 1e-8 default; f32 preserves denorms here but
paired with weight_decay/learning_rate for symmetry.
- `gradient_clip_norm: Option<f64>` — grad norms are f64
accumulators by project convention.
- `min_loss_improvement_pct`, `q_value_floor` — early-stopping
long-horizon stats.
- `cql_alpha` — flows into DQNConfig (ml-dqn) still f64.
- `min_learning_rate`, `lr_min` — paired with learning_rate.
- Family intensity scalars (6 `*_intensity` fields) — f64 PSO
search space; the intensity applies via `as f32` at each call
site in `apply_family_scaling`.
No checkpoint format change: DQNHyperparameters has
`#[derive(Debug, Clone)]` only (not Serialize/Deserialize), so the
TOML-ingest path is the only serde boundary and already casts
explicitly via `hp.field = v as f32;` in
`DqnTrainingProfile::apply_to`. PSO hyperopt bounds stay f64 in
`SearchSpaceSection` and cast at the adapter boundary.
Call-site impact:
- ~30 `as f32` casts removed from hot paths (fused_training
FusedConfig builder, training_loop kernel launches,
constructor DQNConfig builder, action.rs GPU action selector,
trainer/mod.rs WF config). Kernels now receive the hp field
directly via `&hp.x`.
- ~75 `as f32` casts added at the `apply_to` / hyperopt-adapter
ingest boundary — the single conversion point.
- Cross-crate contracts (DQNConfig in ml-dqn, PortfolioTracker
in ml-core, GAECalculator, DropoutScheduler, NoisySigmaScheduler,
KellyOptimizerConfig, RewardConfig) retain their f64 signatures;
ml calls cast at the boundary with `f64::from(hp.x)` so the
contract is explicit and greppable.
Two test-side adjustments:
- `test_kelly_fields_are_public` now asserts `f32` for
kelly_fractional / kelly_max_fraction (these migrated).
- `test_early_stopping_termination` uses `f64::from(...)` to
preserve the f64 threshold computation against the now-f32
`gradient_collapse_multiplier`.
Verified:
- `SQLX_OFFLINE=true CARGO_INCREMENTAL=0 RUSTC_WRAPPER=sccache
cargo check --workspace` — clean, zero new warnings.
- `cargo check --workspace --tests` — clean.
- `cargo test -p ml --lib training_profile::tests` — all 18
migrated tests pass. The one pre-existing failure
(`test_production_profile_applies_all_sections` n_steps
mismatch, 1 vs 5) reproduces on stashed baseline, so it is
unrelated to this change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Final commit toward task #94 val plan_isv parity. Wires the plan
machinery inside the chunked val backtest loop so state positions
[86..92) carry real plan signals matching training instead of zero-
fill. This is the structural distribution-shift fix behind the
observed val trade count of 21 / 214,654 bars (0.0098%).
Phase 6 inside `submit_dqn_step_loop_cublas`, per chunk, after env
step advances portfolio:
1. Forward on the LAST-STEP N rows of chunked_states — a targeted
`compute_q_values_to(last_step_states, N, scratch_q)` rewrites
`save_h_s2` to exactly [N, SH2] for the final step. Needed
because the batched `compute_q_values_to` for N*chunk_len rows
leaves save_h_s2 at an arbitrary last-sub-batch slice.
2. `compute_plan_params(plan_params_buf, N)` runs the trade-plan
MLP on that clean save_h_s2, producing [N, 6] plan_params.
3. `backtest_plan_state_isv` — Flat↔Positioned activation /
deactivation on plan_state[N, 7] and writes plan_isv_buf[N, 6]
consumed by the next chunk's first `launch_gather` for state
positions [86..92).
ch_q_values is reused as the Q-scratch for step (1) — its original
chunk Q-values were already consumed by Phase 4 action-select.
Epoch-boundary diagnostic `launch_plan_diag_and_log` after the step
loop: launches `backtest_plan_diag_reduce` (1-block, 256 threads),
DtoH-copies the 8-float summary, syncs, and emits a single
`tracing::info!(target: "val_plan_diag", ...)` line with the six
labelled plan_isv means plus active_frac / n_active.
Scalars passed with exact i32 types (current_step, max_len,
feat_dim, n_windows) per feedback_cudarc_f64_f32_abi. Kernels
pulled via `plan_state_isv_kernel.as_ref().ok_or_else(...)` — they
are loaded in `ensure_action_select_ready` on first eval.
Not captured inside a CUDA Graph: the backtest step loop uses direct
kernel launches (evaluator comment: graphs SIGSEGV on 500K+ launches).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Second commit toward task #94 val plan_isv parity. Adds structural
wiring in GpuBacktestEvaluator:
- plan_params_buf [N, 6] : live plan MLP output per step
- plan_state_buf [N, 7] : persistent stored plan across bars
- plan_diag_buf [8] : epoch-end diagnostic reduction output
- plan_state_isv_kernel : lazy-loaded CudaFunction from
backtest_plan_kernel.cubin
- plan_diag_reduce_kernel : diagnostic reduction kernel
Buffers allocated zero-initialised in the constructor; kernels loaded
alongside the action_select and scatter_intent kernels in
`ensure_action_select_ready` (same module-load pattern).
Remaining work (follow-up commit):
- Wire QValueProvider::compute_plan_params into the chunked loop to
populate plan_params_buf from the trainer's save_h_s2 after each
forward pass.
- Launch backtest_plan_state_isv after the env_batch_kernel each
chunk to update plan_state_buf and plan_isv_buf for the next
chunk's state gather.
- Launch backtest_plan_diag_reduce + DtoH of plan_diag_buf at epoch
boundary for HEALTH_DIAG emission.
The chunked batch layout (N*chunk_len samples per forward pass)
requires careful slicing to extract the last-step's plan_params for
the plan_state update — handled in the follow-up commit to avoid
rushing a subtle integration.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Val-Flat-collapse fix#5 infrastructure (task #94, 2026-04-24). The
val backtest's `plan_isv_buf` was zero-filled, causing state positions
[86..92) to be OOD relative to training and biasing val argmax toward
Flat/Hold on 99.99% of bars. Research-agent audit confirmed this is a
documented architectural gap, not a bug per se: val's portfolio state
has no plan slots (8 vs training's 30+), so there was no mechanism to
carry plan signals across bars in val.
This commit lays the foundation for closing that gap:
1. `backtest_plan_kernel.cu` (new) — two kernels:
- `backtest_plan_state_isv`: per-window Flat→Positioned plan
activation + plan deactivation on return to Flat + plan_isv[0..5]
computation matching training's formula at
experience_kernels.cu:596-619. Reads live plan_params from the
plan MLP, writes plan_state (persistent across bars) and
plan_isv_out. Extracts raw_close from features buffer for
unrealized P&L ratios.
- `backtest_plan_diag_reduce`: single-block reduction emitting 8
diagnostic floats for plan activity logging (active fraction,
per-slot mean magnitudes, raw active count).
2. `QValueProvider::compute_plan_params` trait method — allows the val
evaluator to run the plan MLP on the trainer's most recent trunk
hidden state (save_h_s2), filling plan_params[N, 6] for use in the
state-update kernel.
3. `FusedTrainingQValueProvider::compute_plan_params` impl — runs
`launch_trade_plan_forward` then DtoD-copies `plan_params_buf` to
the caller's output in the same chunk-loop pattern as
`compute_q_values_to`.
4. build.rs — register `backtest_plan_kernel.cu` for compilation.
Next commit will wire these into `GpuBacktestEvaluator`:
- Add `plan_params_buf[N, 6]`, `plan_state_buf[N, 7]`,
`plan_diag_buf[8]` device allocations.
- Load the two kernels into CudaFunction slots.
- Call sequence per step: `compute_q_values_to` →
`compute_plan_params` → action_select → env_step →
`backtest_plan_state_isv` (reads updated portfolio, writes
plan_state + plan_isv_buf for next step).
- Epoch boundary: `backtest_plan_diag_reduce` + DtoH + HEALTH_DIAG
emission.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
