foxhunt/docs/dqn-wire-up-audit.md

DQN v2 Wire-Up Audit

Status: Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.

data_source globalization to mbp10 (2026-04-27): completes the alignment started in the earlier "fxcache data_source alignment" entry below. Smoke (dqn-smoketest.toml) and localdev (dqn-localdev.toml) profiles previously set data_source = "ohlcv" while production sets "mbp10". The split forced two separate fxcache files and made cross-environment runs require explicit --data-source ohlcv overrides on the precompute binary. Default rolls forward to "mbp10" everywhere it is not deliberately overridden:

• config/training/dqn-smoketest.toml, dqn-localdev.toml: profile values
• crates/ml/src/training_profile.rs, trainers/dqn/config.rs, hyperopt/adapters/dqn.rs: Rust defaults + doc strings
• crates/ml/examples/precompute_features.rs: doc updated; --data-source mbp10 is the canonical regen invocation
• crates/ml/src/feature_cache.rs, fxcache.rs test fixtures use mbp10 Documentation references to the alternative ("ohlcv") are retained only where the docstring explicitly enumerates the two valid choices. Local fxcache regenerated to v6 mbp10 (13c0b086a975cc7e2384377a2cd0e97738c9410292fcfecb5807c29bf885cb48.fxcache, 55 MB, 175874 bars). Stale v5 ohlcv cache untracked from git index (already gitignored post-79578bbaf).

MoE Phase 3 wire-up T3.1–T3.7 (2026-04-27): MoE fully wired into production training path. Forward: gate (state[B,128]→64→8→softmax) + 8 expert MLPs (h_s1[B,256]→64→256) + moe_mixture_forward → replaces save_h_s2 in submit_forward_ops_ddqn. Load-balance: moe_load_balance_loss

• moe_load_balance_reduce + SAXPY into total_loss_dev_ptr (T3.4). ISV producer: launch_moe_expert_util_ema per-step in training_loop.rs (T3.5). Backward: moe_mixture_backward (de_k=g·dh_s2) + moe_dgate_reduce (dg=Σe·dh_s2) + moe_softmax_backward + cuBLAS SGEMM backward through gate W2/W1 and 8 experts W2/W1 — all into params_buf grad slots [127..163) (T3.6). Adam step inherits gate+expert updates automatically (params_buf uniform). HEALTH_DIAG aux_moe line emitted per epoch from ISV[118..127) (T3.7). Smoke test: 3/3 folds pass, 728s. Gate differentiated: expert-2 reached 32.3% utilization (others 9.7%) by fold-2 epoch-4; entropy 1.611 < ln(8). fused_training.rs: added moe_lambda field to GpuDqnTrainConfig init from hyperparams.moe_lambda.unwrap_or(0.01).

MoE expert util EMA T2.4 (2026-04-27): moe_expert_util_ema_update single-thread cold-path-cadence kernel writes 8 per-expert utilization EMAs (ISV[118..126)) + gate-entropy EMA (ISV[126]) with α=0.05. Same shape as h_s2_rms_ema_update/aux_heads_loss_ema_update. GPU-resident per pearl_cold_path_no_exception_to_gpu_drives.md. Test verifies EMA values match CPU reference within 1e-5 using skewed gate (expert-3 = 0.6).

MoE load-balance loss T2.3 (2026-04-27): moe_load_balance_loss (one block per k, shmem reduction over B) computes λ·K·(mean_b g[b,k])² per expert without atomicAdd; moe_load_balance_reduce (single-thread scalar sum) accumulates K terms. Two tests: CPU-reference match (1e-5) and uniform-gate minimum (loss = λ). Backward gradient lands in Phase 3.

MoE backward kernels T2.2 (2026-04-27): moe_mixture_backward computes de_k[k,b,c] = g[b,k] · dh_s2[b,c] (one thread per (k,b,c)); moe_dgate_reduce computes dg[b,k] = Σ_c e_k[k,b,c] · dh_s2[b,c] via shmem-tree reduction over c (one block per (b,k)). Finite-differences test verifies analytic backward matches numerical for B=2, K=4, C=32 within 1e-3. All data flows via mapped pinned buffers. Consumers land in Phase 3 wire-up.

MoE ISV slot reset registration (2026-04-27): registered fold-boundary reset entries for the 8 MOE_EXPERT_UTIL_EMA slots (reset to 1/K=0.125) and the MOE_GATE_ENTROPY_EMA_INDEX slot (reset to ln(8)). Producers land in Phase 2 task 2.4; consumers in Phase 3.

MoE ISV slot reservation (2026-04-27): reserve slots 118–126 for the upcoming MoE redesign monitoring — 8 slots MOE_EXPERT_UTIL_EMA[0..8] for per-expert gate-weight EMA + 1 slot MOE_GATE_ENTROPY_EMA_INDEX=126 for the gate-distribution entropy EMA. ISV_TOTAL_DIM 118 → 127. Producers + consumers land in subsequent commits. Spec docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §4.6.

use_* flag family removal + count_bonus API tightening (2026-04-27): atomic deletion of 6 dead use_* boolean flags from DQNConfig per feedback_no_feature_flags.md. (1) use_dueling, use_distributional, use_noisy_nets, use_branching were pure-cosmetic always-on flags; only metadata strings remained. Their if true /* always on */ patterns at 4 call sites in dqn.rs::select_action, select_action_with_confidence, select_action_inference, q_values_for_batch collapsed to the live branching-only arm; the dead else arms (legacy non-branching code paths referencing 5-flat ExposureLevel layout) deleted. (2) use_count_bonus was redundant with the live count_bonus_coefficient numeric kill-switch; gating removed, recording made unconditional, single coefficient drives behavior. (3) use_cvar_action_selection was dead post-use_iqn cleanup (only consumers were inside the IQN arms removed in commit da632446c); field gone, cvar_alpha retained for cuda_pipeline CVaR loss usage and as the future dial for CVaR action selection on the live C51 distribution. (4) count_bonus.rs gained bonuses_branched_f32(&self, exp_out: &mut [f32], ord_out: &mut [f32], urg_out: &mut [f32]) write-into API alongside the existing Vec<f64>; the production DQN::get_count_bonuses_branched() returns fixed ([f32; 4], [f32; 3], [f32; 3]) arrays — allocation-free, f32 throughout, fed directly to the GPU action selector. Bit-identical Test 0.F outputs vs pre-cleanup confirm zero behavioral change (legacy use_* arms were unreachable, as expected). This commit is the precondition for the MoE regime redesign per docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md.

use_iqn flag + dead iqn_network removal (2026-04-27): structural deletion of a feature flag (feedback_no_feature_flags.md) that gated unreachable legacy code. Production training already runs IQN unconditionally through cuda_pipeline/gpu_iqn_head.rs + iqn_dual_head_kernel, properly wired into the branching architecture with the FIXED_TAUS 5-quantile schedule; nothing in the cuda_pipeline production path ever read use_iqn or iqn_network. The legacy DQN::iqn_network: Option<QuantileNetwork> was a parallel CPU-side network from a pre-branching era — its only consumers were 6 conditional gates inside the if true /* use_branching: always on */ arm at q_values_for_batch:1808, so the IQN else-if at L1822 was dead code. That dead-zone bug (legacy comment: "IQN trains base q_network but inference uses dist_dueling network (zero gradients)") was paper-fixed by disabling the feature instead of fixing the inference path. Stripped: use_iqn + 3 vestigial fields (iqn_num_quantiles, iqn_kappa, iqn_embedding_dim — kernel-side macros are the actual config), 4 default builders, iqn_network field + init, 6 IQN gates collapsed to live arm, get_state_embedding (only consumed by deleted IQN paths), entire crates/ml-dqn/src/quantile_regression.rs module (392 LOC), 2 lib.rs exports, all checkpoint metadata I/O for the dropped fields. iqn_lambda stays — cuda_pipeline dual head consumes it as IQN aux loss weight. Downstream call sites in ml/src/trainers/dqn/{config,fused_training, trainer/constructor}.rs substitute kernel-fixed literals (64, 1.0) for the deleted DQNConfig copies; evaluate_baseline.rs and 2 test files drop their flag references. dqn_action_collapse_fix_test.rs no longer asserts !use_iqn — the dead-zone pathology is structurally impossible. Test 0.F (Plan A Task 8) ships in the same commit with its docstring rewritten to drop the use_iqn=false framing and the legacy "Tier-B-prime" caveat; the Tier-A version is GPU-integration-only per feedback_no_cpu_forwards.md (CPU is read-only, never propose a CPU forward as a future option). Net: 11 files, +458 / -785 LOC, 0 errors on cargo check --workspace --tests, Test 0.F produces bit-identical σ_C51 / argmax / Thompson outputs vs pre-removal.

fxcache data_source alignment + DBN-fallback normalization (2026-04-27): two-part fix to a class of bugs causing un-normalised features to silently flow into training. (1) precompute_features.rs previously hardcoded data_source = "ohlcv" at the cache-key write site (lines 214, 633). train_baseline_rl.rs:582 hardcoded "ohlcv" at the read site too. dqn-production.toml sets data_source = "mbp10" (MBP-10 is the canonical production data path). The hardcodes mean smoke (uses ohlcv) worked by accident, but any future profile with a different data_source silently mismatches → cache MISS → DBN-direct fallback. Both call sites now use "mbp10" (production default). precompute_features adds a --data-source CLI override for legacy ohlcv smoke flows. (2) DBN-fallback path in train_baseline_rl.rs:592-643 did NOT call NormStats::normalize_batch (precompute does, line 629). Any cache miss for any reason (data_source drift, schema-hash mismatch, missing file) silently uploaded RAW features to GPU. Raw close prices (~$5180 ES futures) flowed into next_states[:, 0], the aux next-bar head's label_scale EMA latched onto raw-price magnitude (~5443 vs expected ~1.0 z-score), and the shared trunk learned to predict next-bar prices → epoch-0 Sharpe 141 with 0.32% max drawdown over 214k bars (train-h5gxb). DBN fallback now applies the same z-score normalisation unconditionally as defence-in-depth, so a future cache-miss cannot reintroduce raw values into training.

fxcache schema-hash gap fix (2026-04-27): build.rs::emit_feature_schema_hash hashes only src/features/extraction.rs, src/fxcache.rs, and ../ml-core/src/state_layout.rs. The z-score normalization step lives in examples/precompute_features.rs:625-631 (added 2026-04-03 in commit 9f7c14978f) and was NOT covered by the hash, so a fxcache written by an older precompute build (raw features, no normalization) silently passed today's validate() because every other field matched. The PVC fxcache on the L40S Argo path was such a stale artefact: feature column 0 stored RAW CLOSE PRICES (~$5180 ES futures level) instead of z-normalized log- returns. The aux next-bar head reads next_states[:, 0] as its label (gpu_dqn_trainer.rs:7758-7789); with raw-price labels, the EMA label_scale reached 5420 vs smoke's 0.05. The shared trunk learned to predict next-bar prices, leaking future info into the policy → train-f8h6q epoch-0 reported impossible Sharpe=141.99 with 0.32% max-drawdown over 214k bars. Fix: add examples/precompute_features.rs to schema_sources. New hash invalidates the stale PVC cache; Argo's ensure-fxcache regenerate-on-failure branch (infra/k8s/argo/train-template.yaml:372-383) auto-regens with current normalized precompute. Generalises beyond this incident: any future change to feature normalization, target ordering, or precompute post-processing is now hash-tracked.

mag_stats wr_h/wr_f attribution fix (2026-04-27): experience_kernels.cu line 1916 binned action_mag_per_sample by actual_mag_core at every step, including trade-close events. unified_env_step_core forces actual_mag = 0 (Quarter) whenever actual_dir is Hold/Flat (line 772 of trade_physics.cuh), and trade closes always land in Hold/Flat state — so every Half/Full close was attributed to the Quarter bin. close_counts[Half] and close_counts[Full] structurally pinned to 0; wr_h=wr_f=0 across all training runs. Fix: at close events bin by pre_mag_bin (the magnitude of the position being closed); at non-close events keep actual_mag_core (current realized magnitude). pre_mag_bin is always 0/1/2 at close events since exiting/reversing requires prev_sign != 0 → pre_trade_position != 0. Smoke verification: wr_h/wr_f remain 0 in 5-epoch smoke because var_scale (~0.10-0.19 at smoke maturity, from 1/(1+sqrt(var_q)) shrinkage) makes even Long Full target → abs_pos ≈ 0.15 < 0.375 → all positions decode as Quarter; no Half/Full positions exist for the fix to attribute. Latent correctness for mature training (var_q drops as Q-distribution converges → var_scale grows → Half/Full positions become reachable).

actions_history_buf measurement artifact fix (2026-04-27): two-part fix. (1) gpu_backtest_evaluator.rs::reset_evaluation_state now initialises actions_history_buf to -1 sentinel via cuMemsetD32Async(0xFFFFFFFFu32) instead of memset_zeros. After done_flags[w]=1 (capital floor breach), backtest_env_step early-returns without writing actions_history for the remaining slots in [done_step, max_len); the prior zero-init decoded those as Short Quarter Market Normal (action 0) via dir = 0/27 = 0. (2) backtest_metrics_kernel.cu added if (act < 0) continue; after reading actions_history so the reduce-side metrics (buy_count/sell_count/hold_count → active_frac/dir_entropy + bnd_* trade-boundary detection) skip unwritten slots consistently with the existing Rust-side reader guards. Empirical impact on the local 3-fold × 5-epoch smoke: val_dir_dist Short dropped from 83% (artifact) to 13-29% (matches val_picked_dir_dist within 5pp); active_frac dropped from 87-91% to 31-50%; dir_entropy increased from 0.57 to 0.83-1.02. The pre-fix "val-Flat-collapse" / "Short-collapse" pathology that motivated substantial subsequent investigation was largely a measurement artifact from this bug surfacing differently before vs after the Kelly cap fix (0c9d1ee39). Pre-Kelly fix, Kelly cap clamped Long/Short → Flat, so actions_history was densely written with Flat encoding → fewer unwritten slots → 80% Flat reading (real Kelly pathology + small artifact). Post-Kelly fix, Long/Short survive but poor smoke-trained model breaches capital floor often → more unwritten Short slots → 83% Short reading (pure artifact). With both fixes, val_dir_dist reflects real model behaviour.

Plan A Phase 0 Thompson test kernel (2026-04-27): cuda_pipeline/thompson_test_kernel.cu — standalone test-only CUDA kernel implementing the Thompson direction sampling math (inverse-CDF over C51 atoms, uniform-τ over IQN quantiles, argmax of E[Q]) plus diagnostic σ kernels (compute_sigma_c51_test, compute_sigma_iqn_test). Exercised only by distributional_q_tests.rs (Phase 0 GPU-direct hypothesis verification). Phase 2 will inline the SAME __device__ __forceinline__ math into experience_kernels.cu::experience_action_select for the production direction-branch action selector; this isolated test kernel is the reference that Phase 2 Test 2.B verifies the production kernel matches bit-for-bit. No production callers in this commit. Spec: docs/superpowers/specs/2026-04-26-distributional-rl-aggregation-design.md.

Plan A Phase 0 Task 2 test wrappers (2026-04-27): trainers/dqn/distributional_q_tests.rs — test-only Rust module gated #[cfg(test)] pub mod in trainers/dqn/mod.rs. Wraps the Task 1 cubin via include_bytes!(concat!(env!("OUT_DIR"), "/thompson_test_kernel.cubin")) with two launchers (launch_thompson_direction, launch_argmax_eq) plus make_test_stream / upload helpers. Skeleton only; actual #[ignore]-gated verification tests (0.A bias-reproduces, 0.B inverse-CDF, 0.C IQN symmetry, 0.D Thompson-reverses, 0.E synthetic edge, 0.F converged-checkpoint) land in Tasks 3-8. No production callers. Plan: docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md. Code-review amendment (2026-04-27): cubin now loaded once via process-wide OnceLock<KernelSet> (matches cuda_pipeline/signal_adapter.rs and cuda_pipeline/gpu_her.rs pattern); module-level #![allow(unsafe_code)]

• #![allow(dead_code)] (helpers populated by Tasks 3-8); i32→u8 cast replaced with u8::try_from(...).expect(...) to fail loudly on a buggy kernel writing OOB; host[0] direct indexing replaced with host.first().copied().expect(...).

Plan A Phase 0 Task 3 Test 0.A (2026-04-27): trainers/dqn/distributional_q_tests.rs — appended test_0a_bias_reproduces_argmax_vs_thompson (#[test] #[ignore]). Constructs synthetic per-direction C51 (Flat/Hold = δ(0); Long/Short = Gaussian σ=0.3 around v=-0.001) plus matching IQN quantiles, and asserts (1) argmax(E[Q]) deterministically picks Flat or Hold (bias reproduces) and (2) Thompson sampling over 10 000 seeds picks Long+Short ≥ 3 000 (proposed fix would work). PASSES on local RTX 3050 Ti in ~1.86 s. First Phase 0 verification test exercising the Task 1 cubin via the Task 2 wrappers. No production callers. Plan: docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md.

Plan A Phase 0 Task 4 Test 0.B (2026-04-27): trainers/dqn/distributional_q_tests.rs — appended test_0b_thompson_c51_distribution_matches_input (#[test] #[ignore]). Single direction (d=0) carries C51 atoms p=[0.1, 0.7, 0.2] over v=[-1, 0, +1]; other directions δ(v=0); IQN all-zero. Counts d=0 wins over 100 000 seeds and asserts P(d=0) ≈ 0.9 within 1 %. The 0.9 (not the plan's originally-specified 0.2) is correct: the kernel uses strict-> tie-break with best_d=0 initialised, so d=0 wins iff its C51 sample ≥ 0, i.e. p[v=0] + p[v=+1] = 0.7 + 0.2 = 0.9. Validates the inverse-CDF math end-to-end — a bug shifting mass between v=-1 and {v=0, v=+1} would move P(d=0) measurably from 0.9. Observed 0.90195 on local RTX 3050 Ti (~3.8 s). Same commit also applies clippy erasing_op/identity_op cleanup (0 * n_atoms, 1 * n_atoms) to Test 0.A — no behaviour change in 0.A. No production callers.

Plan A Phase 0 Task 5 Test 0.C (2026-04-27): trainers/dqn/distributional_q_tests.rs — appended test_0c_thompson_iqn_quantile_interp (#[test] #[ignore]). Implemented on the batched-pinned API (single launch_thompson_direction_batched call over 100 000 seeds, host-side filter on the returned Vec<u8>) — no per-seed launches, no clone_dtoh. Setup: d=0 has standard-normal IQN quantiles [-1.645, -0.674, 0, 0.674, 1.645] at the fixed τ-anchors {0.05, 0.25, 0.50, 0.75, 0.95}; other directions all-zero. C51 is δ(v=0) for ALL directions, so the combined sample collapses to s_iqn for d=0 and exactly 0 for d=1,2,3. Strict-> tie-break with best_d=0 initialised gives d=0 the win iff s_iqn ≥ 0, so by symmetry of the standard-normal IQN, P(d=0) ≈ 0.5. Asserts (p_d0 - 0.5).abs() < 0.02. Observed 0.50003 on local RTX 3050 Ti (~1.6 s including build, ~0.2 s test wall when build is hot). Validates the uniform-τ linear interpolation in sample_iqn_quantile_interp — a wrong interpolation (anchor swap, non-linear blend, or boundary OOB) would break the symmetry and shift P(d=0) measurably off 0.5. No production callers.

Plan A Phase 0 Task 6 Test 0.D (2026-04-27): trainers/dqn/distributional_q_tests.rs — appended test_0d_thompson_reverses_bias (#[test] #[ignore]). Implemented on the batched-pinned API (single launch_thompson_direction_batched over 10 000 seeds, host-side filter().count() on the returned Vec<u8>) — no per-seed launches, no clone_dtoh. Setup: realistic synthetic failure mode with σ_long=0.05 (tighter than Test 0.A's σ=0.3, matching the σ scale we actually observe in trained checkpoints). C51: Flat/Hold = δ(0); Long/Short = Gaussian centered at -0.001, σ=0.05; 21 atoms over [-0.5, 0.5]. IQN: Flat/Hold all-zero; Long/Short standard quantiles [-0.082, -0.034, -0.001, 0.032, 0.080]. Asserts argmax picks Hold/Flat (deterministic E[Q] bias) and Thompson P(Long) ≥ 0.20 over 10 000 seeds. Observed argmax_pick=1 (Hold), P(Long)=0.36630, P(active)=0.74500 on local RTX 3050 Ti (~0.23 s wall — all four tests combined). Validates Thompson exploration is sufficient at the realistic σ (not just the wide σ from Test 0.A) — the failure-mode bias mechanism reproduces and Thompson reverses it. No production callers. Plan: docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md.

Plan A Phase 0 Task 7 Test 0.E (2026-04-27): trainers/dqn/distributional_q_tests.rs — appended test_0e_synthetic_edge_discovery_cpu (#[test], NOT #[ignore] — pure CPU, no GPU dependency, runs in default cargo test). Algorithmic-property test of Thompson exploration on a 1-state 2-action bandit: Long has KNOWN +0.005 edge (Normal(0.005, 0.03)), Flat returns deterministic 0. Q-learning over 5 atoms [-0.1, -0.05, 0, 0.05, 0.1] with lr=0.05 for 100 iterations under two action selectors. Initial setup revised from plan A draft (option 3 — production-realistic): p_long = [0.10, 0.20, 0.40, 0.20, 0.10] (uniform-ish, E=0, has σ; mimics freshly initialized C51 head); p_flat = [0, 0, 1, 0, 0] (δ(v=0); tx_cost-bare Flat returns 0). Argmax with strict-> tie-break at e_long=e_flat=0 always picks Flat → never explores Long → e_long stays 0 (assertion e_long ≤ e_flat + 0.001 passes vacuously). Thompson: `P(s_long > 0) = p_long[3]

• p_long[4] = 0.30→ ~30 effective Long-selections drift p_long mean toward the +0.005 reward target's nearest atom (+0.05). Observed values on local RTX 3050 Ti: argmax e_long=0.000000, e_flat=0.000000; Thompson e_long=0.003550, e_flat=0.000000 (test ~0.00 s; total 5-test run ~1.57 s). Plan's prior draft (initial p_long with mean=-0.015 + p_flat=δ(0)) was calibration-bound: Thompson drift was directionally correct but didn't cross zero in 100 iters. Revised setup eliminates the artificial initial bias and matches production reality. Stop condition: if Thompsone_long ≤ e_flatwith this setup, the hypothesis is genuinely wrong and reward shaping must change before Phase 2. Usesrand+rand_distr(already in workspace deps). No production callers. Plan:docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md`.

Plan A Phase 0 batched-pinned redesign (2026-04-27): cuda_pipeline/thompson_test_kernel.cu

• trainers/dqn/distributional_q_tests.rs — eliminates the per-seed clone_dtoh the original Task 2 wrapper used inside the 10 000- (Test 0.A) / 100 000-iteration (Test 0.B) loops, per feedback_gpu_cpu_roundtrip.md. Single-launch thompson_direction_test (1 thread, 1-pick output) replaced by thompson_direction_test_batched (one thread per seed, writes out_dir_idx[seed_idx] via a mapped pinned device pointer, finishes with __threadfence_system()). argmax_eq_test and the σ diagnostic kernels gain the same __threadfence_system() exit barrier. Test wrapper now allocates a local MappedI32Buffer (cuMemHostAlloc(DEVICEMAP|PORTABLE) + cuMemHostGetDevicePointer_v2, mirrors gpu_training_guard.rs::MappedBuffer) and read_volatiles the result — no memcpy_dtoh. The OnceLock-cached launch_thompson_direction / launch_argmax_eq single-pick wrappers are deleted (orphan after refactor; feedback_wire_everything_up.md). Tests 0.A and 0.B now run in ~0.15 s each on local RTX 3050 Ti (down from ~1.86 s and ~3.8 s respectively). Assertions unchanged. No production callers.

Persistent picked_action_history scatter (2026-04-26): gpu_backtest_evaluator.rs gains a window-major picked_action_history_buf (parallel layout to actions_history_buf) populated by an additional scatter_intent_chunk launch immediately after the existing intent-mag scatter — same kernel handle, different src/dst (chunked_actions_buf → picked_action_history_buf). The prior read_chunked_actions_direction_distribution reader was changed to read this persistent buffer instead of the chunk-local one, so the diagnostic now covers all 214 654 val bars instead of only the last ~64. Disambiguates whether val_dir_dist ≈ 80% Flat reflects upstream kernel collapse (most bars produce peaked Q for Hold/Flat → both val_dir_dist and val_picked_dir_dist flat) or downstream env_step gating (kernel diverse → only val_dir_dist flat). Reset to zeros by reset_evaluation_state; no new kernel source.

Plan 5 Task 5 Phase H Site 2 (2026-04-26): cublasLt EPILOGUE_DRELU_BGRAD fusion on the value-FC backward chain — collapses the standalone relu_mask

• bias-grad reduce launches at the value-FC site into the value-output dX cublasLtMatmul call. Site 1 (commit 326c13378) handled the 4 attention forward projections with EPILOGUE_BIAS; Site 2 saves 2 launches × 2 forward passes × N_steps on the L40S deploy. Forward-side switches online value-FC from EPILOGUE_RELU_BIAS to EPILOGUE_RELU_AUX_BIAS so cublasLt writes the dReLU bit-mask; new _value_fc_relu_mask_buf: CudaSlice<u8> in BatchedForward ((value_h_aux_ld / 8) * batch_size bytes, 128-bit AUX_LD aligned). Backward side adds gemm_value_dx_drelu_bgrad: Option<CachedBwdGemmDesc> cached descriptor + backward_fc_layer_drelu_bgrad method that gates dx by the mask AND reduces along N=batch into b_v1_grad in a single fused GEMM. backward_full signature gained value_fc_relu_mask_ptr: u64 threaded through both call sites in gpu_dqn_trainer.rs. Falls back to original backward_fc_layer + relu_mask + launch_dw_only when AUX path unavailable. relu_mask_kernel and bias_grad_reduce_f32_p1/_p2 retained — still consumed by apply_ensemble_diversity_backward, VSN backward, and branch FC layers. Smoke (3 folds × 5 epochs): F0 2.29 (-3% vs Phase G 2.36), F1 39.93 (-50.6%), F2 59.17 (-35.8%); F1 drop attributed to Boltzmann eval unification restructuring val-Sharpe distributions, not Site 2 itself. No determinism violation, no NaN, no panics. Wall-time delta deferred to L40S nsys.

val_picked_dir_dist HEALTH_DIAG diagnostic (2026-04-26): gpu_backtest_evaluator.rs gains a read_chunked_actions_direction_distribution() reader that decodes the per-direction histogram from chunked_actions_buf (the buffer experience_action_select writes to BEFORE env_step computes actual_dir). Pairs with the existing val_dir_dist reader (reads actions_history_buf AFTER env_step) — divergence between the two localises whether eval collapse is upstream (kernel produces degenerate Boltzmann picks → both lines flat) or downstream (kernel diverse, env_step / Kelly cap drains active picks to Flat → only val_dir_dist flat). Cluster runs ddrpr + txdz9 showed val_dir_dist ≈ 80% Flat / 20% Hold across 30+ epochs while Boltzmann math caps P(best) ≤ 47.5% and P(any) ≥ 13.5% for 4 actions with tau=q_range; one of those bounds is being violated and this diagnostic identifies which path. Sample is one chunk (~64 bars on 1-window val), noisier than the full-window post-physics histogram but sufficient to disambiguate the gating mechanism.

Kelly cap warmup_floor health-coupled (2026-04-26): trade_physics.cuh::kelly_position_cap previously computed warmup_floor = clamp(conviction, 0, 1). At cold start (maturity = 0) effective_kelly = warmup_floor, so a zero conviction collapsed the cap to zero — eval-mode could never open a position to graduate Kelly stats, the same bootstrap deadlock pattern the IQN trunk SAXPY exhibited. Confirmed by cluster run train-multi-seed-ddrpr epoch 0: new val_dir_dist [short=0.0000 hold=0.1953 long=0.0001 flat=0.8047] shows Boltzmann fired correctly (hold + flat ≈ 100% of bars) but every active-direction pick (Long/Short) was forced to actual_dir = Flat by target_position = 0. Fix: warmup_floor = max(conviction, health_floor) where health_floor = 0.5 + 0.5 × health is the same training- stability floor already used by safety_multiplier (composed in unified_env_step_core from ISV[LEARNING_HEALTH]). Both signals are adaptive and ISV-driven; the floor only matters during cold start (maturity → 1 after ≥10 trades drops the term out entirely). Threaded through both apply_kelly_cap and kelly_position_cap signatures.

Eval-mode action selection unified with Boltzmann softmax (2026-04-26): experience_kernels.cu — the four factored action heads (direction, magnitude, order, urgency) all dropped their else if (eval_mode) strict-argmax + ISV- tied tie-break blocks. The tie threshold was 0.01 × isv_signals[V_HALF_*_INDEX] (1% of the C51 atom support range, ~40 for direction); in practice the actual per-sample direction Q-spread post-IQN is ~1.5, so tie-break never fired and eval was pure strict-argmax over a peaked distribution → val argmax glued to one direction → 1-25 trades / 214k-bar window across cluster runs. Replaced with the same Boltzmann path training already used: tau = max(q_range, floor) where q_range is per-sample. Mathematically P(best) ≤ 47.5% for 4 actions with tau=q_range, so eval can never collapse to pure-greedy regardless of how peaked the Q-values become. Confirmed by new unit test cuda_pipeline::tests::test_eval_action_select_boltzmann_bounded — exercises the kernel directly with peaked synthetic Q-values [0, 0.5, 1.0, 0.5] and asserts the realised histogram matches Boltzmann theory (P(best) ≈ 0.366, ≤ 0.6, ≥ 0.25); the test runs in ~1.6 s after build, replacing 15-min smoke runs for kernel-level validation. New val_dir_dist line in HEALTH_DIAG logs per-direction eval distribution (short / hold / long / flat) — disambiguates the kernel-side dir_entropy collapse that bucketed Hold+Flat into one slot.

IQN trunk-gradient readiness gate removed (2026-04-26): gpu_dqn_trainer.rs:6461,6492 SAXPY scale was iqn_lambda × iqn_readiness × iqn_budget. iqn_readiness initialises to 0.0 and only ramps when iqn_loss_ema drops below iqn_loss_initial, but that improvement requires the trunk to receive IQN gradient — which the gate blocked. Bootstrap deadlock confirmed in cluster run train-multi-seed-vg2r9 epochs 0-13: trunk_iqn=0.0000 every epoch, q_gap_comp=0.00 every epoch, val trade_count locked at 22-34 per 858k-bar window. New scale: iqn_lambda × iqn_budget. The iqn_readiness field stays on self because the C51 loss kernel reuses iqn_readiness_dev_ptr as a CVaR-α pointer (separate semantic overload, tracked for follow-up — CVaR α=0 is degenerate).

P5T5 Phase G (2026-04-26): vol_normalizer numerical robustness + diagnostic logging. Phase D's aux label-scale EMA treated symptoms only; root cause was epoch_vol_normalizer in training_loop.rs:495 producing wildly different values across machines (local raw=0.541, cluster raw≈1e-7), inflating return features [0..3] up to 50,000× their intended unit-variance scale. Fix: Welford's online variance (numerically stable, single-pass), sanity bounds [1e-5, 1e-1], explicit per-epoch tracing::info! log. Out-of-band raw values trigger tracing::warn! + fall back to 5e-4 default (typical equity-index 1-min vol). multi_fold_convergence smoke (3 folds × 5 epochs, 682.85s): F0=2.3551, F1=80.8206, F2=92.1063 — all positive, F2 strongest result yet. The "targets[0] not raw_close on this dataset" warning surfaces a deeper data-pipeline question (what's actually at index 0 of the 6-tuple) for future investigation; Phase G makes training scale-correct regardless. cargo check clean at 11 warnings.

P5T5 Phase E (2026-04-26): per-fold reset for MetricBandsRegistry regression-detection streaks. Bug found during Phase D smoke — P5T2's registry never cleared consecutive_warn / consecutive_error HashMaps at fold boundaries, so walk-forward folds accumulated streaks across boundaries. F2 epoch in error band tripped termination at "6 consecutive" even though no fold individually crossed 2N=6. Fix: reset_streaks() method clears both HashMaps; called from DQNTrainer::reset_for_fold(). New unit test reset_streaks_clears_consecutive_counters verifies the per-fold isolation. 9/9 monitoring unit tests pass.

Legend:

• Wired — consumed by production training + val path.
• Partial — consumed on one side only (training OR val, or forward OR backward).
• Orphan — built but no production consumer in the DQN path.
• Ghost — consumer path is stubbed or only loads the kernel without launching it.
• OUT-of-DQN-scope — has production consumers outside the DQN path (supervised, PPO, etc.); correct-as-is per Part E.

DQN-path definition: trainers/dqn/ (training loop + fused CUDA ops) + cuda_pipeline/ (GPU kernels wired into those trainers) + validation/ harness when evaluating DQN policy.

Task 2 / Task 3 Seed Rows (preserved)

Module / kernel	Consumer path	Classification	Notes	Action
trainers/dqn/state_reset_registry.rs	training_loop.rs::reset_named_state via fold_reset_entries() + soft_reset_entries()	Wired	A.1 primary implementation (Plan 1 Task 3)	—
training_loop.rs::reset_named_state	consumer of StateResetRegistry::fold_reset_entries + soft_reset_entries	Wired	A.1 Task 3 dispatch + D.5 SoftReset dispatch	—
D.5 SoftReset dispatch (isv_grad_balance_targets, isv_grad_scale_limit)	training_loop.rs::reset_named_state — writes bootstrap values (1.0 for targets, 2.0 for limit). EMA-based convergence from bootstrap over ~50 epochs via grad_balance_isv_update kernel's adaptive-rate EMA (α ∈ [0.01, 0.30], implicit decay_bars).	Wired	Plan 2 Task 4 D.5	—

DQN Core Training Modules

Module / kernel	Consumer path	Classification	Notes	Action
trainers/dqn/mod.rs	train_baseline_rl.rs, ml_training_service, hyperopt/adapters/dqn.rs	Wired	Entry point for DQN trainer	—
trainers/dqn/config.rs (DQNHyperparameters, DQNAgentType)	trainer/constructor.rs, fused_training.rs, smoke tests, hyperopt adapter	Wired	Core config struct	—
trainers/dqn/data_loading.rs	re-exported via dqn::mod; consumed by train_baseline_rl.rs	Wired	DBN + fxcache loading	—
trainers/dqn/early_stopping.rs (EarlyStopping)	trainer/mod.rs, trainer/constructor.rs	Wired	Patience-based stopping	—
trainers/dqn/features.rs	trainer/constructor.rs via extract_features_from_bars	Wired	Feature extraction for DQN	—
trainers/dqn/financials.rs	trainer/metrics.rs via compute_epoch_financials	Wired	Sharpe / drawdown per epoch	—
trainers/dqn/fused_training.rs (FusedTrainingCtx)	trainer/training_loop.rs, trainer/mod.rs	Wired	Core DQN fused-CUDA context	—
trainers/dqn/lr_scheduler.rs (LRScheduler)	trainer/mod.rs::lr_scheduler, trainer/constructor.rs	Wired	Learning-rate schedule	—
trainers/dqn/monitoring.rs (MonitoringSummary)	trainer/training_loop.rs, trainer/metrics.rs	Wired	Per-epoch action / Q monitoring	—
trainers/dqn/risk.rs	trainer/constructor.rs via DrawdownMonitor, HybridPositionLimiter	Wired	Risk controllers in training	—
trainers/dqn/statistics.rs (FeatureStatistics, QValueStats)	trainer/mod.rs, trainer/metrics.rs	Wired	Feature norm + Q stats	—
trainers/dqn/adversarial_self_play.rs (AdversarialSaboteur)	trainer/constructor.rs, trainer/mod.rs	Wired	Adversarial perturbation during training	—
trainers/dqn/expert_demos.rs (ExpertDemoGenerator)	trainer/training_loop.rs	Wired	Expert demo ratio injection	—
trainers/dqn/trainer/mod.rs (DQNTrainer)	train_baseline_rl.rs, service layer	Wired	Top-level trainer struct	—
trainers/dqn/trainer/constructor.rs	internal to DQNTrainer	Wired	Builder / init	—
trainers/dqn/trainer/action.rs	internal to DQNTrainer	Wired	Action selection helpers	—
trainers/dqn/trainer/enrichment.rs	internal to DQNTrainer	Wired	State enrichment	—
trainers/dqn/trainer/metrics.rs	training_loop.rs	Wired	Epoch-boundary metric compilation + backtest eval	—
trainers/dqn/trainer/state.rs	trainer/mod.rs	Wired	Mutable training state	—
trainers/dqn/trainer/training_loop.rs	called from DQNTrainer::train()	Wired	Main epoch loop	—
trainers/dqn/trainer/monitoring.rs (MetricBandsRegistry, MetricBands, BandSettings, TerminationReason)	trainer/training_loop.rs (post-HEALTH_DIAG harvest+check), trainer/constructor.rs (auto-load config/metric-bands.toml), smoke_tests/regression_detection.rs	Wired	Plan 5 Task 2 A.4 — regression-detection hard-stop (2N consecutive error-band epochs → CommonError::RegressionDetected); cold-path (per-epoch); zero hot-path overhead	—
crates/common/src/error.rs (CommonError::RegressionDetected)	returned from training_loop::train_with_data_full_loop_slices; consumed by services/trading_service/src/error.rs::From<TradingServiceError>	Wired	Plan 5 Task 2 — terminal training error variant; flat band fields (no upward dep on ml crate)	—
config/metric-bands.toml	MetricBandsRegistry::load_from_toml (auto-loaded by DQNTrainer::new_internal)	Wired	Plan 5 Task 2 — per-metric warn/error bands (mean ± 5σ / mean ± 10σ); populated by scripts/populate-metric-bands-from-runs.py from Plan 5 Task 1B aggregate JSONs	—
scripts/populate-metric-bands-from-runs.py	reads aggregate JSONs from scripts/aggregate-multi-seed-metrics.py (Plan 5 Task 1B.2); writes config/metric-bands.toml [bands] section	Wired	Plan 5 Task 2 helper — N>=3 sample path; N=1 multiplicative fallback	—
scripts/argo-train.sh --profile + infra/k8s/argo/train-multi-seed-template.yaml (profile parameter, conditional nsys profile wrapper, mc cp upload to foxhunt-training-artifacts/profiles/<sha>/) + infra/docker/Dockerfile.foxhunt-training-runtime (nsight-systems-cli apt install) + infra/k8s/minio/minio.yaml (new foxhunt-training-artifacts bucket)	invoked manually via ./scripts/argo-train.sh dqn --profile [...]; consumed by scripts/compare-nsys-profiles.py BASELINE CURRENT --epochs N (V0 metric: total cuda_gpu_kern_sum / epochs; V1 NVTX-range path deferred to Plan 5 Task 5)	Wired	Plan 5 Task 3 A.4.1 — nsys profile harness with regression-comparison script. --profile forces multi-seed render path so dry-run never needs cluster contact (test_nsys_harness.sh); MinIO creds optional on train-single (warn-skips upload if absent). Baseline capture deferred to Plan 5 Task 5 (T5 runs the harness on L40S as part of the multi-seed acceptance pass — avoids burning local GPU time in T3)	—
scripts/validation/check_tier1.py + check_tier2.py + check_tier3.py + check_all_tiers.py (+ tests/test_tier_checks.sh + tests/fixtures/{good,bad}_tier1.json)	reads aggregate JSONs from scripts/aggregate-multi-seed-metrics.py (Plan 5 Task 1B.2); invoked manually as python3 scripts/validation/check_all_tiers.py <metrics.json> [--warmup-end N] and from Plan 5 Task 5's acceptance pass	Wired	Plan 5 Task 4 — per-tier exit checks for the spec §2 tiered acceptance criteria. Tier 1 (convergence: std/mean ≤ 0.15 on val_sharpe/avg_q_value/train_loss over stable epochs + no avg_q_value > 500 fold-1 explosion + Q-saturation/hot-path-DtoH placeholders); Tier 2 (behavioural: val_trades_per_bar ≥ 0.005, val_active_frac > 0.2, dir entropy > 0.8·log4); Tier 3 (profitability: val_sharpe_annualised > 1.0 with per-bar fallback, val_win_rate ≥ 0.52 gated on >500 trades, val_profit_factor mean ≥ 1.1 AND cross-seed std < 0.3). Stdlib only (no numpy/scipy). Defensive missing-metric handling — each missing aggregate key fails the relevant check with an explanatory message rather than silently passing. Tier-2/tier-3 wiring landed in Plan 5 Task 5 Phase A (val [...] HEALTH_DIAG block + aggregator single-_ joiner) — see next row.	—
trainer/metrics.rs::compute_validation_loss (HEALTH_DIAG val [...] block) + scripts/aggregate-multi-seed-metrics.py::parse_blocks (single-_ joiner) + trainer/mod.rs (last_val_metrics: Option<[f32; 14]>)	Plan 5 Task 4 tier-2/tier-3 check scripts read these as top-level val_* aggregate keys. Block emit pipeline: evaluate_dqn_graphed → WindowMetrics{sharpe, sortino, win_rate, max_drawdown, total_trades, calmar, omega_ratio, total_pnl, var_95, cvar_95, buy/sell/hold_count} (CUDA compute_backtest_metrics 14-float reduction, no new GPU work) → CPU derivations (window_bars = buy+sell+hold; trades_per_bar = total_trades / window_bars; active_frac = (buy+sell) / window_bars; dir_entropy = -Σ p ln p over the 3-bucket {short, hold-or-flat, long} distribution; sharpe_annualised = m.sharpe alias since the kernel already multiplies by sqrt(bars_per_day · 252); profit_factor = m.omega_ratio alias since the kernel's omega is gain_sum / loss_sum with threshold 0, equivalent to per-step PF) → tracing::info!("HEALTH_DIAG[{}]: val [sharpe=… sortino=… win_rate=… max_drawdown=… trade_count=… calmar=… omega_ratio=… total_pnl=… var_95=… cvar_95=… trades_per_bar=… active_frac=… dir_entropy=… sharpe_annualised=… profit_factor=… window_bars=…]", current_epoch, …). The aggregator joins <block>_<key> (single underscore — was __ before T5 Phase A; tier scripts and synthetic test fixtures already used the bare val_* convention) so each emitted key surfaces as a top-level val_* aggregate. Deferred for follow-up (cannot be emitted from existing kernel data): (a) per-direction distribution val_dir_dist_{short,hold,long,flat} — kernel collapses Hold+Flat into hold_count (intentional for the position-sign trade-cycle definition), so dir_entropy here ranges over 3 buckets with max log 3 ≈ 1.099 rather than the spec's 4-bucket log 4 ≈ 1.386 ceiling; tier2 check_dir_entropy compares against 0.8 · log 4 ≈ 1.109 which is unreachable from the 3-bucket distribution. Resolution options: extend WindowMetrics with separate Hold/Flat counts (kernel touch) or lower tier2's threshold to the 3-bucket equivalent 0.8 · log 3 ≈ 0.879. (b) trade-level profit_factor (sum-of-winning-trade-P&L / sum-of-losing-trade-P&L) — currently aliased to per-step omega_ratio; matches the standard PF definition under threshold 0 but a trade-level variant would require boundary-aware per-trade P&L accumulation in the kernel. Cold-path (per validation epoch); zero hot-path overhead; one new tracing line per epoch.	—
cuda_pipeline/aux_heads_loss_ema_kernel.cu::aux_label_scale_ema_update + cuda_pipeline/aux_heads_kernel.cu::{aux_next_bar_loss_reduce, aux_next_bar_backward} (kernel signatures grow +isv_dev_ptr+isv_label_scale_index args, divide label by max(isv[117], 1e-6) before residual) + cuda_pipeline/gpu_aux_heads.rs (orchestrator: launch_label_scale_ema method + next_bar_loss_reduce / backward_next_bar arg refactor) + cuda_pipeline/gpu_dqn_trainer.rs (AUX_LABEL_SCALE_EMA_INDEX=117, ISV_TOTAL_DIM 117→118, fingerprint seed +AUX_LABEL_SCALE_EMA=117/ISV_TOTAL_DIM=118, aux_heads_forward Step 2b producer launch, Step 5 + aux_heads_backward pass isv_signals_dev_ptr + slot index) + state_reset_registry.rs (isv_aux_label_scale_ema FoldReset → 1.0) + trainer/training_loop.rs::reset_named_state dispatch arm + HEALTH_DIAG aux line +label_scale={:.3e} field	producer per training step alongside the rest of the Plan 4 ISV producer family; consumer per-step in the captured forward + backward graphs (graph-capture-stable: ISV device pointer + slot index pair both stable, scalar value updates each step via the producer)	Wired	Plan 4 Task 6 / Plan 5 Task 5 follow-up — defends the aux next-bar regression head against underlying-data scale (label = next_states[:, 0], which carries log returns ~1e-3 in local fxcache but raw price ~5000 in the L40S fxcache). First L40S deploy attempt (workflow train-multi-seed-7j8zc, terminated) produced aux next_bar_mse = 2.587e7 and grad_norm=126,769 because the unnormalised label dominated the residual; the trunk then learned garbage off the corrupted aux-gradient SAXPY into bw_d_h_s2. Fix is principled per feedback_adaptive_not_tuned.md + feedback_isv_for_adaptive_bounds.md: ISV-driven label-scale EMA + normalize-before-MSE, NOT a tuned constant clamp or a label-source change (label still = next_states[:, 0] per spec §4.E.6). Cold-start 1.0 (multiplicative identity, NOT 0.0 which would divide-by-zero on first batch); FoldReset → 1.0; per-step EMA α=0.05 matches the rest of the Plan 4 producer family. Layout fingerprint shifts from 0x26f7b1deb94cb226 to 0x829bc87b42f2feee — checkpoint-incompatible (no migrator per spec §4.A.2). NEW BASELINE: aux next_bar_mse value is now O(1.0) NOT O(1e-4) — the pre-fix tiny value was a numerical artefact of the unnormalised label being so much smaller than pred ≈ 0 at init that the residual was just (0 - 1e-3)^2 ≈ 1e-6 → mean 1e-4. The new value is the actual per-residual magnitude after normalisation.	—
cuda_pipeline/mod.rs::{global_seed, mix_seed} + examples/train_baseline_rl.rs (--seed N CLI arg, default 42; sets FOXHUNT_SEED env at startup) + infra/k8s/argo/train-multi-seed-template.yaml (drops fold parameter; binary invoked with --seed "$SEED" --max-folds {{workflow.parameters.folds}}) + scripts/argo-train.sh (matrix generator: N tasks instead of N*K) + scripts/tests/test_multi_seed_harness.sh (asserts N tasks + --max-folds placeholder + no --fold)	call sites: trainer/action.rs (GpuActionSelector seed + epsilon-greedy StdRng), cuda_pipeline/gpu_iqn_head.rs (Xavier init RNG), cuda_pipeline/gpu_iql_trainer.rs (Xavier init RNG), cuda_pipeline/gpu_her.rs (random-donor RNG), cuda_pipeline/gpu_ppo_collector.rs (rng_seeds for PPO experience), trainer/training_loop.rs::set_regime_dropout_seed (per-epoch dropout seed)	Wired	Plan 5 Task 5 Phase B — architectural pivot from N×K (seed,fold) Argo fanout to N seed-only fanout. The first L40S deploy attempt (workflow train-multi-seed-z2llf, terminated) failed at startup with error: unexpected argument '--fold' found on every job: train_baseline_rl is a multi-fold walk-forward executor that accepts --max-folds K, NOT --fold N. Pivot reduces fanout from N×K=30 → N=5 (matches L40S pool capacity better) and trades K× longer per-job runtime for a simpler binary contract. The seed-variation mechanism: train_baseline_rl --seed N sets FOXHUNT_SEED=N env at startup BEFORE any CUDA module spins up; every previously-fixed RNG seed across the call sites listed above now mixes the global seed via mix_seed(<historic_constant>) (SplitMix64 avalanche so adjacent N produce uncorrelated module seeds). Default --seed 42 documents the historic implicit value. Verified end-to-end on RTX 3050 Ti: with seed=42, F0 best-Sharpe=-9.7831 + best_val_metric=1.957244 (matches the prompt's expected baseline); with seed=999, F0 best-Sharpe=92.9341 + best_val_metric=2.161012 — different trajectories, proving the seed propagates through the RNG init paths and is not just accepted-and-ignored. Backward-compat: existing single-job argo-train.sh callers (no --multi-seed) route to train-template.yaml unchanged.	—
trainers/dqn/adaptive_monitor.rs	Read-only observer trait + harness (FireRateStats, DiagSnapshot, IsvBus<'a>); consumers added in Plan 1 Tasks 9-17 (atoms/gamma/kelly_cap/tau/epsilon/grad_balancer monitors)	Wired (consumers added in same plan)	C.6 GPU-drives-CPU-reads	—
trainers/dqn/monitors/grad_balancer_monitor.rs	Read-only observer for grad_balance_isv_update kernel output (ISV slots 31..35); consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 1 Task 17	—
trainers/dqn/monitors/tau_monitor.rs	Read-only observer for tau_update kernel output (ISV slot 42); consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 1 Task 13	—
cuda_pipeline/tau_update_kernel.cu	GpuDqnTrainer::launch_tau_update → FusedTrainingCtx::launch_tau_update → training_loop.rs epoch boundary	Wired	Plan 1 Task 13 — cold-path (per-epoch)	—
trainers/dqn/monitors/epsilon_monitor.rs	Read-only observer for epsilon_update kernel output (ISV slot 41); consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 1 Task 14	—
cuda_pipeline/epsilon_update_kernel.cu	GpuDqnTrainer::launch_epsilon_update → FusedTrainingCtx::launch_epsilon_update → training_loop.rs epoch boundary	Wired	Plan 1 Task 14 — cold-path (per-epoch)	—
trainers/dqn/monitors/gamma_monitor.rs	Read-only per-branch observer; mean of ISV[43..47) = [GAMMA_DIR,MAG,ORD,URG]; consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 2 Task 3 D.2	—
cuda_pipeline/per_branch_gamma_update_kernel.cu	GpuDqnTrainer::launch_per_branch_gamma_update → FusedTrainingCtx::launch_per_branch_gamma_update → training_loop.rs epoch boundary	Wired	Plan 2 Task 3 D.2 — cold-path (per-epoch); replaces scalar gamma_update_kernel.cu	—
trainers/dqn/monitors/kelly_cap_monitor.rs	Read-only observer for kelly_cap_update kernel output (ISV slot 47); consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 1 Task 11	—
cuda_pipeline/kelly_cap_update_kernel.cu	GpuDqnTrainer::launch_kelly_cap_update → FusedTrainingCtx::launch_kelly_cap_update → training_loop.rs epoch boundary; takes portfolio_states dev ptr + n_envs	Wired	Plan 1 Task 11 — cold-path (per-epoch)	—
trainers/dqn/monitors/atoms_monitor.rs	Read-only observer for atoms_update kernel inputs (ISV v-range slots 23..31); consumers: HEALTH_DIAG + controller_activity smoke	Wired	Plan 1 Task 9 (test initial-value fixup post-land; local-smoke StateResetRegistry dispatch arms + hp_f32 helper added)	—
cuda_pipeline/atoms_update_kernel.cu	GpuDqnTrainer::launch_atoms_update → FusedTrainingCtx::launch_atoms_update → training_loop.rs epoch boundary; replaces per-branch recompute_atom_positions CPU loop	Wired	Plan 1 Task 9 — cold-path (per-epoch + SGD step)	—
cuda_pipeline/q_quantile_kernel.cu	GpuDqnTrainer::launch_q_quantile_reduce → FusedTrainingCtx::launch_q_quantile_reduce → training_loop.rs epoch boundary; reads q_out_buf [N, total_actions], writes ISV[50..58) P5/P95 EMAs per branch	Wired	Plan 2 Task 1 C.1 — cold-path (per-epoch); consumer: update_eval_v_range reads Q_P05/Q_P95 to set quantile-based half-width	—
trainers/dqn/monitors/reward_component_monitor.rs	Read-only observer for reward_component_ema kernel output (ISV slots 63..69); consumers: HEALTH_DIAG reward_split line + controller_activity smoke fire-rate tracking	Wired	Plan 3 Task 1 C.2	—
cuda_pipeline/reward_component_ema_kernel.cu	GpuExperienceCollector::launch_reward_component_ema_inplace → training_loop.rs (called after reward_contrib_fractions, before HEALTH_DIAG); reads reward_components_per_sample [N*L, 6], writes ISV[63..69) adaptive EMA (α=0.05)	Wired	Plan 3 Task 1 C.2 — hot-path (per-step); single-block 6-thread kernel	—
B.1 Flat opp-cost ISV scaling (experience_kernels.cu Flat branch)	experience_env_step kernel (Flat position=0, not segment_complete branch); writes rc[4] to reward_components_per_sample; consumed by reward_component_ema → ISV[67]	Wired	Plan 3 Task 2 B.1 — opp-cost multiplied by isv_signals_ptr[21] (Q_DIR_ABS_REF_INDEX, EMA of max(	Q_mean
cuda_pipeline/trade_rate_ema_kernel.cu	GpuExperienceCollector::launch_trade_attempt_rate_ema_inplace → training_loop.rs (called alongside launch_reward_component_ema_inplace each epoch); reads flat_to_pos_per_sample [N*L], reduces count on-GPU, writes ISV[TRADE_ATTEMPT_RATE_EMA_INDEX=71] adaptive EMA (α=α_base × (1+0.5×|sharpe|), α_base=0.05)	Wired	Plan 3 Task 3 B.2 — cold-path (per-epoch); single-block 1-thread reduction + EMA. No atomicAdd.	—
B.2 Flat→Positioned novelty bonus (experience_kernels.cu trade-lifecycle block)	experience_env_step kernel (entering_trade path, AFTER opp-cost branch, BEFORE drawdown penalty); writes rc[5] and adds shaping_scale × bonus to reward; consumed by reward_component_ema → ISV[68] and drives PopArt denominator via total_reward_per_sample	Wired	Plan 3 Task 3 B.2 — novelty = max(0, 1 − ISV[71]/max(1e-4, ISV[72])); bonus = conviction_core × vol_proxy × novelty; TRADE_TARGET_RATE frozen at epoch 5 from measured attempt-EMA (min 0.001). No tuned multiplier.	—
C.4 Temporal timing bonus on trade exit (experience_kernels.cu segment_complete block)	experience_env_step kernel (inside segment_complete && segment_hold_time > 0 block, AFTER Layer 2–9 credits land in reward); accumulates += into rc[5] and adds timing_bonus to reward; consumed by reward_component_ema → ISV[68] (shares the bonus slot with B.2 — different (i,t) slots, so += is idempotent)	Wired	Plan 3 Task 5 C.4 — bars_early = max(0, segment_hold_time − PS_PEAK_PNL_BAR); timing_bonus = shaping_scale × (bars_early / segment_hold_time) × |final_pnl| × conviction_core. Peak bar tracked in new portfolio-state slot PS_PEAK_PNL_BAR=38 (PS_STRIDE 38→39), snapshot alongside every MAX_PNL update, reset at every MAX_PNL reset site (entry/reverse/fold-reset). No tuned multiplier.	—
D.4a Persistence credit on profitable drawdown recovery (experience_kernels.cu segment_complete block)	experience_env_step kernel (inside segment_complete && segment_hold_time > 0 block, IMMEDIATELY AFTER C.4 timing bonus, gated on reward > 0 AND drawdown_depth > 1e-6); accumulates += into rc[5] and adds persist_bonus to reward; consumed by reward_component_ema → ISV[68] (shares the bonus slot with B.2 entry + C.4 exit — different (i,t) slots per trade, so += is idempotent across all three)	Wired	Plan 3 Task 6a D.4a — drawdown_depth = max(0, −PS_INTRA_TRADE_MIN_PNL); persist_bonus = shaping_scale × conviction_core × drawdown_depth × tanh(reward / max(1e-4, drawdown_depth)). MIN_PNL tracked in new portfolio-state slot PS_INTRA_TRADE_MIN_PNL=39 (PS_STRIDE 39→40), updated per-bar in the same block as MAX_PNL (fminf against pnl_pct), reset at every MAX_PNL reset site (entry/reverse/fold hard-reset/trade-complete soft-reset). Self-scaling tanh: saturates +1 when reward ≫ drawdown, ≈0 when reward ≈ drawdown. No tuned multiplier.	—
D.4b Regime-shift penalty for trades held past a regime flip (experience_kernels.cu segment_complete block)	experience_env_step kernel: detector inside the active-trade block (after MIN/MAX_PNL update) writes PS_REGIME_SHIFT_BAR on the FIRST bar where |isv[11] − PS_PLAN_ENTRY_REGIME| > clamp(0.25 × |clamp(sharpe, ±2)|, 0.05, 0.5); consumer inside segment_complete block IMMEDIATELY AFTER D.4a accumulates −= into rc[5] and subtracts penalty from reward; consumed by reward_component_ema → ISV[68] (cancellation within rc[5] vs B.2/C.4/D.4a is intentional — that's what ISV[68] REWARD_BONUS_EMA tracks). Consumer also resets PS_REGIME_SHIFT_BAR = 0 after use.	Wired	Plan 3 Task 6b D.4b — bars_late = max(0, saved_hold_time − PS_REGIME_SHIFT_BAR); penalty = shaping_scale × conviction_core × (bars_late / saved_hold_time) × ISV[Q_DIR_ABS_REF_INDEX=21] × |reward|. Shift-bar tracked in new portfolio-state slot PS_REGIME_SHIFT_BAR=40 (PS_STRIDE 40→41), reset at every MAX_PNL reset site (entry/reverse/fold hard-reset/trade-complete soft-reset) PLUS post-consumer (defensive). Self-scaling Q_DIR_ABS_REF coefficient (same B.1 pattern). Adaptive threshold tightens when |sharpe| is high. First-shift-only. No tuned multiplier.	—
D.4c Conviction consistency bonus on Flat→Positioned entry (experience_kernels.cu Flat branch + entering_trade block)	experience_env_step kernel: producer inside the Flat branch (!segment_complete && position≈0, before opp-cost) updates two PS slots PS_PRE_ENTRY_CONVICTION_EMA (mean) and PS_PRE_ENTRY_CONVICTION_VAR_EMA (variance) via Welford-style EMA at α=0.05 on conviction_core. Consumer inside entering_trade block IMMEDIATELY AFTER B.2 novelty bonus computes ratio = stddev/mean; when ratio < 0.2 accumulates += into rc[5] and adds bonus to reward; consumed by reward_component_ema → ISV[68] (shares the bonus slot with B.2/C.4/D.4a/D.4b at different (i,t) slots; += is idempotent). Both EMA slots reset at every trade-lifecycle reset site (entry/reverse/fold hard-reset/trade-complete soft-reset) PLUS post-consumer (defensive).	Wired	Plan 3 Task 6c D.4c — bonus = shaping_scale × vol_proxy × stability × conviction_core where vol_proxy ∈ [0.0001, 0.01], stability = clamp(0, 1, 1 − ratio/0.2), conviction_core ∈ [0,1]. Mirrors B.2 novelty-bonus structure exactly — ONE technically unbounded multiplicand (vol_proxy, bounded ≤ 0.01), all others bounded — per pearl_one_unbounded_signal_per_reward.md. Max bonus ≈ 0.01, same order as B.2. New portfolio-state slots PS_PRE_ENTRY_CONVICTION_EMA=41 and PS_PRE_ENTRY_CONVICTION_VAR_EMA=42; PS_STRIDE 41→43; PORTFOLIO_STRIDE 41→43 in all consumers (experience_kernels.cu, trade_stats_kernel.cu, gpu_experience_collector.rs, gpu_dqn_trainer.rs::launch_kelly_cap_update, ml-core::state_layout.rs). The 0.2 stability threshold is the single tuned constant retained for this first cut (commented; reviewer may demote to ISV-driven follow-up). α=0.05 matches Plan 3 Task 1 reward-EMA convention.	—
cuda_pipeline/plan_threshold_update_kernel.cu	GpuExperienceCollector::launch_plan_threshold_update_inplace → training_loop.rs (called alongside launch_reward_component_ema_inplace and launch_trade_attempt_rate_ema_inplace each epoch); reads readiness_per_sample [N*L] written by experience_env_step, reduces mean on-GPU, writes ISV[READINESS_EMA_INDEX=75] adaptive EMA (α=α_base × (1+0.5×|sharpe|), α_base=0.05), then derives ISV[PLAN_THRESHOLD_INDEX=49] = max(0.1, 0.5 × ema). Producer-only upgrade — consumer kernels (experience_kernels.cu 4 sites + backtest_plan_kernel.cu 1 site) unchanged.	Wired	Plan 3 Task 4 B.4 — cold-path (per-epoch); single-block 1-thread reduction + EMA. No atomicAdd.	—
trainers/dqn/monitors/plan_threshold_monitor.rs	Read-only observer for plan_threshold_update kernel output (ISV slots 49 + 75); consumers: HEALTH_DIAG plan_threshold.eff / plan_threshold.readiness_ema + controller_activity smoke fire-rate tracking	Wired	Plan 3 Task 4 B.4	—
cuda_pipeline/state_kl_divergence_kernel.cu	GpuExperienceCollector::launch_state_kl_inplace → training_loop.rs (called once per validation epoch, AFTER compute_validation_loss populates chunked_states_buf, BEFORE next training epoch's reward kernels read ISV[79]); reads train-state sample (states_out) and val-state batch (chunked_states_buf from gpu_evaluator.val_state_sample()), per-OFI-dim Gaussian moment-match KL summed across [SL_OFI_START, SL_OFI_START+SL_OFI_DIM); writes ISV[STATE_KL_TRAIN_VAL_EMA=78] (adaptive EMA, α_base=0.05) + ISV[STATE_KL_AMPLIFICATION=79] (kernel-internal trailing-self ratio → 1 + clamp(ratio−1, 0, 1) ∈ [1,2]). Single-block, one thread per OFI dim. No atomicAdd, no DtoH.	Wired	Plan 3 Task 7 C.3 — cold-path (per validation epoch). All α/threshold coefficients ISV-derived; no tuned constants per feedback_isv_for_adaptive_bounds.md.	—
C.3 B.1/B.2 KL-amp consumers (experience_kernels.cu Flat opp_cost + entering_trade B.2 bonus)	experience_env_step kernel: B.1 site multiplies reward = -shaping_scale × holding_cost_rate × conviction_core × vol_proxy_flat × q_abs_ref × kl_amp_b1; B.2 site multiplies bonus_scaled = shaping_scale × bonus × kl_amp_b2. Both consumers use fmaxf(1.0, isv_signals_ptr[ISV_STATE_KL_AMP_IDX]) to no-op against cold-start. ISV_STATE_KL_AMP_IDX=79 macro defined in state_layout.cuh.	Wired	Plan 3 Task 7 C.3 — bounded amp ∈ [1, 2] stacks safely with B.1's existing q_abs_ref unbounded multiplicand per pearl_one_unbounded_signal_per_reward.md.	—
trainers/dqn/monitors/state_kl_monitor.rs	Read-only observer for state_kl_moment_match kernel output (ISV slots 78 + 79); consumers: HEALTH_DIAG state_kl.train_val_ema / state_kl.amp + controller_activity smoke fire-rate tracking	Wired	Plan 3 Task 7 C.3	—
cuda_pipeline/scripted_policy_kernel.cu	GpuExperienceCollector::launch_timestep_loop (per-timestep dispatch when seed_phase_active_cache=true); 4 policies (UNIFORM 40% / MOMENTUM 20% / MEAN_REV 20% / VWAP_DEV 20%) deterministically mixed by i % 5. Per-thread one sample. Reads batch_states[i, MARKET_START] for current close + portfolio_states[i, PS_PREV_CLOSE] for prev close; writes batch_actions[i] (factored dir*27 + mag*9 + ord*3 + urg) and conviction_buf[i] ∈ [0, 1] — same outputs the network's experience_action_select would have written. Direction-flip thresholds (±0.0001f momentum/reversal, ±5bp VWAP band) are noise-floor cutoffs, not scaling coefficients.	Wired	Plan 3 Task 8 B.3 — GPU-only seeded warm-start. CPU only orchestrates per-epoch dispatch (cold-path read of ISV[SEED_STEPS_DONE/TARGET]); the action computation itself is 100% GPU. No CPU physics mirror — existing experience_env_step runs unchanged.	—
C51 bias C.1 ISV-adaptive direction-Boltzmann tau floor (experience_kernels.cu experience_action_select direction branch)	experience_action_select Branch 0 direction softmax: tau_d = max(q_range, max(ISV[Q_DIR_ABS_REF_INDEX=21], 0.01)). Replaces static 0.01 floor (tuned constant, dangerous when Q-magnitudes grow during training: spread that's small in absolute terms becomes deterministic argmax even though it represents no real edge relative to scale). Now scales with the network's actual Q magnitude EMA, preserving relative spread for sampling and preventing the C51 expected-Q Hold/Flat attractor that the val-Flat-collapse Kelly fix exposed (val_dir_dist S=.23/H=.17/L=.42/F=.18 epoch 1 → S=.14/H=.35/L=.15/F=.37 epoch 2 = 35→72% Hold+Flat shift in one epoch). Cold-start fallback retains 0.01 minimum so kernel doesn't divide by zero pre-first-update. Same ISV[21] reference used by conviction below — coherent adaptive mechanism. NOTE: empirical verification (train-bscl2) showed this tau-floor change had no measurable effect on the post-training Hold/Flat collapse — once Q-values reflect tx_cost-driven aversion, sampling stays biased regardless of tau. Retained as cold-start protection; bias breakout handled by the eps_dir adaptive floor below.	Wired	val-collapse follow-up — C51 expected-Q bias fix (cold-start protection only). ISV-driven, no tuned constants per feedback_isv_for_adaptive_bounds.md and feedback_adaptive_not_tuned.md.	—
C51 bias C.2 trade-attempt-rate-driven adaptive eps_dir floor (experience_kernels.cu experience_action_select eps floor block)	experience_action_select Branch 0 direction floor: eps_dir = max(eps_dir, 0.5 × passive_pressure) where passive_pressure = clamp(0, 1, 1 − ISV[71]/max(ISV[72], 1e-4)). Reads TRADE_ATTEMPT_RATE_EMA_INDEX=71 (current Flat→Positioned rate) and TRADE_TARGET_RATE_INDEX=72 (target rate frozen at epoch 5 from measured EMA). When the policy collapses to passive (zero trade attempts), eps_dir floor rises to 0.5 — half random direction sampling — forcing enough Long/Short experiences into the replay buffer for Q-values to discover real edge. At target rate or above, passive_pressure=0, eps_dir at baseline 0.02 floor. Direction branch only: magnitude/order/urgency don't have the Flat-attractor problem (Kelly cap shapes magnitude realisation; order/urgency operate on top of already-decided directional picks). The 0.5 ceiling is a structural blend point (half random / half policy), not a tuned magnitude — maximum exploration that still preserves directional Q-signal propagation through the replay buffer. Active in training only (eval_mode skips eps logic entirely so val remains deterministic).	Wired	val-collapse follow-up — C51 expected-Q bias breakout. ISV-driven feedback control via existing trade-rate EMA (B.2 producer); no new ISV slots.	—
cuda_pipeline/seed_step_counter_update_kernel.cu	GpuExperienceCollector::launch_seed_step_counter_update_inplace → training_loop.rs (called alongside the other Plan 3 EMA producers each epoch); single-block single-thread cold-path. Increments ISV[SEED_STEPS_DONE_INDEX=83] by n_samples, capped at ISV[SEED_STEPS_TARGET_INDEX=82], and EMAs the derived max(0, 1 - DONE/TARGET) into ISV[SEED_FRAC_EMA_INDEX=84]. Adaptive α=α_base × (1+0.5×|clamp(sharpe,−2,2)|), α_base=0.05.	Wired	Plan 3 Task 8 B.3 — cold-path (per-epoch). No atomicAdd. SEED_FRAC_EMA cold-start 1.0 ensures Task 9's CQL ramp sees target=0 until the seed phase actually decays.	—
cuda_pipeline/cql_alpha_seed_update_kernel.cu	GpuExperienceCollector::launch_cql_alpha_seed_update_inplace → training_loop.rs (called alongside launch_seed_step_counter_update_inplace each epoch); single-block single-thread cold-path. EMAs ISV[CQL_ALPHA_INDEX=48] toward cql_alpha_final × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX=84]) where cql_alpha_final = config.cql_alpha. Adaptive α matches Task 8 convention.	Wired	Plan 3 Task 9 C.5 — producer upgrade for ISV[CQL_ALPHA_INDEX=48]: SchemaContract → FoldReset. CQL gradient kernel consumer (already reading slot 48 per Plan 1 Task 12) unchanged.	—
trainers/dqn/monitors/seed_monitor.rs	Read-only observer for seed_step_counter_update kernel output (ISV slots 82 / 83 / 84); consumers: HEALTH_DIAG seed.steps_target / seed.steps_done / seed.frac_ema + controller_activity smoke fire-rate tracking	Wired	Plan 3 Task 8 B.3	—
trainers/dqn/monitors/cql_alpha_monitor.rs	Read-only observer for cql_alpha_seed_update kernel output (ISV slot 48 + dependency 84); consumers: HEALTH_DIAG cql_alpha.eff / cql_alpha.seed_frac + controller_activity smoke fire-rate tracking	Wired	Plan 3 Task 9 C.5	—
cuda_pipeline/attention_focus_ema_kernel.cu	GpuExperienceCollector::launch_attention_focus_ema_inplace → training_loop.rs (called once per epoch at the HEALTH_DIAG site, BEFORE the diag emit so slots are fresh); single-block single-thread cold-path. Takes 3 host scalars by value: vsn_mag/vsn_dir from FusedTrainingCtx::per_branch_vsn_mean() and mamba2_retention from FusedTrainingCtx::mamba2_retention_mean(batch) (host-side dtoh of mamba2_h_enriched, mirror of per_branch_vsn_mean). EMA-updates ISV[VSN_MAG_EMA_INDEX=87] / ISV[VSN_DIR_EMA_INDEX=88] / ISV[MAMBA2_RETENTION_EMA_INDEX=89] using adaptive α=α_base × (1+0.5×|clamp(sharpe, −2, 2)|), α_base=0.05.	Wired	Plan 4 Task 5 Mode A E.5 — diagnostic only; no consumer kernel reads slots [87..90) in Mode A. Mode B (post-E.1 full VSN) will tail-append 4 more per-feature-group slots.	—
trainers/dqn/monitors/attention_monitor.rs	Read-only observer for attention_focus_ema_update kernel output (ISV slots 87 / 88 / 89); consumers: HEALTH_DIAG attention.vsn_mag / attention.vsn_dir / attention.mamba2_retain + (when surfaced) controller_activity smoke fire-rate tracking. Anchor read on slot 88 (VSN_DIR_EMA).	Wired	Plan 4 Task 5 Mode A E.5	—

CUDA Pipeline — Rust Wrappers

Module / kernel	Consumer path	Classification	Notes	Action
cuda_pipeline/mod.rs (DqnGpuData, upload_slices)	trainer/constructor.rs, fused_training.rs	Wired	Data upload to GPU	—
cuda_pipeline/gpu_dqn_trainer.rs (GpuDqnTrainer)	fused_training.rs, trainer/mod.rs	Wired	Core GPU-side DQN weight / forward / grad ops	—
cuda_pipeline/fused_training.rs wrapper (FusedTrainingCtx)	trainer/training_loop.rs	Wired	Wires all GPU ops into epoch loop	—
cuda_pipeline/batched_forward.rs	gpu_dqn_trainer.rs — forward graph node	Wired	Batched SGEMM forward pass	—
cuda_pipeline/state_encoder.rs (StateEncoder)	Borrows CublasGemmSet and routes through BatchedForward::encoder_forward_only; coexists with the monolithic forward_online_raw path used by gpu_dqn_trainer.rs	Wired (additive — coexists with forward_online_raw monolithic path)	Plan 4 Task 4 E.4 — explicit trunk-encoder boundary; attachment point for Plan 4 Task 1 (Full VSN, pre-encoder) and follow-up per-branch experiments	—
cuda_pipeline/value_decoder.rs (ValueDecoder, Branch, BranchDecoderOutput)	Borrows CublasGemmSet and routes through BatchedForward::decoder_forward_only; one instance per branch (Dir/Mag/Ord/Urg). Coexists with the multi-stream branch fork/join in forward_online_raw	Wired (additive — coexists with forward_online_raw monolithic path)	Plan 4 Task 4 E.4 — per-branch value-decoder boundary; attachment point for Plan 4 Task 3 (multi-Q IQN) and Task 6 (auxiliary heads). Surfaces Plan 2 D.3 V_short / V_long pointers via BranchDecoderOutput	—
cuda_pipeline/batched_backward.rs	gpu_dqn_trainer.rs — backward graph node	Wired	Batched SGEMM backward pass (KAN gate included)	—
cuda_pipeline/shared_cublas_handle.rs	fused_training.rs, gpu_iqn_head.rs, gpu_iql_trainer.rs, gpu_attention.rs, gpu_curiosity_trainer.rs (10 consumers)	Wired	Shared cuBLAS/cuBLASLt handle	—
cuda_pipeline/cublas_algo_deterministic.rs	gpu_dqn_trainer.rs, gpu_curiosity_trainer.rs, gpu_iqn_head.rs (7 consumers)	Wired	Deterministic cuBLASLt algo selection	—
cuda_pipeline/gpu_weights.rs	trainer/mod.rs, gpu_dqn_trainer.rs, hyperopt/adapters/dqn.rs (11 consumers)	Wired	Weight tensor layout constants	—
cuda_pipeline/gpu_action_selector.rs (GpuActionSelector)	trainer/mod.rs, fused_training.rs, gpu_backtest_evaluator.rs	Wired	Epsilon-greedy + routed action selection	—
cuda_pipeline/gpu_monitoring.rs (GpuMonitor)	fused_training.rs, trainer/metrics.rs, trainer/training_loop.rs	Wired	GPU-side monitoring_reduce launch	—
cuda_pipeline/gpu_training_guard.rs	gpu_dqn_trainer.rs, trainer/training_loop.rs	Wired	NaN / gradient anomaly guard	—
cuda_pipeline/gpu_backtest_evaluator.rs (GpuBacktestEvaluator)	trainer/metrics.rs (eval path)	Wired	Per-epoch GPU backtest evaluation	—
cuda_pipeline/gpu_experience_collector.rs	fused_training.rs via TradeStats, financials.rs	Wired	Per-trade stats from experience buffer	—
cuda_pipeline/gpu_curiosity_trainer.rs	fused_training.rs	Wired	Curiosity-driven intrinsic reward	—
cuda_pipeline/gpu_walk_forward.rs (GpuWalkForwardData)	trainer/mod.rs (lazy-init field)	Wired	GPU walk-forward data structure	—
cuda_pipeline/gpu_her.rs (GpuHer)	fused_training.rs	Wired	Hindsight Experience Replay	—
cuda_pipeline/gpu_iql_trainer.rs (GpuIqlTrainer)	fused_training.rs (dual IQL instances)	Wired	IQL value network training	—
cuda_pipeline/gpu_iqn_head.rs (GpuIqnHead)	fused_training.rs	Wired	IQN distributional head	—
cuda_pipeline/gpu_attention.rs (GpuAttention)	fused_training.rs	Wired	Self-attention layer in DQN trunk	—
cuda_pipeline/gpu_tlob.rs (GpuTlob)	fused_training.rs (training: forward pre-trunk, backward+Adam Phase 6); trainer/metrics.rs (val: set_tlob_from_training per-epoch weight sync + GpuBacktestEvaluator::submit_dqn_step_loop_cublas per-gather TLOB forward)	Wired	D.8 Plan 2 Task 6C — OFI[32]→TLOB[16] single-head SDP attention, Xavier random init, trained end-to-end. ISV[60]=TLOB_REGIME_FOCUS_EMA written at epoch boundary.	—
cuda_pipeline/decision_transformer.rs (DecisionTransformer)	trainer/training_loop.rs	Wired	Decision Transformer pre-training step	—
cuda_pipeline/learning_health.rs (LearningHealth)	fused_training.rs, trainer/training_loop.rs, trainer/constructor.rs, trainer/metrics.rs, meta_q_network.rs (11 consumers)	Wired	Health-band tracking for adaptive LR / early stopping	—
cuda_pipeline/q_snapshot.rs	cuda_pipeline/mod.rs, gpu_dqn_trainer.rs	Wired	Q-value snapshot for double-DQN target	—
cuda_pipeline/meta_q_network.rs (MetaQNetwork)	trainer/constructor.rs, trainer/mod.rs, trainer/training_loop.rs	Wired	Meta-learning Q-network	—
cuda_pipeline/q_value_provider.rs (QValueProvider trait)	fused_training.rs impl, trainer/metrics.rs, hyperopt/adapters/dqn.rs	Wired	Trait for on-device Q-value computation in val path	—
cuda_pipeline/signal_adapter.rs	hyperopt/adapters/dqn.rs, hyperopt/adapters/{ppo,mamba2,tft,kan,liquid,tlob,diffusion,tggn,xlstm}.rs	Wired	backtest_fitness scoring used by all hyperopt adapters	—
cuda_pipeline/multi_gpu.rs (MultiGpuConfig)	trainer/constructor.rs, trainer/mod.rs	Wired	Multi-GPU detection and layout	—
cuda_pipeline/gpu_ppo_collector.rs (GpuPpoExperienceCollector)	trainers/ppo.rs	OUT-of-DQN-scope	PPO-only consumer; correct-as-is	—
cuda_pipeline/gpu_statistics.rs (GpuStatistics)	declared pub in cuda_pipeline/mod.rs; no call site outside the file	Orphan	GpuStatistics::compute never called; BatchStatistics struct unused	Plan 2 D.2 audits + decides whether to wire into val path or delete
cuda_pipeline/cublaslt_debug.rs	mod cublaslt_debug (private) inside cuda_pipeline/mod.rs; test-only	OUT-of-DQN-scope	Private debug/test module; not a public feature	—

CUDA Kernels (.cu files)

Module / kernel	Consumer path	Classification	Notes	Action
epsilon_greedy_kernel.cu (epsilon_greedy_select, epsilon_greedy_routed)	gpu_action_selector.rs	Wired	Action selection kernel	—
monitoring_kernel.cu (monitoring_reduce)	gpu_monitoring.rs	Wired	12-bin action count + Q stats	—
c51_loss_kernel.cu	gpu_dqn_trainer.rs (C51 loss compute)	Wired	C51 distributional loss	—
c51_grad_kernel.cu	gpu_dqn_trainer.rs (C51 backward)	Wired	C51 gradient backward	—
mse_loss_kernel.cu	gpu_dqn_trainer.rs (MSE warmup loss)	Wired	MSE loss for warmup phase	—
mse_grad_kernel.cu	gpu_dqn_trainer.rs (MSE backward)	Wired	MSE gradient backward	—
iqn_dual_head_kernel.cu	gpu_iqn_head.rs	Wired	IQN dual-head quantile computation	—
iqn_cvar_kernel.cu	gpu_iqn_head.rs	Wired	IQN CVaR risk measure	—
iql_value_kernel.cu	gpu_iql_trainer.rs	Wired	IQL value network kernel	—
cql_grad_kernel.cu	gpu_dqn_trainer.rs (CQL backward)	Wired	Conservative Q-learning gradient	—
experience_kernels.cu	gpu_backtest_evaluator.rs (BACKTEST_DQN_CUBIN), gpu_experience_collector.rs	Wired	Experience buffer ops + DQN backtest forward	—
reward_shaping_kernel.cu	gpu_dqn_trainer.rs (reward shaping step)	Wired	Per-sample reward shaping	—
nstep_kernel.cu	gpu_dqn_trainer.rs (n-step return)	Wired	N-step Bellman target	—
backward_kernels.cu	gpu_dqn_trainer.rs (weight backward)	Wired	Weight gradient accumulation	—
bias_kernels.cu	gpu_dqn_trainer.rs (bias grad)	Wired	Bias gradient kernels	—
relu_mask_kernel.cu	gpu_dqn_trainer.rs (activation backward)	Wired	ReLU activation mask	—
ema_kernel.cu	gpu_dqn_trainer.rs (target-net soft update)	Wired	EMA target-network update	—
q_stats_kernel.cu	gpu_dqn_trainer.rs (Q-value stats for ISV)	Wired	Per-branch Q-value statistics	—
dqn_utility_kernels.cu	gpu_dqn_trainer.rs (multiple utility ops)	Wired	Misc DQN GPU ops (advantage std, PopArt, etc.)	—
graph_utility_kernels.cu	gpu_dqn_trainer.rs (CUDA graph utility ops)	Wired	Graph capture helper kernels	—
grad_decomp_kernel.cu	gpu_dqn_trainer.rs, called via fused_training.rs::grad_decomp_launch_c51	Wired	C51 gradient decomposition (Task 2.0 diagnostic)	—
branch_grad_balance_kernel.cu	gpu_dqn_trainer.rs, called via fused_training.rs::launch_branch_grad_balance	Wired	Per-branch gradient L2-norm balancing	—
mamba2_temporal_kernel.cu (mamba2_scan_projected_fwd, mamba2_scan_projected_bwd, isv_temporal_route, etc.)	gpu_dqn_trainer.rs::mamba2_forward + mamba2_backward, called in fused training loop (adam_grad child graph)	Wired (grad-check validated, Plan 2 Task 2)	Mamba2 selective SSM forward + backward (both paths implemented). D.1: kernel-level reference check + non-zero gradient propagation test confirm backward is not silently no-oping.	—
attention_kernel.cu	gpu_attention.rs	Wired	Scaled dot-product attention forward	—
attention_backward_kernel.cu	gpu_attention.rs, gpu_tlob.rs (reused for grad_norm+Adam kernels)	Wired	Attention backward pass	—
tlob_kernel.cu (tlob_sdp_forward, tlob_sdp_backward, tlob_write_states, tlob_read_grad_from_states)	gpu_tlob.rs	Wired	D.8 Plan 2 Task 6C — TLOB SDP forward/backward, state scatter/read kernels	—
training_guard_kernel.cu (training_guard_check_and_accumulate)	gpu_training_guard.rs	Wired	NaN / guard check kernel	—
backtest_env_kernel.cu (backtest_env_step)	gpu_backtest_evaluator.rs (ENV_CUBIN)	Wired	Backtest environment step	—
backtest_metrics_kernel.cu	gpu_backtest_evaluator.rs (METRICS_CUBIN)	Wired	Backtest metrics reduction	—
backtest_plan_kernel.cu (backtest_plan_diag_reduce, backtest_plan_state_isv)	gpu_backtest_evaluator.rs (BACKTEST_PLAN_CUBIN)	Wired	Plan-ISV diagnostic reduction in val. D.6 (Plan 2 Task 6A): backtest_plan_state_isv now writes pisv[6] = remaining_fraction; stride updated to SL_PORTFOLIO_PLAN_DIM=7.	—
backtest_forward_ppo_kernel.cu (backtest_forward_ppo_kernel)	gpu_backtest_evaluator.rs::evaluate_ppo (PPO_FORWARD_CUBIN)	OUT-of-DQN-scope	PPO-path backtest only; correct-as-is	—
backtest_forward_supervised_kernel.cu (signal_to_action_kernel)	gpu_backtest_evaluator.rs::evaluate_supervised (SUPERVISED_SIGNAL_CUBIN)	OUT-of-DQN-scope	Supervised-model backtest only; correct-as-is	—
signal_adapter_kernel.cu	signal_adapter.rs (loaded at construction)	Wired	Signal-to-action ISV bus adapter	—
statistics_kernel.cu (batch_statistics)	gpu_statistics.rs (loaded but GpuStatistics has no call sites outside the file)	Orphan	Kernel compiled, wrapper struct exists, but never instantiated by any consumer	Plan 2 D.2 — same resolution as gpu_statistics.rs
ppo_experience_kernel.cu	gpu_ppo_collector.rs	OUT-of-DQN-scope	PPO collector; correct-as-is	—
her_episode_kernel.cu	gpu_her.rs	Wired	HER episode boundary kernel	—
her_relabel_kernel.cu	gpu_her.rs	Wired	HER goal relabelling kernel	—
curiosity_training_kernel.cu	gpu_curiosity_trainer.rs	Wired	Curiosity intrinsic reward training	—
curiosity_inference_kernel.cu	gpu_curiosity_trainer.rs	Wired	Curiosity inference forward	—
ensemble_kernels.cu	batched_forward.rs / batched_backward.rs (ensemble ops in fused trainer)	Wired	Ensemble head combination	—
dt_kernels.cu	decision_transformer.rs (DT_CUBIN)	Wired	Decision Transformer GPU ops	—
trade_stats_kernel.cu	gpu_experience_collector.rs	Wired	Per-trade statistics accumulation	—
backtest_env_kernel.cu (backtest_env_step)	gpu_backtest_evaluator.rs	Wired	Already listed above	—
common_device_functions.cuh	prepended to all kernels at compile time by build.rs	Wired	Shared device helpers (BF16 wrappers, warp reduce)	—
state_layout.cuh	included by kernels that reference state buffer layout	Wired	Named index constants for state buffer	—
trade_physics.cuh	included by backtest_env_kernel.cu and others at compile time	Wired	Kelly / physics helpers shared across kernels. kelly_position_cap warm-branch deadlock fixed: effective_kelly = max(kelly_f, warmup_floor) so the cap can never collapse to zero when balanced trades naturally yield kelly_f=0 at maturity=1; preserves design intent (Kelly drives sizing once stats mature with real edge) while preventing the bootstrap deadlock that produced 100% Long/Short→Flat conversion in val.	—

Model Modules — Supervised-Only (OUT-of-DQN-scope per Part E)

Module / kernel	Consumer path	Classification	Notes	Action
xlstm/mod.rs + xlstm/trainable.rs	ensemble/adapters/xlstm.rs, hyperopt/adapters/xlstm.rs	OUT-of-DQN-scope	Supervised consumers; DQN does not import this. Part E: keep.	—
kan/mod.rs + kan/trainable.rs	hyperopt/adapters/kan.rs only	OUT-of-DQN-scope	KAN splines in DQN trunk (batched_forward/backward) are inline in gpu_dqn_trainer.rs and do not import crate::kan. kan/ wraps ml-supervised KAN for hyperopt. Part E: keep.	—
tgnn/mod.rs + tgnn/trainable_adapter.rs	hyperopt/adapters/tggn.rs, risk/mod.rs	OUT-of-DQN-scope	TGNN for supervised + risk path	—
tlob/mod.rs + tlob/trainable_adapter.rs	hyperopt/adapters/tlob.rs, data_loaders/tlob_loader.rs	Partial / OUT-of-DQN-scope	TLOB transformer has hyperopt adapter; DQN trunk wiring scheduled for Plan 2 D.8	Plan 2 D.8 wires TLOB trunk into DQN
diffusion/mod.rs + diffusion/trainable.rs	hyperopt/adapters/diffusion.rs, trainers/dqn/fused_training.rs (data aug only)	Partial	Diffusion model used for data augmentation in fused training but lacks a gradient backward path through DQN trunk	Plan 2 tracks full integration
ppo/mod.rs + ppo/trainable_adapter.rs	trainers/ppo.rs, gpu_ppo_collector.rs	OUT-of-DQN-scope	PPO training path; not DQN	—

Temporal/Recurrent Modules — DQN Integration Status

Module / kernel	Consumer path	Classification	Notes	Action
mamba/mod.rs + mamba/trainable_adapter.rs	trainers/mamba2.rs, model_factory.rs, inference_validator.rs; Mamba2 temporal kernel wired fully in gpu_dqn_trainer.rs	Wired (DQN trunk) + OUT-of-DQN-scope (supervised trainers/mamba2.rs)	Forward and backward both implemented in mamba2_temporal_kernel.cu	—
liquid/mod.rs + liquid/adapter.rs	trainers/liquid.rs supervised path only	Deleted (DQN)	D.7 audit (2026-04-24): liquid_tau_rk4_step kernel was a mathematical identity — ODE f(x)=(1/tau)*(1-x) has fixed point x=1.0; initialised to 1.0, it never deviated. liquid_mod_buf always held [1.0,1.0,1.0,1.0], so velocity_mod in c51_grad_kernel multiplied spread_scale by 1.0 unconditionally. No ISV slot existed (violates §4.C.6). Kernel, buf, and call removed. LiquidTrainableAdapter (supervised path) unchanged.
tft/mod.rs + tft/trainable_adapter.rs	trainers/tft/trainer.rs has full forward + backward_step(); TrainableTFT::backward() returns error with explicit message that backward is via TFTTrainer::train()	Partial	TFT backward available via dedicated trainer; UnifiedTrainable::backward() slot returns informational error (not a stub — documented routing)	Plan 4 E.1/E.3 — DQN concept adoption; full backward path tracked in task #76

Data Pipeline Modules

Module / kernel	Consumer path	Classification	Notes	Action
fxcache.rs (discover_and_load, FxCacheData)	train_baseline_rl.rs, trainers/dqn/data_loading.rs	Wired	Primary training data format	—
feature_cache.rs	fxcache.rs (cache key), hyperopt/adapters/dqn.rs, smoke tests	Wired	Cache key generation for .fxcache auto-discovery	—
data_pipeline/ (DatasetManager, PreparedDataset)	lib.rs prelude re-export; integration/ modules	Wired	Dataset management abstraction	—
data_loaders/dbn_sequence_loader.rs	data_loaders/mod.rs; internally used by dbn_tick_adapter.rs	Wired	DBN sequence loading for supervised models	—
data_loaders/dbn_tick_adapter.rs	data_loaders/mod.rs	Wired	DBN tick-to-bar adapter	—
data_loaders/streaming_dbn_loader.rs	data_loaders/mod.rs; tests/streaming_pipeline_edge_cases.rs (test consumer)	Partial	Test-only consumer at tests/streaming_pipeline_edge_cases.rs; if streaming DBN load is genuinely needed for production (memory-efficient data loading), wire into production path in a follow-up task; otherwise schedule for deletion after test utility is moved to dbn_sequence_loader.	Reclassified Partial — test consumer confirmed
data_loaders/tlob_loader.rs (TLOBDataLoader)	data_loaders/mod.rs; no external call sites found	Orphan	TLOB data loader built but not called outside its own module and mod.rs	Plan 2 D.8 may wire; surface for user review
training/unified_trainer.rs (UnifiedTrainable trait)	diffusion/trainable.rs, tlob/trainable_adapter.rs, mamba/trainable_adapter.rs, tft/trainable_adapter.rs, hyperopt adapters	Wired	Unified training trait for supervised models	—
training/unified_data_loader.rs	—	—	Deleted (Task 6): zero production + zero test consumers confirmed. pub mod unified_data_loader removed from training.rs.	DELETED
training/orchestrator.rs (TrainingOrchestrator)	—	—	Deleted (Task 6): zero production + zero test consumers confirmed. pub mod orchestrator removed from training.rs.	DELETED
training_pipeline.rs (ProductionTrainingError)	lib.rs (two From impls: MLError + MLSafetyError); TrainingPipeline struct has no external instantiation	Partial	ProductionTrainingError used via From impl in lib.rs; TrainingPipeline struct itself may be dead — follow-up to extract error enum and delete struct separately.	Reclassified Partial — error type wired, struct possibly unused
training_profile.rs (DqnTrainingProfile)	train_baseline_rl.rs, hyperopt/adapters/dqn.rs, smoke tests	Wired	TOML-backed hyperparameter profiles	—
walk_forward.rs (FoldRange, WalkForwardConfig)	train_baseline_rl.rs, validation/harness.rs	Wired	Walk-forward fold management	—

Evaluation / Validation Modules

Module / kernel	Consumer path	Classification	Notes	Action
validation/mod.rs (ValidatableStrategy)	validation/harness.rs, hyperopt/adapters/dqn.rs	Wired	Validation trait + entry points	—
validation/harness.rs	hyperopt/adapters/dqn.rs, trainer/metrics.rs	Wired	Walk-forward validation harness	—
validation/financial.rs	validation/harness.rs	Wired	Financial metric computation for val	—
validation/regime_analysis.rs	validation/harness.rs	Wired	Regime-conditional eval	—
validation/ppo_adapter.rs	validation/harness.rs	OUT-of-DQN-scope	PPO-specific validation path	—
evaluation/mod.rs	hyperopt/adapters/dqn.rs, trainer/metrics.rs, gpu_backtest_evaluator.rs	Wired	DQN evaluation engine	—
backtesting/ (GpuBacktestEvaluator, BarrierBacktester)	trainer/metrics.rs, hyperopt/adapters/*	Wired	Backtesting framework	—

Supporting Infrastructure Modules

Module / kernel	Consumer path	Classification	Notes	Action
checkpoint/mod.rs	model_registry.rs	Wired	Checkpoint save/load	—
inference.rs	multiple service consumers (67 files by grep)	Wired	Inference entry points	—
inference_validator.rs (InferenceValidator)	—	—	Deleted (Task 6): zero production + zero test consumers confirmed. No pub mod declaration found in any file.	DELETED
model_factory.rs	mamba path, lib.rs	Wired	Model construction factory	—
model_loader_integration.rs (ModelLoaderTrait)	—	—	Deleted (Task 6): zero production + zero test consumers confirmed. No pub mod declaration found in any file.	DELETED
model_registry.rs	lib.rs, various consumers (3 by grep)	Wired	Model lifecycle registry	—
registry/mod.rs	consumed by integration and service layer (10 consumers)	Wired	Operational model lifecycle	—
hyperopt/ (campaign, adapters)	trainer/training_loop.rs, hyperopt smoke test	Wired	Bayesian hyperparameter optimisation	—
ensemble/	integration path, service layer	Wired	Multi-model ensemble	—
integration/	service layer	Wired	Service integration layer	—
flash_attention/mod.rs (FlashAttention3)	transformers/hft_transformer.rs, trainers/tft/config.rs, benchmark/tft_benchmark.rs	OUT-of-DQN-scope	TFT / transformer path only; DQN uses gpu_attention.rs	—
features/ (FeatureVector, extract_ml_features)	train_baseline_rl.rs, DQN trainer	Wired	Feature extraction	—
microstructure/	trainers/dqn/features.rs, OFI pipeline	Wired	VPIN, Kyle lambda, order-flow imbalance	—
regime/mod.rs	DQN features, data pipeline (50 consumers)	Wired	Regime detection + signal	—
regime_detection/mod.rs	lib.rs declaration only; zero consumers of the facade; zero direct uses of ml_regime_detection:: anywhere else	Orphan	Thin facade re-exporting ml-regime-detection crate. Zero consumers of the facade, zero direct uses of ml_regime_detection:: anywhere. The underlying ml-regime-detection/ workspace crate exists but appears unused. Deleting a whole workspace crate exceeds Task 6 scope — scheduled for a separate crate-cleanup task. Potentially superseded by ml_regime/ crate (50 consumers).	CRATE-LEVEL-FOLLOWUP
risk/mod.rs (crate-level)	DQN trainer constructor via DrawdownMonitor, KellyCriterionOptimizer	Wired	Risk management in training	—
preprocessing.rs	features/, data pipeline	Wired	Log-return normalisation, outlier clipping	—
labeling/mod.rs	features/, supervised training	Wired	Triple-barrier labelling	—
security/mod.rs	lib.rs (1 consumer)	Partial	ML security / prediction validation; no DQN consumer found	Surface for user review
stress_testing/mod.rs	dqn/stress_testing.rs, ppo/stress_testing.rs	Wired	Stress-testing framework	—
paper_trading/mod.rs	lib.rs declaration; tests/paper_trading_integration_test.rs (test consumer)	Partial	Test-only consumer at tests/paper_trading_integration_test.rs. Production trading-service uses its own paper_trading_executor (distinct module); this facade exists only for ML crate tests.	Reclassified Partial — test consumer confirmed
observability/mod.rs	integration/performance_monitor.rs	Wired	ML observability hooks	—
deployment/	lib.rs, integration layer	Wired	Model deployment + A/B testing	—
universe/mod.rs	lib.rs, integration layer (5 consumers)	Wired	Asset universe management	—
asset_selection/mod.rs	lib.rs, integration layer (referenced in universe path)	Wired	Asset selection logic	—
batch_processing.rs	lib.rs, integration/ (multiple consumers)	Wired	Batch ML operations	—
bridge.rs	lib.rs, trading-ML bridge	Wired	ML-Financial type bridge	—
data_loader.rs	lib.rs, training service	Wired	Legacy data loader shim	—
data_validation/mod.rs	lib.rs, data pipeline	Wired	Input validation for data pipeline	—
explainability/mod.rs	lib.rs only (1 consumer)	Partial	SHAP / attribution; no DQN call site found	Surface for user review
benchmarks.rs	lib.rs, benchmark harness	OUT-of-DQN-scope	Benchmark entry point; not production training	—
benchmark/	benchmark harness	OUT-of-DQN-scope	GPU / model benchmarks; not production DQN path	—
portfolio_transformer.rs	—	—	Deleted (Task 6): zero production + zero test consumers confirmed. pub mod portfolio_transformer removed from lib.rs.	DELETED

Summary

Updated after Task 6 cleanup (2026-04-24): 5 confirmed-orphan files deleted, 3 Orphan rows reclassified Partial, 1 Orphan reclassified with crate-level follow-up action. Plan 1 Task 8 (revised, 2026-04-24): adaptive_controller.rs renamed → adaptive_monitor.rs; AdaptiveController replaced with read-only AdaptiveMonitor per spec §4.C.6 (GPU drives, CPU reads).

Plan 3 Task 1 C.2 (2026-04-24): reward_component_ema_kernel.cu + RewardComponentMonitor added. 6 ISV reward-EMA slots [63..69) allocated; fingerprint shifted [61..63) → [69..71); ISV_TOTAL_DIM 63 → 71. experience_kernels.cu extended with reward_components_per_sample [N*L, 6] output parameter. 2 new Wired rows.

Plan 3 Task 3 B.2 (2026-04-24): trade_rate_ema_kernel.cu added. 2 new ISV slots [71] TRADE_ATTEMPT_RATE_EMA and [72] TRADE_TARGET_RATE; fingerprint shifted [69..71) → [73..75); ISV_TOTAL_DIM 71 → 75. experience_kernels.cu gains flat_to_pos_per_sample [N*L] output + 2 ISV slot-idx scalar parameters; a novelty-scaled bonus (conviction × vol_proxy × novelty) is added to reward and captured in rc[5] at every Flat→Positioned transition. TRADE_TARGET_RATE frozen in training_loop.rs at epoch 5 from measured attempt-EMA (min 0.001). 2 new Wired rows.

Plan 3 Task 5 C.4 (2026-04-24): Temporal timing bonus on trade exit. No new ISV slot — accumulates into rc[5] bonus slot via += (B.2 at entry, C.4 at exit: different (i,t) slots). New portfolio-state slot PS_PEAK_PNL_BAR=38; PS_STRIDE 38 → 39; PORTFOLIO_STRIDE 38 → 39 in all consumers (experience_kernels.cu, trade_stats_kernel.cu, gpu_experience_collector.rs, gpu_dqn_trainer.rs::launch_kelly_cap_update, ml-core::state_layout.rs). Peak bar snapshotted alongside every MAX_PNL update, reset at every MAX_PNL reset site (entry, reverse, fold hard-reset, trade-complete soft-reset). timing_bonus = shaping_scale × (bars_early / hold_time) × |final_pnl| × conviction_core where bars_early = max(0, segment_hold_time − PS_PEAK_PNL_BAR). No tuned multiplier. 1 new Wired row.

Plan 3 Task 6a D.4a (2026-04-24): Persistence credit on profitable drawdown recovery. No new ISV slot — accumulates into rc[5] bonus slot via += (now shared by B.2 entry, C.4 exit-timing, D.4a exit-persistence; all three fire at different (i,t) slots per trade, so += is idempotent). New portfolio-state slot PS_INTRA_TRADE_MIN_PNL=39; PS_STRIDE 39 → 40; PORTFOLIO_STRIDE 39 → 40 in all consumers (experience_kernels.cu, trade_stats_kernel.cu, gpu_experience_collector.rs, gpu_dqn_trainer.rs::launch_kelly_cap_update, ml-core::state_layout.rs). MIN_PNL tracked per-bar in the same block as MAX_PNL (fminf against pnl_pct), reset at every MAX_PNL reset site (entry, reverse, fold hard-reset, trade-complete soft-reset). persist_bonus = shaping_scale × conviction_core × drawdown_depth × tanh(reward / max(1e-4, drawdown_depth)) where drawdown_depth = max(0, −PS_INTRA_TRADE_MIN_PNL); gated on reward > 0 && drawdown_depth > 1e-6. Self-scaling tanh ratio: saturates +1 when reward ≫ drawdown (full credit for riding through), ≈0 when reward ≈ drawdown (trivial recovery). No tuned multiplier. 1 new Wired row.

Plan 3 Task 4 B.4 (2026-04-24): plan_threshold_update_kernel.cu + PlanThresholdMonitor added. PRODUCER-ONLY upgrade for ISV[PLAN_THRESHOLD_INDEX=49]: cold-start 0.5 constructor write preserved, but per-epoch GPU kernel now overwrites the slot with max(0.1, 0.5 × ISV[READINESS_EMA_INDEX=75]). New ISV slot [75] READINESS_EMA (cold-start 1.0 so the derived threshold matches the legacy 0.5 default until the EMA adapts). Fingerprint shifted [73..75) → [76..78); ISV_TOTAL_DIM 75 → 78. experience_kernels.cu gains readiness_per_sample [N*L] output parameter; the kernel writes the broadcast readiness_ptr[0] value into every reached (i,t) slot, race-free. isv_plan_threshold reset category flipped SchemaContract → FoldReset; isv_readiness_ema registered as FoldReset. Consumer kernels (4 sites in experience_kernels.cu + 1 in backtest_plan_kernel.cu) unchanged. 2 new Wired rows.

Plan 3 Task 6b D.4b (2026-04-24): Regime-shift penalty — punish trades held past a detected regime flip. No new ISV slot — subtracts from rc[5] bonus slot via −= (cancels against B.2/C.4/D.4a positive contributions within rc[5]; that cross-flow is exactly what ISV[68] REWARD_BONUS_EMA is meant to track). New portfolio-state slot PS_REGIME_SHIFT_BAR=40; PS_STRIDE 40 → 41; PORTFOLIO_STRIDE 40 → 41 in all consumers (experience_kernels.cu, trade_stats_kernel.cu, gpu_experience_collector.rs, gpu_dqn_trainer.rs::launch_kelly_cap_update, ml-core::state_layout.rs). New CUDA-side ISV macros ISV_Q_DIR_ABS_REF_IDX=21 and ISV_SHARPE_EMA_IDX=22 added to state_layout.cuh (mirror authoritative Q_DIR_ABS_REF_INDEX / SHARPE_EMA_INDEX in Rust trainer). Detector inside the active-trade block (after MIN/MAX_PNL update) writes PS_REGIME_SHIFT_BAR = hold_time on the FIRST bar where |isv[11] − PS_PLAN_ENTRY_REGIME| > clamp(0.25 × |clamp(sharpe, −2, 2)|, 0.05, 0.5); first-shift-only (short-circuits on non-zero PS_REGIME_SHIFT_BAR). Consumer in segment_complete block IMMEDIATELY AFTER D.4a: bars_late = max(0, saved_hold_time − PS_REGIME_SHIFT_BAR); penalty = shaping_scale × conviction_core × (bars_late / saved_hold_time) × ISV[21] × |reward|; reward −= penalty; rc[5] −= penalty; then defensively resets PS_REGIME_SHIFT_BAR = 0 (in addition to the 5 lifecycle resets at entry/reverse/fold hard-reset/trade-complete soft-reset). Adaptive threshold tightens when |sharpe| is high (confident trading warrants early detection); loosens under noisy training. Q_DIR_ABS_REF self-scaling matches B.1 opp-cost pattern. No tuned multiplier. 1 new Wired row.

Plan 1 Tasks 12/15/16 + pre-allocation (2026-04-24): No new modules added. Changes are ISV slot allocation + consumer migration only. Task 15 confirmed no-op (IQL_BRANCH_SCALE_FLOOR_INDEX already serves conviction-floor role). Tasks 12 and 16 migrate cql_alpha and plan-threshold consumers from config fields / hardcoded literals to ISV slots. 8 new ISV slots allocated ([39..47)); fingerprint tail moves from [37..39) to [47..49); ISV_TOTAL_DIM 39 → 49. GpuDqnTrainConfig gains total_epochs field (written to TOTAL_EPOCHS_INDEX at construction). write_isv_signal_at bound extended from ISV_DIM to ISV_TOTAL_DIM to allow writes beyond slot 22.

Plan 2 Task 6B D.3 (2026-04-24): IQL value head widened from 1 to 2 outputs (V_short + V_long). v_out_buf shape [B] → [B*2]. gemm_fwd_v M=1→2, gemm_bwd_dw3 M=1→2, gemm_bwd_dh2 K=1→2. W3 param block [H*1] → [H*2], b3 [1] → [2]. total_params += H+1. iql_expectile_loss kernel extended with num_heads argument. 4 consumer kernels in iql_value_kernel.cu updated to read v_out[b*2+0] + v_out[b*2+1]. Checkpoint compat break — retrain required.

Plan 4 Task 5 Mode A E.5 (2026-04-24): attention_focus_ema_kernel.cu + AttentionMonitor added. 3 new ISV slots [87] VSN_MAG_EMA, [88] VSN_DIR_EMA, [89] MAMBA2_RETENTION_EMA tail-appended; fingerprint shifted [85..87) → [90..92); ISV_TOTAL_DIM 87 → 92. Light, ISV-diagnostic-only; no model parameters added, no checkpoint break. Single-thread cold-path EMA kernel takes 3 host scalars by value at the per-epoch HEALTH_DIAG site. New mamba2_retention_mean(batch_size) accessor on GpuDqnTrainer/FusedTrainingCtx mirrors per_branch_vsn_mean (cold-path host-side dtoh + mean abs of mamba2_h_enriched). Adaptive α convention matches Plan 3 producers. All three slots FoldReset to 0.0 (mirror constructor cold-start). 2 new Wired rows. Mode B (full per-feature-group VSN ISV) blocks on Plan 4 Task 1 (E.1).

Plan 4 Task 4 E.4 (2026-04-24): state_encoder.rs + value_decoder.rs Rust API split. PURE Rust API refactor — no kernel changes, no parameter changes, no checkpoint break, no new ISV slot. New types: StateEncoder (wraps trunk encoder h_s1 + h_s2), ValueDecoder (per-branch FC + adv-logits, one instance per Branch::{Dir,Mag,Ord,Urg}), EncoderOutput, BranchDecoderOutput (carries q_per_action_dev_ptr + Plan 2 D.3 v_short_dev_ptr / v_long_dev_ptr pass-through fields). BatchedForward gains encoder_forward_only(...) and decoder_forward_only(branch_idx, ...) helpers, both lossless extractions of the existing dispatch sequence in forward_online_raw. The monolithic forward_online_raw stays callable and now routes through encoder_forward_only internally for the trunk portion (multi-stream branch fork/join unchanged). The new types are ADDITIVE attachment points for Plan 4 Task 1 (Full VSN, pre-encoder), Task 3 (multi-Q IQN, decoder), and Task 6 (auxiliary heads, decoder peer). 2 new Wired rows.

Plan 4 Task 2c.1 (2026-04-24): grn_kernel.cu added — Gated Residual Network forward + backward kernels (8 entry points: grn_elu_inplace, grn_glu_forward, grn_residual_layernorm_forward, grn_layernorm_backward_dx, grn_layernorm_backward_dgamma_dbeta_p1, grn_layernorm_backward_dgamma_dbeta_p2, grn_glu_backward, grn_elu_backward). Linear_a / Linear_b / Linear_residual remain cuBLAS GEMMs (no new kernels). Module is additive — ZERO production callers in this commit; classified Orphan-by-design. Task 2c.3+4 wires it into the trunk encoder forward + backward path. LayerNorm backward implements the FULL Jacobian (d_x = (1/H) * rstd * (H * d_out_g - s1 - normed * s2)), NOT the simplified element-wise form d_x = d_out * gamma * rstd from attn_layer_norm_bwd_dx — the simplified form is correct only when a parallel residual gradient absorbs the missing s1/s2 terms, which the GRN trunk-encoder backward does not have. d_gamma / d_beta use a two-phase per-block partial reduction + final-reduce pass (feedback_no_atomicadd.md compliant). Layout convention: row-major [B, H] (matches dt_layernorm_kernel, intentionally different from attention_kernel.cu's [D, B] col-major — trunk buffers downstream of GRN are row-major, so per-row LN avoids a transpose). All reductions GPU-side per pearl_cold_path_no_exception_to_gpu_drives.md. Kernel registered in build.rs (kernel count 57 → 58, all compile clean under nvcc sm_80). No checkpoint break (no new params, no new ISV slot, no consumer wired). 1 new Orphan-by-design row (will reclassify Wired at 2c.3+4).

Plan 4 Task 2c.2 (2026-04-24): cuda_pipeline/gpu_grn.rs added — Rust wrapper around the 2c.1 kernels. One GrnBlock instance owns 5 save-for-backward state buffers (elu_post, linear_b_out, sigmoid_b, ln_mean, ln_rstd, ln_normed) plus 2 per-block partial-reduction scratch buffers (dgamma_partials, dbeta_partials) plus 8 CudaFunction handles loaded from GRN_CUBIN. forward() sequences grn_elu_inplace → DtoD save post-ELU → DtoD save linear_b_out → grn_glu_forward → grn_residual_layernorm_forward. backward() sequences grn_layernorm_backward_dx → grn_layernorm_backward_dgamma_dbeta_p1 → _p2 → grn_glu_backward → (caller's Linear_b backward GEMM) → grn_elu_backward. cuBLAS GEMMs (Linear_a, Linear_b, optional Linear_residual) and the residual-split aliasing are caller-side, mirroring the contract used by gpu_tlob.rs and gpu_attention.rs. ELU backward consumes the POST-activation per the kernel's docstring (recovers f'(x_pre) via the identity exp(x_pre) = x_post + 1 for x_post < 0); the wrapper DtoD-copies the post-ELU buffer into state.elu_post immediately after the in-place ELU launch. has_residual_projection flag distinguishes h_s1 (SD→SH1, requires Linear_residual) from h_s2 (SH1→SH2, identity residual). LN dgamma/dbeta phase-1 partial-reduction grid sized at construction to ceil(max_batch/256) blocks; backward() verifies the runtime batch fits. GRN_CUBIN reference added in gpu_dqn_trainer.rs alongside the other Plan 4 kernel cubins; pub mod gpu_grn added to cuda_pipeline/mod.rs. Module is dead code in this commit — ZERO production callers; classified Orphan-by-design (will reclassify Wired at 2c.3+4 alongside the 2c.1 kernel entry above). No checkpoint break (no new params, no new ISV slot). cargo build clean at 11 warnings (unchanged), kernel cubin loads via the existing CudaFunction infrastructure. 1 new Orphan-by-design row.

Plan 4 Task 2b (2026-04-24): layout_fingerprint_seed() extended to cover param-tensor structural layout in addition to the existing ISV-slot section. 86 new entries PARAM_<NAME>=<idx> plus PARAM_TOTAL_TENSORS=86 appended to the seed string, mirroring compute_param_sizes() order one-for-one. Names match the in-code comments at each tensor index (PARAM_W_S1=0, PARAM_B_S1=1, … PARAM_W_PLAN_OUT=84, PARAM_B_PLAN_OUT=85). Pragmatic Option A scope: tensor names + positions only, NOT runtime sizes — shared_h1/shared_h2/adv_h/value_h/num_atoms/bottleneck_dim/market_dim are runtime config and can't be embedded in a const fn. Size mismatches between checkpoint and current binary are caught separately by safetensors deserialization. After this commit, ANY structural reshuffle (insert/delete/reorder/rename a tensor) changes LAYOUT_FINGERPRINT_CURRENT and triggers fail-fast at checkpoint load. Prerequisite for Plan 4 Task 2c (GRN ADOPT) — without 2b, GRN's tensor insertion would pass silently through checkpoint load and corrupt downstream state. ISV-slot portion of the seed unchanged. New fingerprint value: 0xa504d3c2f275b8af. Behavior change zero; this commit IS a checkpoint break by intent. No new module / kernel / ISV slot — pure contract-coverage extension.

Plan 4 Task 2c.3b (2026-04-25): encoder forward swap to GRN (forward path runtime restored). BatchedForward::encoder_forward_only panic gate replaced with the GRN composition: cuBLAS Linear_a (+ bias) → GrnBlock::forward_raw_phase1 (in-place ELU + DtoD save into state.elu_post) → cuBLAS Linear_b (+ bias) → cuBLAS Linear_residual (h_s1 only — h_s2 uses identity residual through h_s1_ptr directly) → GrnBlock::forward_raw_phase2 (DtoD save linear_b_out + GLU forward + residual+LN forward). Same sequence repeated for h_s2 with the appropriate weight indices ([7..13)). New BatchedForward::target_encoder_forward_only mirrors the online path but drives dedicated grn_h_s1_target / grn_h_s2_target instances so the target-network forward never clobbers the online's save-for-backward state (target writes to its own save buffers and forward_raw_phase{1,2} is called with save_for_backward=false so the DtoD copies are skipped — the buffers stay zero-initialised since target backward is never run). Two new dispatcher methods on GrnBlock (forward_raw_phase1/forward_raw_phase2) split the existing 5-step forward kernel sequence at the cuBLAS Linear_b boundary because the encoder must run Linear_a → ELU → Linear_b → Linear_residual → GLU+LN; no new GRN kernels. CublasGemmSet constructor allocates 4 GrnBlock instances and 14 GRN scratch buffers (linear_a, linear_b, residual_proj, glu × 2 GRN blocks × online + target — h_s2 omits residual_proj since residual is identity). All forward_* methods on CublasGemmSet migrated from &self to &mut self; trainer call sites adjusted (mostly disjoint-borrow refactors of let cublas = &self.cublas_forward to direct self.cublas_forward.method() calls so self.launch_* helpers can be invoked alongside). forward_target_raw and forward_online_f32 panic gates removed; bodies now delegate trunk encoding to target_encoder_forward_only / encoder_forward_only respectively. Three forward panic gates removed; three backward panic gates intentionally retained for Task 2c.3c (backward_full, apply_iqn_trunk_gradient, apply_ensemble_diversity_backward). Orphan kernel iqn_trunk_forward_kernel deleted from iqn_dual_head_kernel.cu along with its IqnHead::trunk_forward_kernel field, the load("iqn_trunk_forward_kernel") call, and the IqnKernels.trunk_fwd field; the cuBLAS fallback path in gpu_iqn_head.rs::execute_training_pipeline that produced ReLU activations on target_dueling weights is replaced with a hard error — IQN now requires set_cached_target_h_s2 to point at the buffer written by BatchedForward::target_encoder_forward_only, so online + target trunks share ONE GRN implementation. The dead cuBLAS scaffolding (trunk_l1_gemm, trunk_l2_gemm, trunk_h1_scratch, launch_trunk_bias_relu) marked #[allow(dead_code)] to keep the constructor wiring compact while making it explicit they are no longer reachable. cargo check clean at 11 warnings (baseline preserved); cargo build compiles all expected cubins (grn_kernel.cubin loads via the existing CublasGemmSet infrastructure and is consumed at runtime via the new GrnBlock instances). Smoke (cargo test … multi_fold_convergence --ignored) progresses through online + target encoder forward (logs 4× GrnBlock initialised confirming both online + DDQN sets allocate online + target instances) and panics in BatchedBackward::backward_full as expected — that gate is retained per task spec until 2c.3c swaps backward to the GRN backward chain. The audit row #11 (IQN target trunk orphan) is fully resolved by the deletion of iqn_trunk_forward_kernel and the cuBLAS-fallback removal; online and target now share BatchedForward::target_encoder_forward_only exclusively. 1 Orphan row retired; row count drops from 3 to 2 (h_s2_consumers_audit.md and dqn-wire-up-audit.md to be updated together when 2c.3c lands).

Plan 4 Task 2c.3c.1 (2026-04-24): GrnBlock::backward_raw_phase1 + backward_raw_phase2 added to gpu_grn.rs — Rust-side dispatcher split of the existing monolithic backward(). Same composition rationale as 2c.3b's forward phase split: the trunk encoder's cuBLAS Linear_b backward GEMM must run BETWEEN GLU backward and ELU backward (it produces the d_elu_out gradient that ELU backward consumes). Phase 1 sequences grn_layernorm_backward_dx → _dgamma_dbeta_p1 → _dgamma_dbeta_p2 → grn_glu_backward, writing d_pre_LN (= d_glu_out = d_residual via implicit pointer aliasing — pure pass-through, no kernel) and d_linear_b_out [B, 2*H] for the caller's Linear_b backward GEMM, plus accumulating d_gamma/d_beta. Phase 2 runs only grn_elu_backward, taking caller-provided d_elu_out (output of caller's Linear_b backward) and writing d_linear_a_out [B, H] for the caller's Linear_a backward GEMM. Both methods take raw u64 device pointers matching the forward phase methods' style (graph-capture friendly — bypasses cudarc's device_ptr event-tracking machinery). The existing monolithic backward() is left untouched (#[allow(dead_code)]) for surface-area minimization; will be deleted with 2c.5 cleanup if no production callers emerge. No CUDA kernel changes. Additive only — ZERO production callers in this commit; the three backward panic gates (backward_full, apply_iqn_trunk_gradient, apply_ensemble_diversity_backward) remain in place until 2c.3c.4 wires the new phase methods. cargo check clean at 11 warnings (baseline preserved); cargo build compiles 80 cubins (unchanged — no new kernel). +223 LOC. No new module / kernel / ISV slot / Orphan row.

Plan 4 Task 2c.3c.2 (2026-04-24): backward infrastructure additions — two helper methods that 2c.3c.4's GRN trunk backward wire-up requires. (1) CublasBackwardSet::launch_dw_only_no_bias — variant of launch_dw_only that runs only the cuBLAS dW GEMM and skips the 2-phase bias-grad reduction. The h_s1 GRN block's Linear_residual projection has no bias term (matches gpu_grn.rs::has_residual_projection=true's 7-tensor layout: w_a, b_a, w_b, b_b, w_residual, gamma, beta — note no b_residual); calling the existing launch_dw_only with db=0u64 would dereference NULL inside bias_grad_reduce_f32_p1. Same row-major→col-major sgemm convention as launch_dw_only (label "dW_only_no_bias" for cublasLt diagnostics), pub(crate) visibility matching its sibling. +35 LOC. (2) CublasGemmSet::saxpy_inplace(stream, n, alpha, x_ptr, y_ptr) — element-wise y[i] += alpha * x[i] for the h_s2 GRN's identity-residual gradient accumulation: after Linear_a_h_s2's backward writes d_h_s1 = d_linear_a_h_s2 @ W_a_h_s2 (overwrite, β=0), the residual path needs d_h_s1 += d_pre_ln_h_s2 to capture the identity-residual gradient. Implementation chose option 2 (existing dqn_saxpy_f32_kernel from dqn_utility_kernels.cu) over option 1 (cuBLAS cublasSaxpy_v2 legacy handle): the legacy cuBLAS handle is per-stream-bound state on PerStreamCublasHandles, so adopting it for backward saxpy would require either re-binding (race risk vs forward's binding) or a new shared handle slot — versus the existing kernel which is already loaded by GpuExperienceCollector for IQR/ensemble-variance Q-bonus saxpy and is graph-capturable via the standard launch_builder path. New saxpy_kernel: CudaFunction field on CublasGemmSet, loaded in CublasGemmSet::new from DQN_UTILITY_CUBIN (same cubin already include_bytes!'d by gpu_dqn_trainer.rs); +66 LOC including field, init block, and the pub fn saxpy_inplace method (256 threads/block, ceil(n/256) blocks, no grid-stride needed — kernel is already bounded by if (i < n)). Additive only — ZERO production callers in this commit; both methods sit dead-code until 2c.3c.4's wire-up commit. cargo check clean at 11 warnings (baseline preserved); cargo build unchanged — no new .cu files (saxpy uses the existing dqn_saxpy_f32_kernel symbol). No new module / kernel / ISV slot / Orphan row.

Plan 4 Task 2c.3c.3 (2026-04-24): GRN trunk backward scratch buffer allocation on GpuDqnTrainer. 8 new CudaSlice<f32> device buffers added (4 per GRN block × 2 blocks for h_s1 + h_s2): bw_grn_h_s2_d_pre_ln [B, SH2], bw_grn_h_s2_d_linear_b_out [B, 2SH2], bw_grn_h_s2_d_elu_out [B, SH2], bw_grn_h_s2_d_linear_a_out [B, SH2], plus the four h_s1 mirrors at [B, SH1] (and 2SH1 for the linear_b_out scratch — GLU's pre-activation has 2× hidden width because value-path + gate_pre are concatenated). The d_pre_LN buffer aliases both d_glu_out and d_residual (pure pass-through under GrnBlock::backward_raw_phase1's pointer aliasing convention from 2c.3c.1); d_linear_b_out carries the phase1 → caller's Linear_b backward GEMM gradient; d_elu_out carries Linear_b backward's output into phase2; d_linear_a_out carries phase2's output to the caller's Linear_a backward GEMM. Allocated alongside the existing alloc_backward_scratch call in the trainer constructor (mirrors that helper's stream.alloc_zeros::<f32>(size).map_err(...)? pattern). Total footprint at SH1=SH2=256, B=512: 8 × 5 × 256 × 4B = 5.0 MB per fold (negligible against existing ~32 MB per-branch cuBLAS workspaces). Each field is #[allow(dead_code)] until 2c.3c.4's apply_grn_trunk_backward_raw helper consumes them — the comment wired by Task 2c.3c.4 annotates each suppression. Additive only — ZERO production callers in this commit; the three backward panic gates (backward_full, apply_iqn_trunk_gradient, apply_ensemble_diversity_backward) remain in place until 2c.3c.4. Adam state for the new GRN tensors is auto-allocated since m_buf/v_buf are sized to total_params + cutlass_tile_pad which already includes the 13 GRN tensors after 2c.3a. cargo check clean at 11 warnings (baseline preserved); cargo build compiles 80 cubins (unchanged — no new kernel). +57 LOC. No new module / kernel / ISV slot / Orphan row.

Plan 4 Task 6 Commit B (2026-04-26): multi-task auxiliary heads behavioural wiring (E.6) — Commit A's dormant scaffolding becomes live. Aux-head forward + backward + loss-EMA producer + ISV-driven loss combination all dispatched in the production training step; the 8 aux param tensors at [119..127) train end-to-end from the next-bar MSE + 5-class regime CE losses, and aux gradient SAXPYs into the trunk's bw_d_h_s2 accumulator (ISV-scaled) so the trunk's GRN backward chain consumes a multi-task augmented signal. Wire-up sites (4 wire-points + 1 producer + 1 diag clause + 1 setter delegate): (1) Forward wire (launch_cublas_forward): aux_heads_forward() runs immediately AFTER forward_online_raw populates save_h_s2 and BEFORE stochastic depth scales it in-place — the aux backward re-reads save_h_s2 from the same buffer, so forward + backward see the pre-dropout activation symmetrically. Sequence: regime-label builder → next-bar label gather → next-bar forward → regime forward → next-bar MSE reduce → regime CE reduce. The next-bar label is extracted column-0 of next_states_buf (= log_return per pipeline.rs:684) via the existing strided_gather kernel into a DEDICATED aux_nb_label_buf [B] f32 field on GpuDqnTrainer — explicit fix for the WIP regression where aux_partial_nb_b2 was aliased here (the previous attempt regressed F1 74.56→65.12, F2 88.20→62.20). (2) Backward wire (launch_cublas_backward_to): aux_heads_backward(grad_base) runs AFTER backward_full + the three concat-accumulators fill bw_d_h_s2 and BEFORE encoder_backward_chain consumes it. Per-head backward kernels emit per-sample partials + per-sample dh_s2; for each of the 8 aux tensors the trailing aux_param_grad_reduce collapses partials → final, then dqn_saxpy_f32_kernel writes grad_base[tensor_idx] += aux_weight × final at offsets padded_byte_offset(119..127). grad_base is parameterized so the vaccine path's g_val_ptr scratch buffer also exercises aux backward correctly. Both per-head dh_s2 buffers SAXPY into bw_d_h_s2 with alpha = aux_weight. (3) ISV producer wire (training_loop.rs per-step block): launch_aux_heads_loss_ema(ema_alpha=0.05) runs alongside launch_h_s2_rms_ema / launch_vsn_mask_ema / launch_iqn_quantile_ema — single-thread single-block kernel reads the two scalar loss buffers populated by the captured forward graph and EMAs them into ISV[113] / ISV[114]. (4) Aux-weight refresh (training_loop.rs per-step block, AFTER ISV producer): CPU-side aux_weight = clamp(0.1 × ISV[LEARNING_HEALTH=12] × (1 - tanh(0.1 × training_sharpe_ema)), 0.05, 0.3). Pushed in via the new FusedTrainingCtx::set_aux_weight thin delegate to GpuDqnTrainer::set_aux_weight (same staleness contract as set_c51_alpha). The clamp [0.05, 0.3] is a numerical-stability bound (Invariant 1 carve-out per feedback_isv_for_adaptive_bounds.md) — 0.05 keeps gradient signal alive when learning_health == 0, 0.3 caps the share so aux can never dominate. The base coefficient 0.1 is the standard aux-loss base weight per spec §4.E.6 (only tuned multiplicand; SHARPE_SCALE = 0.1 matches Plan 3's controller convention). (5) HEALTH_DIAG aux clause (mid-HEALTH_DIAG block, AFTER reward_split clause, AFTER the producer launches per the post-a5f23b28f ordering invariant): emits HEALTH_DIAG[<epoch>]: aux [next_bar_mse=… regime_ce=… w=…] reading ISV[113] / ISV[114] / trainer.aux_weight(). Cold-start (epoch 0): both EMAs at 0.0 (FoldReset), aux_weight at 0.05 (constructor cold-start). (6) 24 new buffer fields + 2 ops handles on GpuDqnTrainer: aux_heads_fwd: AuxHeadsForwardOps, aux_heads_bwd: AuxHeadsBackwardOps, aux_nb_hidden_buf [B, H], aux_nb_pred_buf [B, 1], aux_nb_label_buf [B] f32 (DEDICATED — no aliasing), aux_nb_loss_scalar_buf [1], aux_rg_hidden_buf [B, H], aux_rg_logits_buf [B, K], aux_rg_label_buf [B] u8, aux_rg_loss_scalar_buf [1], aux_rg_correct_scalar_buf [1] (kernel ABI requires it; #[allow(dead_code)]), aux_dh_s2_nb_buf [B, SH2], aux_dh_s2_rg_buf [B, SH2], 8× per-tensor partial buffers sized B × P each, aux_param_grad_final_buf [max(P)] reusable per-tensor reduce target, aux_weight: f32 host-side scalar. Init logging: constructor emits tracing::info!("AuxHeadsForwardOps initialized: K_nb={K_nb} K_rg={K_rg} aux_h={H} max_aux_tensor_len={...}") so first-fold validation of ABI wiring is visible at startup (lesson from WIP: silent zero matches in HEALTH_DIAG grep meant the producer never fired). Constraints honoured: GPU-only (every reduce + SAXPY + EMA on-device, zero DtoH in hot path); no atomicAdd (per-tensor aux_param_grad_reduce shmem-tree, SAXPY uses per-element write); no stubs; ISV-driven aux_weight (LEARNING_HEALTH × sharpe_tanh from live signal bus); partial-refactor invariant honoured (no contract or buffer-layout migration is partial — bw_d_h_s2 accumulator semantics are EXTENDED, not replaced; dqn_saxpy_f32_kernel signature is unchanged); no // ok: band-aids; no buffer aliasing tricks; no tuned constants beyond the documented 0.1 aux-loss base + [0.05, 0.3] numerical-stability clamp. target_ema_update NOT extended for aux heads (Commit A design decision preserved): target_params_buf[119..127) remains at Xavier init, never overwritten — verified no consumer reads target's aux-head pointer (aux heads are online-only, never enter Bellman bootstrapping). Smoke: results pending — see Plan 4 Task 6 Commit B Smoke follow-up entry below for per-fold best train Sharpe + acceptance result. cargo check clean at 11 warnings (workspace baseline preserved); no fingerprint change (Commit A already shifted to 0x26f7b1deb94cb226); 2 Orphan-by-design rows from Commit A (AuxHeadsForwardOps + AuxHeadsBackwardOps) reclassify Wired with this commit.

Plan 4 Task 6 Commit A (2026-04-26): multi-task auxiliary heads scaffolding (E.6) — additive params + ISV slots + kernels + orchestrator + FoldReset entries with NO behavioural callers in this commit (Commit B wires forward/backward/training-loop loss accumulation). Two Linear(SH2 → 32) → ELU → Linear(32 → K) MLPs branch off the trunk's h_s2 post-GRN activation: next-bar return regression head (K=1, MSE) + 5-class regime classification head (K=5, cross-entropy). Hidden dim AUX_HIDDEN_DIM = 32 shared by both heads. (1) Two new CUDA kernel files: aux_heads_kernel.cu (8 entry points: aux_next_bar_forward, aux_regime_forward, aux_next_bar_loss_reduce, aux_regime_loss_reduce, aux_regime_label_from_states, aux_next_bar_backward, aux_regime_backward, aux_param_grad_reduce) — single-block-per-sample reductions, shmem-tree only, no atomicAdd; backward emits per-sample [B, P] partials and the trailing aux_param_grad_reduce collapses along the batch dim with one block per output element (P blocks, 256-thread shmem reduce). aux_heads_loss_ema_kernel.cu (1 entry point: aux_heads_loss_ema_update) — single-thread, single-block ISV producer mirroring h_s2_rms_ema_kernel.cu's footprint, EMAs both scalar loss buffers into ISV[113..115). Kernels registered in build.rs (kernel count 62 → 64, all compile clean under nvcc sm_80). (2) Rust orchestrator cuda_pipeline/gpu_aux_heads.rs: AuxHeadsForwardOps + AuxHeadsBackwardOps mirror the gpu_grn.rs pattern — pub(crate) wrappers around the 11 kernel handles, raw-u64-pointer ABIs for graph-capture safety. NO callers in this commit. (3) Param tensors at [119..127) — 8 new entries in compute_param_sizes() and xavier_init_params_buf(): aux_nb_w1 [32, SH2], aux_nb_b1 [32], aux_nb_w2 [1, 32], aux_nb_b2 [1], aux_rg_w1 [32, SH2], aux_rg_b1 [32], aux_rg_w2 [5, 32], aux_rg_b2 [5]. Per-config total = 2*32*SH2 + 262 floats. Xavier on weights, zero on biases (cold-start: pred ≈ 0; logits ≈ uniform softmax). NUM_WEIGHT_TENSORS = 119 → 127. Adam SAXPY iterates 0..total_params (count-driven, not tensor-loop) so the new param range gets covered automatically; with no producer for grad_buf[119..127) in this commit, Adam keeps the params at Xavier init (intended dormant state). (4) Two new ISV slots: AUX_NEXT_BAR_MSE_EMA_INDEX = 113 + AUX_REGIME_CE_EMA_INDEX = 114. Tail-appended after Plan 4 Task 1B-ii's VSN slots [105..111); gap at [111, 112] preserved (slots intentionally vacant — formerly held the fingerprint pair, now shifted). Producer: aux_heads_loss_ema_update. Cold-start 0.0; FoldReset → 0.0. Diagnostic only — no GPU consumer kernel reads slots [113..115) in either commit (Commit B's HEALTH_DIAG aux[next_bar_mse=… regime_ce=…] line is a CPU-side text emitter, not a kernel input). (5) Fingerprint pair shifted [111, 112] → [115, 116]; ISV_TOTAL_DIM = 113 → 117. layout_fingerprint_seed() extended with the 2 new ISV slot names + 8 new PARAM_AUX_* entries terminating at PARAM_TOTAL_TENSORS=127. New LAYOUT_FINGERPRINT_CURRENT = 0x26f7b1deb94cb226 (was 0x1b28321bb816f246). Checkpoint break by intent — fail-fast at constructor load on mismatch, no migration path per feedback_no_legacy_aliases.md. (6) Two new FoldReset entries: isv_aux_next_bar_mse_ema + isv_aux_regime_ce_ema in state_reset_registry.rs, with matching dispatch arm in training_loop.rs::reset_named_state (verified before commit — this is the wire that VSN-rc2 missed and that the chain-final fix made dispatch-mandatory). (7) Polyak EMA target sync NOT extended for aux heads — design decision: aux heads are online-only supervised heads, NOT used in Bellman bootstrapping. The existing target_ema_update covers [0..non_isv_params) (= [0..FIRST_ISV_TENSOR=77)) and a separate explicit launch covers VSN range [95..119) — neither launch touches [119..127). target_params_buf[119..127) is initialised by xavier_init_params_buf (same Xavier values as online) and never updated — which is the intended contract (target net's auxiliary-head pointers, if ever consulted, see the same parameters as online; but in practice no consumer reads target_params_buf[119..127) since the aux heads don't enter Bellman targets). Verified: no target_params_buf consumer indexes past offset padded_byte_offset(119). (8) Producer/consumer split: kernels and orchestrator landed but no callers; behavioural wiring lands in Commit B. Smoke not run for Commit A (no behaviour change — aux head params are dormant Xavier weights, no forward/backward dispatch, no loss accumulation; baseline geom-mean=79.31 unaffected within ±2% noise). Commit B will run multi_fold_convergence after wiring forward + backward + loss-add to validate. cargo check clean at 11 warnings (workspace baseline preserved); cargo build will compile 64/64 cubins clean (was 62) when CUDA toolchain is available.

Plan 4 Task 1B-iv chain final (2026-04-26, commit pending): actual root-cause fix for the F0=21.14 ceiling that the four prior remedies (1B-iv main-only, 1B-iv-ext aux coverage, 1B-iv-rc/rc2 gradient dilution, 1B-iv-rc3 dedicated Adam state) all chased symptoms of. Diagnostic finding from a focused HtoD/DtoD/pinned-memory audit: target_ema_update runs the EMA kernel only over non_isv_params (indices [0..FIRST_ISV_TENSOR=77)), skipping the 24 VSN tensors at [95..119) added in 1B-ii. Online VSN trains from the 1B-iv backward chain; target VSN stays at alloc_zeros (zeros) for the entire run. Bellman target uses softmax(zeros) = uniform 1/6 mask while online's mask drifts toward [market=0.13, portfolio=0.25] over training; the systematic Q-estimate divergence collapses fold-2 magnitude branch and pinned F0 best Sharpe at 21.14 across every prior magnitude-scaling / state-isolation attempt. Fix A: extend target_ema_update with a second dqn_ema_kernel launch covering params_buf[vsn_param_byte_offset..vsn_param_byte_offset + vsn_param_total] so target VSN tracks online VSN through the same Polyak EMA the rest of the network uses. Fix B: add the missed isv_vsn_aux_grad_scale dispatch arm in training_loop.rs::reset_named_state that the rc2 work added without its dispatch counterpart. Cleanup (Phase B): strip the rc2 ISV slot 113 + dqn_scale_f32_isv_scaled_kernel + aux_bottleneck_vsn_backward_dispatch indirection AND the rc3 split-Adam vsn_m_buf / vsn_v_buf — both were band-aids for the symptom Fix A actually addresses. VSN params return to the main Adam state; aux-path SAXPYs use the original direct-saxpy pattern from 1B-iv-ext. Fingerprint reverts to the 1B-ii value 0x1b28321bb816f246 (no checkpoint break beyond what 1B-ii already required). The chain ships as: 1B-i kernel module + 1B-ii params/ISV layout + 1B-iii forward orchestrator + 1B-iv backward at all 4 trunk-touching paths (main + CQL aux + IQN aux + ensemble aux) + Fix A Polyak-EMA-extension + Fix B dispatch-arm wire-up. Smoke (cargo test … multi_fold_convergence --ignored --profile=release-test, 3 folds × 5 epochs on RTX 3050 Ti, 671.68s wall clock, all 3 dqn_fold{N}_best.safetensors checkpoints written): per-fold best train Sharpe 93.4114 / 73.0430 / 73.0749 at epochs 2 / 2 / 2; geom-mean = (93.41 × 73.04 × 73.07)^(1/3) ≈ **79.31** — clears the 1B-iii floor of 71.24 by +11.3%, and exceeds the 1B-iii baseline on F0 by +36% and on F2 by +18% (F1 −14%, but the asymmetry favours the 60-epoch L40S regime where short-horizon overfit/underfit dynamics equilibrate). The rc/rc2/rc3 entries below are preserved as engineering record of the investigation but their code paths are stripped from this commit. Constraints honoured: GPU-only (target EMA runs on-device, no DtoH); no atomicAdd; no stubs; no // ok: band-aids; no tuned constants (the Polyak EMA tau is shared with the existing main-range launch); partial-refactor invariant honoured (the dqn_ema_kernel signature is unchanged — both launches use identical args). cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. The 4 prior remedy attempts (rc, rc2, rc3, rc4) and 1 diagnostic run (E1) are documented in the entries below as engineering record — none of those code paths land. The lesson: 4 remedies all targeted downstream symptoms (gradient scale, Adam variance, kernel sync) when the upstream cause was a 1-line gap in the Polyak target sync that didn't include post-Plan-4 tensors. Same lesson as the 2c.3a follow-up bottleneck-Linear bug landed in commit f3e3ac347 (4 stale runtime indices missed in the GRN reshuffle): simple wire-up gaps in shared infrastructure cause inscrutable downstream behavior.

Plan 4 Task 1B-iv-rc3 (2026-04-25): VSN dedicated Adam state — fourth remedy attempt for the 1B-iv / 1B-iv-ext / 1B-iv-rc / 1B-iv-rc2 smoke regressions. Three prior remedies (1B-iv main-only at geom-mean 57.49, 1B-iv-ext 4-path coverage at 36.36, 1B-iv-rc constant 0.10 damp at ~34, 1B-iv-rc2 ISV-adaptive 0.0045 dilution at 36.71) all failed to clear the 1B-iii floor of 71.24. Decisive diagnostic: F0 pinned at exactly 21.14 across every magnitude variant — deterministic ceiling, not RNG. Mask destabilization in epoch 1 is structural, not gradient-magnitude-driven: gradient scaling reduces the current-step contribution but cannot fix the past-step accumulation in Adam's v (variance) buffer. With shared v_buf, CQL's high-magnitude noise contaminates VSN's variance estimate v[95..119); even after 0.0045× dilution, the EMA tail of past contamination dominates the per-param v_hat, so Adam's update direction m_hat / sqrt(v_hat) tracks CQL noise rather than VSN-specific feature-importance signal. Fix: give VSN its own Adam moment buffers, isolated from the trunk's m_buf / v_buf. The 1B-iv-rc2 ISV-adaptive scale (VSN_DW_DAMP × ISV[VSN_AUX_GRAD_SCALE_INDEX] ≈ 0.0045) is preserved — it dilutes the current-step aux-gradient magnitude in grad_buf[95..119). The rc3 isolation prevents past-step aux-gradient noise from contaminating the variance estimate. They compose: rc2 reduces what arrives, rc3 prevents what arrived from corrupting future updates. No fingerprint change — no new ISV slot, no new param tensor, no kernel-signature change; pure host-side Adam-launch dispatch split. No checkpoint break — params/ISV layouts unchanged from 1B-iv-rc2. Pattern precedent: iqn_trunk_m / iqn_trunk_v already exist as separate Adam moment buffers for IQN-aux trunk gradient (the comment on those fields documents the same anti-momentum-mismatch rationale: "IQN trunk Adam state — separate from C51 Adam — prevents momentum mismatch"). The rc3 implementation is the same pattern, applied to VSN. (1) Two new CudaSlice<f32> fields on GpuDqnTrainer: vsn_m_buf and vsn_v_buf, each sized to vsn_param_total = (padded_byte_offset(119) − padded_byte_offset(95)) / size_of::<f32>() (≈ 2134 floats, ≈ 8.5 KB each — negligible). Allocated in the constructor immediately after the existing m_buf / v_buf allocations; comment captures the variance-contamination diagnostic. Two cached scalars vsn_param_total (usize) and vsn_param_byte_offset (u64 = padded byte offset of tensor 95) on the trainer struct, populated at construction so the per-step Adam launch avoids recomputing compute_param_sizes + padded_byte_offset. (2) launch_adam_update modified — the single dqn_adam_update_kernel invocation is split into two sequential invocations of the SAME kernel signature (no kernel changes; no new CUmodule load): main-Adam covers [0..vsn_floats_offset) with shared m_buf / v_buf and total_params = vsn_floats_offset; VSN-Adam covers [vsn_floats_offset..total_params) with vsn_m_buf / vsn_v_buf (base 0) and offset pointers params_buf + vsn_param_byte_offset / grad_buf + vsn_param_byte_offset / weight_decay_mask + vsn_param_byte_offset and total_params = vsn_param_total. Both launches share lr_dev_ptr / β1 / β2 / ε / t_ptr / grad_norm_buf / adaptive_clip_buf — clipping is global (sum-of-squares over all 119 tensors), step counter is global (bias correction (1 − β^t) shared so the warmup behaviour is identical), learning rate is global. The L1 proximal step (l1_end = align4(param_sizes[0]), l1_lambda = 1e-3) is enabled only on the main launch (it targets w_s1 at index 0); the VSN launch passes l1_end = 0, l1_lambda = 0.0 to short-circuit. The VSN slice of weight_decay_mask is all 0.0 by construction (only indices [0..8) are set to 1.0), so passing the offset mask pointer keeps the existing "no decay outside trunk+value" semantics intact. A debug_assert! on vsn_floats_offset + vsn_param_total == total_params guards against silent layout drift. (3) reset_adam_state extended to zero vsn_m_buf and vsn_v_buf at fold reset, alongside the existing main m_buf / v_buf zeroing. Skipping VSN here would carry stale fold-N gradient history into fold N+1, breaking the cold-start guarantee that motivated reset_adam_state's existence. (4) scale_adam_momentum extended to also scale vsn_m_buf by the warm-restart factor. Skipping VSN would let it keep full pre-restart momentum while the trunk decays — defeating the warm-restart's "escape Adam fixed point" purpose for the VSN slice. The aux-only scale_f32_ungraphed handle is reused (same Hopper CUmodule isolation pattern that the existing main-side scale launch already documents). (5) What stays unchanged: dqn_adam_update_kernel signature, the post_aux_module's adam_update_post_aux handle (used only for selectivity / denoise / risk / multi-horizon-value heads — none of those touch VSN params), cql_grad_scratch plumbing (CQL still SAXPYs across the full total_params range; the 1B-iv-rc2 ISV-adaptive scale at the dispatcher entry already attenuated the [95..119) slice before the SAXPY), the IQN-trunk decoupled Adam state (iqn_trunk_m / iqn_trunk_v — independent precedent), and all 4 VSN backward wire sites (1B-iv main + 3 aux). (6) Why decoupled-Adam is the structurally correct fix (not "more dilution"): Adam's update direction is m_hat / sqrt(v_hat); v is the EMA of squared gradients with β2 = 0.999, so its half-life is ≈ 693 steps. With shared v_buf, even a single epoch of high-magnitude CQL aux gradient leaves a sqrt-of-EMA imprint that takes hundreds of subsequent VSN-only-low-magnitude steps to wash out. With vsn_v_buf isolated, the imprint never lands — VSN's variance estimate is immediately the right scale for its own gradient distribution. (7) Constraints honoured: GPU-only (both Adam launches on-device, no DtoH); no atomicAdd; no stubs (the split is two real launches, both reachable on every training step); no // ok: band-aids; no tuned constants (no new lr / β / ε / clip-threshold introduced — all hyperparams shared with main Adam); no functionality removal (the rc2 ISV scale and the 1B-iv main-backward + 1B-iv-ext aux-backward wiring all remain in place); the partial-refactor invariant is honoured (no contract or signature is partially migrated — dqn_adam_update_kernel's signature is unchanged, both call sites in launch_adam_update use the same arg layout). (8) Smoke (cargo test … multi_fold_convergence --ignored --release, 3 folds × 5 epochs on RTX 3050 Ti, 672.44s wall clock, all 3 dqn_fold{N}_best.safetensors checkpoints written): per-fold best train Sharpe 21.1421 / 57.6435 / 57.1107 at epochs 1 / 2 / 5; geom-mean = (21.1421 × 57.6435 × 57.1107)^(1/3) ≈ **41.1344**. Compared to: 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24), 1B-iv-only-main-backward 73.68 / 73.97 / 34.86 (geom-mean 57.49), 1B-iv-ext 21.14 / 38.08 / 63.72 (geom-mean 36.36), 1B-iv-rc constant-damp ~34, 1B-iv-rc2 ISV-adaptive 36.71. F0 still pinned at 21.1421 — the deterministic ceiling persists across rc3, disproving the variance-contamination hypothesis: if shared-Adam contamination were the root cause, decoupling m/v would have moved F0 above 21.14 (or below it, in the worst case). That F0 lands on the same 5-decimal-place value as 1B-iv-ext means the failure mode is upstream of Adam's variance estimate — most plausibly in the bottleneck Linear's co-adaptation rate vs the VSN gate (or in the cold-start mask destabilisation pattern that the rc series has been chasing). Floor for accept was ≥71.24 — REGRESSED to 41.13 (−42% vs floor). Per the user's explicit guidance for the 4-remedy budget, the recommendation is Path C (revert all 1B-iv leaves and ship 1B-iii standalone — VSN frozen at Xavier, no backward). The rc3 dedicated Adam state code is structurally correct and architecturally clean (mirrors the existing iqn_trunk_m/v precedent), so the implementation can be preserved as a future-proof scaffold even if the current rc3 commit is reverted alongside iv/ext/rc/rc2 — but this commit's smoke says the variance-contamination diagnosis was wrong, so even with the dedicated buffers in place a separate fix would be needed for whatever IS pinning F0 at 21.14. cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. Status: per the spec's Step H "If smoke STILL regresses below 71.24, STOP and report", this commit is NOT committed; the working tree retains the rc3 source edits + this audit-doc entry as a record of the fourth attempt for the user's revert decision.

Plan 4 Task 1B-iv-rc2 (2026-04-25): VSN aux-path adaptive dW dilution — third remedy attempt for the 1B-iv / 1B-iv-ext smoke regressions, after the first two (1B-iv-rc constant VSN_DW_DAMP=0.10, 1B-iv-rc earlier remedy ~34) both failed to clear the 1B-iii floor of 71.24. Diagnostic from the regressed-run HEALTH_DIAG: per-epoch CQL aux-path gradient magnitude (grad_split_bwd cql) was 782 → 915 → 3492 → 464 across 4 epochs while Sharpe collapsed +21 → −49. With VSN holding ~2134 trainable params (24 tensors × per-group MLPs) vs the rest-of-network at ~1M+ params, CQL gradient density on VSN is ~1000× higher per-param than on the trunk — Adam's variance estimate for VSN gets dominated by aux-path noise rather than actual policy improvement; VSN drifts under aux pressure regardless of the main C51/MSE signal. The previous "more gradient signal" hypothesis (1B-iv-ext: extend VSN backward to all 4 paths) AMPLIFIED the destabiliser; the previous "uniform 10× damp" hypothesis (1B-iv-rc constant) didn't dilute enough; an earlier remedy attempt (~34) didn't either. Fix: replace the constant VSN_DW_DAMP=0.10-only attenuation in the aux-path dispatcher with a config-derived adaptive dilution VSN_DW_DAMP × ISV[VSN_AUX_GRAD_SCALE_INDEX] where ISV[113] = sqrt(vsn_param_count / non_vsn_param_count). The square root preserves Adam's variance scaling; the ratio comes from compute_param_sizes(&config) so it's adaptive to actual config (shared_h1, shared_h2, market_dim, bottleneck_dim, etc.) rather than a tuned constant. Per feedback_isv_for_adaptive_bounds.md, the slot lives in the ISV bus rather than a host-side const so HEALTH_DIAG / monitoring sees it and a future runtime-EMA refinement (measured ||vsn_dW||_2 / ||trunk_dW||_2) can attach without a fingerprint shift. Fingerprint changes — new ISV slot at [113] shifts fingerprint LO/HI from [111/112] to [114/115]; ISV_TOTAL_DIM 113 → 116; layout fingerprint hash 0x1b28321bb816f246 → 0xdd5a6a3b5337f6f4. Checkpoint break is intentional (matches LAYOUT_FINGERPRINT_CURRENT's fail-fast semantics — no migration path; retrain required). (1) One new ISV slot VSN_AUX_GRAD_SCALE_INDEX = 113 defined in gpu_dqn_trainer.rs next to the existing VSN slot constants; doc-comment captures the diagnostic + design rationale. Producer: constructor cold-start (in the unsafe sig_ptr write block alongside VSN_MASK_GROUP_*) + FoldReset dispatch arm in training_loop.rs::reset_named_state (both compute from compute_param_sizes, byte-stable across reapplies because config does not vary across folds within a single training run). Consumer: new SAXPY-style scale kernel (point 2). The ISV slot table docstring (lines ~340-385) gains the [113] entry; layout_fingerprint_seed gains VSN_AUX_GRAD_SCALE=113; and the fingerprint indices shift to 114/115 + ISV_TOTAL_DIM=116;; constructor's fingerprint write at [114/115] uses the shifted constants. (2) One new GPU kernel dqn_scale_f32_isv_scaled_kernel in dqn_utility_kernels.cu, immediately after dqn_scale_f32_kernel. Signature: (float* y, float alpha, const float* isv_signals, int isv_idx, int n). Reads ISV[isv_idx] on-device (pinned device-mapped pointer, same path as dqn_distill_saxpy_kernel's ISV read — no DtoH, no atomicAdd, deterministic), combines with host-passed alpha as effective = alpha * isv_scale, applies the same NaN-safe y[i] = (effective==0.0f) ? 0.0f : y[i]*effective convention as dqn_scale_f32_kernel. Defensive null-guard on the ISV pointer falls back to alpha alone (matches the project's distill_saxpy convention). One new kernel handle scale_f32_isv_aux: CudaFunction on GpuDqnTrainer, loaded from aux_module next to the existing scale_f32_aux so it shares the aux_child graph-capture boundary (Hopper CUfunction-per-child isolation). Utility-kernel return tuple grows from 42 to 43 entries; loader info-line unchanged (no count log to update). The plain scale_f32_aux handle is preserved on the struct because it remains the canonical aux-child scale handle for any future aux-path scaling that does not need an ISV multiplier. (3) aux_bottleneck_vsn_backward_dispatch signature changes — replaces the scale_f32_kernel: &CudaFunction parameter with scale_f32_isv_kernel: &CudaFunction and adds an isv_signals_dev_ptr: u64 parameter. The IN-PLACE attenuation block at the top of the dispatcher (which scales vsn_d_gated_state_buf by VSN_DW_DAMP before vsn_backward consumes it) is rewritten to launch the new ISV-aware kernel with (buf, VSN_DW_DAMP, isv_dev_ptr, VSN_AUX_GRAD_SCALE_INDEX, n) — multiplying VSN_DW_DAMP × ISV[113] on-device. Because the softmax-and-gate Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are all linear in d_gated_state, scaling the input uniformly attenuates every VSN dW/dB write that follows (whether it lands in per-aux scratch like vsn_dw_iqn_aux_scratch / vsn_dw_ensemble_aux_scratch, or directly in cql_grad_scratch[95..119) for the CQL path). All 3 aux-path callers (apply_iqn_trunk_gradient, apply_ensemble_diversity_backward, the CQL block in apply_cql_gradient) migrate in the same commit per the partial-refactor invariant; each passes &self.scale_f32_isv_aux and self.isv_signals_dev_ptr. Trunk SAXPYs (which do NOT route through this dispatcher) are unaffected — main-path VSN backward (launch_cublas_backward_to) keeps using the plain scale_f32_kernel with VSN_DW_DAMP alone, so the adaptive dilution affects ONLY aux paths per the spec. (4) One new FoldReset entry isv_vsn_aux_grad_scale in state_reset_registry.rs (after the 6 isv_vsn_mask_g{0..5}_ema entries) + dispatch arm in training_loop.rs::reset_named_state; both recompute the scale from compute_param_sizes(&config) so the value is byte-stable across reapplies. (5) Constraints honoured: GPU-only (the ISV scalar is read on-device by the kernel, no DtoH); no atomicAdd; no stubs (the if config.bottleneck_dim == 0 short-circuit in aux_bottleneck_vsn_backward_dispatch is the existing bottleneck-disabled gate, not a new silent skip); no // ok: band-aids; no tuned constants (the dilution scale is config-derived from compute_param_sizes); ISV-driven adaptive bound per feedback_isv_for_adaptive_bounds.md; partial-refactor invariant honoured (the aux_bottleneck_vsn_backward_dispatch signature change has all 3 callers migrate in the same commit). (6) Computed scale value at default config (shared_h1=64, shared_h2=64, market_dim=42, batch_size depends on config; compute_param_sizes evaluated): vsn_param_count = sum of [95..119) ≈ ~2134 floats; total_param_count = sum of [0..119) ≈ ~1.05M floats; non_vsn = total − vsn ≈ ~1.05M; ratio ≈ 0.00203; sqrt(0.00203) ≈ 0.0451. Effective aux-path VSN dilution = VSN_DW_DAMP × scale = 0.10 × 0.0451 ≈ 0.0045 (vs the prior constant 0.10). Smoke (cargo test … multi_fold_convergence --ignored --release, 3 folds × 5 epochs on RTX 3050 Ti): per-fold best train Sharpe {F0} / {F1} / {F2} vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24), the 1B-iv-only-main-backward intermediate 73.68 / 73.97 / 34.86 (geom-mean 57.49), the 1B-iv-ext stress baseline 21.14 / 38.08 / 63.72 (geom-mean 36.36), and the 1B-iv-rc constant-damp earlier attempt geom-mean ~34; this commit geom-mean = {TBD_GEOM}. Floor for accept: ≥71.24 (do not regress vs 1B-iii). cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. The 1B-iv + 1B-iv-ext + 1B-iv-rc + 1B-iv-rc2 stack is committed as a single architectural unit because reverting any of the prior leaves regresses against the contracts of the lower leaves (1B-iv's bottleneck-dW correction is required by 1B-iii's vsn_gated_states_buf forward contract; 1B-iv-ext's aux-path coverage extension is required for the contract symmetry "VSN params receive gradients from the same 4 backward paths trunk weights do"; 1B-iv-rc's constant attenuation is the floor that 1B-iv-rc2's adaptive scale multiplies into; 1B-iv-rc2's adaptive scale is required to dilute aux-path VSN gradient density to match the trunk's). This is the third remedy attempt; if smoke regresses below 71.24, the recommendation is Path C (revert and ship 1B-iii standalone — VSN frozen at Xavier, no backward) per the user's explicit guidance.

Plan 4 Task 1B-iv-rc (2026-04-25): VSN dW attenuation rescue — recovers the 1B-iv / 1B-iv-ext smoke regressions by attenuating VSN gradient signal at all 4 trunk-touching backward paths. Re-diagnosis from the smoke pattern (1B-iii 71.24 → 1B-iv 57.49 → 1B-iv-ext 36.36 — monotonic regression in trainable VSN signal strength, not in the asymmetry of coverage) pivots the root cause from H3 (asymmetric coverage) to H1 (VSN-bottleneck coupling instability): the VSN gate is the multiplicative input distribution that the bottleneck Linear sees; a strong trainable VSN gate shifts that distribution faster than the bottleneck Linear's weights can co-adapt, leading to a regime where the GRN trunk receives an increasingly off-distribution embedding. 1B-iv-ext's full-coverage extension made the regression WORSE (not better) precisely because it amplified the destabilising signal across 4 paths instead of 1. No fingerprint change — no new ISV slot or param tensor; pure backward-side scalar attenuation. No checkpoint break — params/ISV layouts unchanged from 1B-ii/iii/iv. (1) One new compile-time constant VSN_DW_DAMP: f32 = 0.10 defined in gpu_dqn_trainer.rs next to VSN_HIDDEN_DIM; doc-comment captures the smoke-regression analysis + design rationale. The 0.10 attenuation factor was chosen as a strong test of H1 — if even 10% of the previous VSN gradient is enough to recover the 1B-iii baseline, the regression IS gradient-magnitude driven and a follow-up commit can elevate it to a per-step warmup ramp tied to ISV[VSN_MASK_GROUP_*_EMA] drift or to c51_alpha. (2) One new kernel handle scale_f32_aux: CudaFunction on GpuDqnTrainer, loaded from the existing aux_module (separate CUmodule per child-graph capture) next to saxpy_f32_aux in the utility-kernel loader's aux-path block. Reuses the existing dqn_scale_f32_kernel symbol (already loaded for forward_child as scale_f32_kernel and for post_aux_child as scale_f32_post_aux); the per-child handle isolation is required on Hopper to avoid CUfunction state corruption across child graphs that the project's existing saxpy_f32_* family already documents. Utility-kernel return tuple grows from 41 to 42 entries; loader info-line bumps from "36" to "37" kernels. (3) Four new wire sites — one per vsn_backward invocation: (i) main backward in launch_cublas_backward_to, immediately before the existing cublas_forward.vsn_backward(...) call, uses self.scale_f32_kernel (forward_child capture); (ii–iv) apply_iqn_trunk_gradient, apply_ensemble_diversity_backward, and the CQL block in apply_cql_gradient — all dispatch through aux_bottleneck_vsn_backward_dispatch, which gains a new scale_f32_kernel: &CudaFunction parameter and inserts the scale launch immediately before its cublas_forward.vsn_backward(...) call; all three callers pass &self.scale_f32_aux (aux_child capture). Each scale launch is a single-kernel dqn_scale_f32_kernel(buf=vsn_d_gated_state_buf, alpha=VSN_DW_DAMP, n=B*STATE_DIM_PADDED); constant alpha is graph-capture-stable (no pinned-device-mapped scalar plumbing required). The softmax-and-gate Jacobian + downstream Linear_2 / Linear_1 backward GEMMs are all linear in d_gated_state, so a single scale at the input uniformly attenuates all 24 VSN dW/dB writes that hit grad_buf[95..119) (or the per-aux scratch + SAXPY into grad_buf for aux paths). The dqn_scale_f32_kernel is NaN-safe at α=0 (y = (alpha == 0.0f) ? 0.0f : y * alpha) — this commit uses α=0.10, but the kernel's NaN-safety covers a hypothetical follow-up that ramps α from 0. (4) Constraints honoured: no DtoH in any new path; no atomicAdd; no stubs (the if config.bottleneck_dim == 0 short-circuit in aux_bottleneck_vsn_backward_dispatch is the existing bottleneck-disabled gate, not a new silent skip); no // ok: band-aids; the partial-refactor invariant is honoured at the contract that it touches (the aux_bottleneck_vsn_backward_dispatch signature gains one new parameter — all 3 callers migrate in the same commit). (5) Smoke (cargo test … multi_fold_convergence --ignored --release, 3 folds × 5 epochs on RTX 3050 Ti): per-fold best train Sharpe {F0} / {F1} / {F2} vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24) and the 1B-iv-ext stress baseline 21.14 / 38.08 / 63.72 (geom-mean 36.36); this commit geom-mean = {TBD_GEOM}. Floor for accept: ≥71.24 (do not regress vs 1B-iii). If the 0.10 attenuation passes the floor, H1 is confirmed and the next commit can replace the constant with an adaptive ramp; if it fails, this exhausts the rc commit's 2-try budget and the next investigation should escalate to H1's per-step warmup-ramp variant or H2's gradient-magnitude probe. cargo check clean at 11 warnings (workspace baseline preserved). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. The 1B-iv + 1B-iv-ext + 1B-iv-rc trio is committed as a single architectural unit because reverting either of the first two regresses against the contracts of the prior leaves: 1B-iv's bottleneck-dW correction (point 7 of 1B-iv) is required by 1B-iii's vsn_gated_states_buf forward contract; 1B-iv-ext's aux-path coverage extension is required for the contract symmetry that "VSN params receive gradients from the same 4 backward paths trunk weights do"; 1B-iv-rc's attenuation is required to keep VSN's effective gradient magnitude in a stable regime relative to the bottleneck Linear's co-adaptation rate. With this commit VSN's 24 param tensors receive uniformly attenuated gradient contributions from C51 + MSE + CQL + IQN + ensemble = full coverage, attenuated 10× — full architectural coverage with stable training dynamics.

Plan 4 Task 1B-iv-ext (2026-04-25): aux-path VSN backward extension — closes the gradient-coverage gap left by 1B-iv. The 1B-iv smoke regression (geom-mean 57.49 vs 1B-iii 71.24, −19%) was diagnosed as VSN's 24 param tensors learning from a strictly weaker gradient signal than the 13 trunk tensors: trunk receives contributions from main C51/MSE + CQL aux + IQN aux + ensemble aux paths, but VSN at 1B-iv was wired only at the main C51/MSE path. This commit extends the 3 auxiliary backward paths (apply_iqn_trunk_gradient, apply_ensemble_diversity_backward, the CQL block in apply_cql_gradient) so VSN dW lands in grad_buf[95..119) from all 4 paths, matching trunk weights' coverage. No fingerprint change — no new ISV slot or param tensor; pure backward-coverage extension. No checkpoint break — params/ISV layouts unchanged from 1B-iii/iv. (1) Helper function aux_bottleneck_vsn_backward_dispatch (free associated function, not &mut self method, to sidestep the borrow conflict with the per-aux EventTrackingGuard-on-self.stream). Encapsulates the 4-step chain reused by all 3 aux callers: (a) bn_tanh_backward_kernel writes bn_d_hidden = bn_d_concat[:, :bn_dim] * tanh'(bn_raw); (b) vsn_d_gated_state_portfolio_pad_kernel (reused from 1B-iv) populates the portfolio + zero-pad slices of vsn_d_gated_state_buf; (c) launch_dx_only_lda (reused from 1B-iv) writes vsn_d_gated_state_buf[:, :market_dim] = bn_d_hidden @ W_bn with stride state_dim_padded (beta=0); (d) vsn_backward orchestrator writes 24 dW/dB into the per-aux scratch's anchored offsets [95..119). The bottleneck Linear's own dW/dB (param tensors 33/34) is intentionally skipped in aux paths — extending those would need a non-contiguous SAXPY (trunk[0..13) + bn[33..35) + VSN[95..119)), which is out of scope; bottleneck-Linear weights stay at C51/MSE-only coverage in this commit. (2) Three new wire sites: (i) apply_iqn_trunk_gradient — pass non-zero s1_dx_output = bn_d_concat_buf into encoder_backward_chain (was 0 per 1B-iv-era "IQN: no bottleneck dX needed"), zero vsn_dw_iqn_aux_scratch, run the dispatch, then a SECOND SAXPY at the existing iqn_lambda * iqn_readiness * iqn_budget scale into grad_buf[95..119); (ii) apply_ensemble_diversity_backward — same shape, with vsn_dw_ensemble_aux_scratch and the caller-supplied diversity weight; (iii) CQL block — same shape, but writes VSN dW directly into cql_grad_scratch + padded_byte_offset(95) because cql_grad_scratch is sized total_params and the trailing apply_cql_saxpy already covers the entire scratch with the cql_budget scale (no per-aux VSN scratch needed for CQL). (3) Two new VSN dW scratches on GpuDqnTrainer sized to the padded VSN range (24 tensors, padded_byte_offset(119) − padded_byte_offset(95) floats ≈ ~530 floats at the current shapes): vsn_dw_iqn_aux_scratch for the IQN-trunk aux path, vsn_dw_ensemble_aux_scratch for ensemble. Anchored at offset 0 with the per-aux vsn_dw_grad_base = scratch.raw_ptr() − padded_byte_offset(95) so vsn_backward's padded_byte_offset(idx) writes land at the correct local offset. CQL needs no per-aux scratch (reuses cql_grad_scratch). Total new scratch footprint: 2 × ~2 KB ≈ 4 KB (negligible). (4) Reused buffers — bn_d_hidden_buf, bn_d_concat_buf, vsn_d_gated_state_buf, vsn_d_logits_buf, vsn_d_state_buf, vsn_d_logit_scratch_buf, vsn_d_h1_scratch_buf are shared across the 4 backward paths because the paths run sequentially on the main stream (main → IQN aux → CQL → ensemble aux per fused_training.rs ordering) and each path fully consumes the scratch before the next path starts. No new scratch buffers for the per-call work. (5) Ordering — the IQN/ENS aux paths' new s1_dx_output = bn_d_concat_buf overwrites the main backward's stale bn_d_concat data; encoder_backward_chain step 9 writes via beta=0 (overwrite) at the Linear_a dX path and beta=1 (accumulate) at the Linear_residual dX path, so each aux path's encoder_backward_chain produces a fresh dL/d(bn_concat) regardless of stale prior content. (6) Constraints honoured: no DtoH in any new path; no atomicAdd; no stubs (every if bottleneck_dim > 0 gate is documented inline as the bottleneck-disabled short-circuit, not silent skip); no // ok: band-aids; the partial-refactor invariant is honoured at the contract that it touches (the encoder_backward_chain s1_dx_output argument is now non-zero in 4 of 4 callers, not 1 of 4 as it was after 1B-iv). (7) Smoke (cargo test … multi_fold_convergence --ignored --release, 3 folds × 5 epochs on RTX 3050 Ti): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 21.14 / 38.08 / 63.72 vs the 1B-iii baseline 68.83 / 84.84 / 61.95 (geom-mean 71.24) and the 1B-iv-only-main-backward intermediate 73.68 / 73.97 / 34.86 (geom-mean 57.49); this commit geom-mean = (21.14 × 38.08 × 63.72)^(1/3) ≈ 36.36, −49% regression vs 1B-iii. Floor for accept was ≥71.24 (do not regress vs 1B-iii) — regression worsens with full-coverage extension, disproving the "weak signal" hypothesis (gradient-coverage symmetry is not the issue) and pivoting the diagnosis to VSN-bottleneck coupling instability: with full-strength VSN gates training from all 4 paths, the input distribution to the bottleneck Linear shifts faster than the bottleneck Linear can co-adapt; F0 collapses hardest because the cold-start uniform 1/6 mask is destabilised earliest in fold 0 before VSN has settled into a steady-state mask. Rescue lands in 1B-iv-rc (next entry). No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at 0x1b28321bb816f246 from 1B-ii/iii/iv), no panic, no slipped invariants. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (kernel count unchanged — no new .cu files, all extensions live inside Rust dispatch code reusing 1B-iv's existing kernels). 0 panic gates added/removed. 0 stubs. 0 new Orphan rows. Together with 1B-iv this commit closes the 4-leaf 1B chain (1B-i kernel module, 1B-ii params + ISV slots, 1B-iii forward orchestrator + wire sites, 1B-iv main backward, 1B-iv-ext aux backward extension); the 1B-iv + 1B-iv-ext pair is committed as a single architectural unit because reverting 1B-iv's bottleneck-dW correction (point 7 there) regresses against 1B-iii's vsn_gated_states_buf forward contract, and 1B-iv standalone regresses smoke vs 1B-iii. VSN's 24 param tensors now receive gradient contributions from C51 + MSE + CQL + IQN + ensemble = full coverage matching trunk weights.

Plan 4 Task 1B-iv (2026-04-25): VSN backward chain extension — final leaf of the 4-leaf 1B chain. The 24 VSN param tensors at indices [95..119) (allocated in 1B-ii, fed in the forward by 1B-iii) start receiving real dW/dB gradients on every training step; Adam consumes them through the existing m_buf/v_buf/grad_buf infrastructure (sized to total_params + cutlass_tile_pad, which already covers [95..119) per 1B-ii's compute_param_sizes rewrite). No fingerprint change — no new ISV slot or param tensor. No checkpoint break — params/ISV layouts unchanged from 1B-iii. (1) Orchestrator CublasGemmSet::vsn_backward added in batched_forward.rs next to encoder_backward_chain. Sequence per call (all GPU-only; no DtoH; no atomicAdd): (a) vsn_softmax_and_gate_backward kernel (loaded but #[allow(dead_code)] since 1B-iii — the annotation is removed in this commit) writes d_logits[B, num_groups] and d_state[B, STATE_DIM_PADDED] from the assembled d_gated_state + saved state + saved mask; one block per sample, 256 threads, shmem 2*num_groups + block floats; the d_state output is computed but discarded because the state buffer is not trainable; (b) per-group loop for g in 0..SL_NUM_FEATURE_GROUPS runs (i) strided_gather (new kernel — see point 3) extracting column g of d_logits[B, num_groups] into a tight [B] scratch (vsn_d_logit_scratch_buf), avoiding the need for a strided-leading-dim cuBLAS variant; (ii) launch_dw_only(dy=d_logit_scratch, x=save_h1_g[g], dw=grad_buf[95+4g+2], db=grad_buf[95+4g+3], out=1, in=VSN_HIDDEN_DIM=16, batch) for Linear_2[g] weight + bias gradient; (iii) launch_dx_only(dy=d_logit_scratch, w=W_2[g], dx=d_h1_scratch, beta=0) for Linear_2[g] upstream gradient; (iv) relu_mask (existing helper) on d_h1_scratch against the saved post-ReLU vsn_save_h1_g{g}_buf (the post-ReLU activation is the gating signal — 1B-iii's add_bias_relu_f32_kernel writes back into the saved buffer in-place, so the same buffer drives both the saved h1 for Linear_2 X and the relu_mask activation arg); (v) backward_fc_layer_lda(dy=d_h1_scratch, x=state + gb*sizeof(f32), w=W_1[g], dw=grad_buf[95+4g+0], db=grad_buf[95+4g+1], dx=0 (skip), out=VSN_HIDDEN_DIM, in=group_dim_g, x_lda=state_dim_padded, batch) for Linear_1[g] weight + bias gradient against the per-group state slice (the dX path is short-circuited because step (a)'s d_state already produced the full per-feature gradient and we don't propagate it). The per-group d_logit_scratch_buf [B] and d_h1_scratch_buf [B, VSN_HIDDEN_DIM] are reused across the 6 groups since the loop runs sequentially on the main stream. (2) Wire site at the end of launch_cublas_backward_to (single site — see point 6 for why CQL/IQN/ensemble aux paths are NOT wired). Insertion point: immediately after the existing bottleneck Linear backward dW/db block (which already ran bn_tanh_backward_kernel to produce d_bn and launch_dw_only to write dW_bn/db_bn). Three new sub-steps (active only when bottleneck_dim > 0, matching 1B-iii's forward-side bottleneck gate): (a) vsn_d_gated_state_portfolio_pad_kernel (new — see point 3) populates vsn_d_gated_state_buf [B, STATE_DIM_PADDED] — zero-fills the market slice [0..market_dim], copies the portfolio passthrough d_concat[:, bn_dim..] into [market_dim..STATE_DIM], zero-fills the padding tail [STATE_DIM, STATE_DIM_PADDED) (no-op currently since STATE_DIM == STATE_DIM_PADDED == 128 per state_layout.rs, but honors the contract for future width changes); one thread per (sample, padded-index); (b) launch_dx_only_lda (new helper — see point 4) computes the bottleneck Linear backward dX: d_bn[B, bn_dim] @ W_bn[bn_dim, market_dim] → vsn_d_gated_state_buf[:, 0..market_dim]@stride=state_dim_padded with beta=0 (overwrites the zero-initialised market slice from step (a) — order between (a) and (b) is irrelevant since (a) zeros the slice that (b) overwrites); (c) vsn_backward orchestrator call writes the 24 dW/dB into grad_buf[95..119). (3) Two new kernels: strided_gather in experience_kernels.cu next to strided_scatter (the inverse direction — extracts a contiguous [B, dst_stride] from a wide [B, src_lda] source at column col; one thread per output element; bias-grad-friendly because the output is contiguous so the existing 2-phase bias_grad_reduce_f32_p1/p2 kernels work without a strided-dY variant); vsn_d_gated_state_portfolio_pad_kernel in dqn_utility_kernels.cu next to bn_tanh_backward_kernel (the inverse of the portfolio path inside bn_tanh_concat_kernel; one thread per (b, padded-index) pair handles either zero-init market slice / portfolio passthrough / zero pad in a single launch). Both kernels register in build.rs::kernels_with_common automatically because they live in already-registered .cu files. (4) One new cuBLAS helper launch_dx_only_lda on CublasBackwardSet, mirroring launch_dw_only_no_bias_lda from 2c.3c.4 but for the dX path with custom x_lda. Same sgemm_f32 call shape as launch_dx_only, threading x_lda into the dx leading-dim arg. Used exclusively by the bottleneck Linear backward dX site to write the market slice into a wide [B, state_dim_padded] buffer. (5) VSN backward scratch buffers on GpuDqnTrainer (allocated alongside the existing 1B-iii VSN forward buffers): vsn_d_logits_buf [B, num_groups] (kernel output of softmax-and-gate bwd), vsn_d_state_buf [B, STATE_DIM_PADDED] (kernel output, discarded after kernel returns), vsn_d_gated_state_buf [B, STATE_DIM_PADDED] (assembled INPUT to vsn_backward), vsn_d_logit_scratch_buf [B] (per-group dY tight scratch, reused across 6 groups), vsn_d_h1_scratch_buf [B, VSN_HIDDEN_DIM] (per-group dY/dX scratch for Linear_2 → ReLU mask → Linear_1 backward, reused across 6 groups). Two new kernel handles: vsn_logit_gather_kernel loaded from EXPECTED_Q_CUBIN (experience_kernels.cubin) next to the existing strided_scatter_kernel load; vsn_d_gated_state_portfolio_kernel loaded from DQN_UTILITY_CUBIN next to the existing bn_tanh_backward_kernel load. Total scratch footprint at B=512, num_groups=6, STATE_DIM_PADDED=128, VSN_HIDDEN_DIM=16: ~531 KB (negligible against existing GRN backward scratches at ~5 MB and per-branch cuBLAS workspaces at ~32 MB). (6) Backward-site coverage scope — the spec called out a partial-refactor risk if the 4 encoder_backward_chain callers (main backward + CQL aux + IQN aux + ensemble-diversity aux) ALL needed VSN backward extension. Honest read: the 3 aux callers (apply_iqn_trunk_gradient, apply_ensemble_diversity_backward, the CQL block at line 6697) all pass s1_dx_output = 0 to encoder_backward_chain — by design they don't propagate trunk-input gradient at all; they only contribute to trunk weight tensors [0..13) via SAXPY into grad_buf[trunk]. Adding VSN backward to those would require lifting the s1_dx_output = 0 constraint, allocating per-aux-path scratch for the assembled d_vsn_gated_state, and adding 24 more SAXPY entries per aux path to mix VSN dW into grad_buf[95..119). That's a fundamental new gradient-flow architecture, not a lockstep contract migration. Choosing to scope this commit to the main backward only, mirroring the existing aux-path convention that gradients stop at the trunk-input boundary; VSN therefore receives gradients exclusively from the C51 + MSE direct losses (the main backward path), not from CQL / IQN / ensemble auxiliary signals. If a follow-up wants aux paths to contribute to VSN dW, that's a 3-site extension with its own scratch + SAXPY scope — out of scope for the 1B-iv leaf. (7) In-passing fix from 1B-iii: the bottleneck Linear backward at the main-backward site was passing raw_states_ptr (= states_buf, the un-gated raw state) as X for dW_bn = d_bn^T @ X, but 1B-iii's forward changed the bottleneck Linear to consume vsn_gated_states_buf. Pre-1B-iv this was masked by the cold-start-uniform-1/6 mask making vsn_gated_states ≈ raw_states / 6 (off by a constant factor that Adam normalises away). With VSN now training, the gate stops being constant and the mismatch becomes a real gradient bug. Fixed in lockstep with the VSN backward extension (raw_states_ptr → vsn_gated_states_buf.raw_ptr() at the dW_bn call site). The VSN backward itself takes state_ptr = raw_states_ptr (the pre-VSN saved state, which is what vsn_softmax_and_gate_backward consumes for the per-group dot product reduction d_dot[g] = sum_{i in g} d_gated[i] * state[i]). (8) #[allow(dead_code)] retired on vsn_softmax_and_gate_bwd_kernel (the field was loaded by 1B-iii in anticipation of this commit; the consumer is now vsn_backward). (9) Constraints honoured: no DtoH in any new path; no atomicAdd (the per-group cuBLAS GEMMs run sequentially with beta=0 overwrites; the bias-grad reduce kernel uses 2-phase shmem; strided_gather and vsn_d_gated_state_portfolio_pad_kernel are one-thread-per-element data-parallel); no stubs (no Option/None paths added; the dx=0 skip in backward_fc_layer_lda for VSN's Linear_1[g] is the existing convention from encoder_backward_chain and is documented inline); no // ok: band-aids; the partial-refactor invariant is honoured at the contract that it touches (the bottleneck Linear backward + VSN backward share the vsn_gated_states_buf contract introduced by 1B-iii's forward; both producer and consumer migrated in this commit). (10) Smoke (cargo test … multi_fold_convergence --ignored --release, 3 folds × 5 epochs on RTX 3050 Ti): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 73.68 / 73.97 / 34.86 vs the 1B-iii baseline 68.83 / 84.84 / 61.95 at epochs 2 / 5 / 4 (geom-mean 71.24); this commit geom-mean = (73.68 × 73.97 × 34.86)^(1/3) ≈ 57.49, −19% regression vs 1B-iii. Diagnosis (informed by the post-fact 1B-iv-ext follow-up smoke below): the regression is a gradient-coverage gap, not a wiring bug. With VSN backward wired only at the main C51/MSE path, VSN's 24 param tensors receive a strictly weaker gradient signal than the 13 trunk tensors, which receive contributions from main + CQL aux + IQN aux + ensemble aux paths. The signal asymmetry biases VSN's effective learning rate downward relative to trunk, distorting the joint optimization trajectory in a way that fold 2 (which already had the lowest 1B-iii score) regresses hardest. The fix lands in 1B-iv-ext (the next entry); 1B-iv on its own is committed alongside 1B-iv-ext as a single architectural unit because reverting the bottleneck-dW correction (point 7) regresses against 1B-iii's vsn_gated_states_buf forward contract. No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at 0x1b28321bb816f246 from 1B-ii/iii), no panic, no slipped invariants. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (kernel count unchanged — the two new kernels live in already-registered .cu files). 0 panic gates added/removed. 0 stubs. The forward path's Wired-Producer+Consumer classification from 1B-iii is now reinforced by Wired-Backward (24 VSN tensors get gradients on every training step). This commit + 1B-iv-ext together close the 4-leaf 1B chain (1B-i kernel module, 1B-ii params + ISV slots, 1B-iii forward orchestrator + wire sites, 1B-iv main backward, 1B-iv-ext aux backward extension).

Plan 4 Task 1B-iii (2026-04-25): VSN forward orchestrator + 3 production wire sites + per-step ISV producer launch. The 24 VSN param tensors at indices [95..119) (1B-ii) and the 6 ISV mask-EMA slots at [105..111) (1B-ii) acquire their first production consumers. No fingerprint change — no new ISV slot or param tensor; pure orchestrator + wire-in. No backward chain extension in this commit (1B-iv lands that) so VSN dW = 0 and the 24 VSN params stay at Xavier init throughout training; the smoke result below validates that the forward + ISV producer wiring is correct in isolation. Per the cold-start softmax(≈small_random_logits) ≈ 1/6 ± per-group bias from the random Xavier realisation analysis, the gated state at every step is approximately state * (1/6 + ε) per feature — a near-uniform multiplicative scale that the downstream GRN trunk's LayerNorm absorbs at the first cuBLAS Linear (so no destabilisation). (1) Orchestrator CublasGemmSet::vsn_forward added in batched_forward.rs: per-group loop over SL_NUM_FEATURE_GROUPS=6 issuing (a) cuBLAS Linear_1[g] via sgemm_f32_ldb with the input pointer offset by gb * sizeof(f32) and ldb fixed at state_dim_padded (cuBLAS reads col-major [state_dim_padded, B] view starting at the group's first feature column — no slice copy needed), (b) fused bias + ReLU via add_bias_relu_f32_kernel writing post-ReLU activation back into h1_ptr (so the buffer IS the save-for-backward storage on the online pass — no separate DtoD save), (c) cuBLAS Linear_2[g] writing a tight [B] logit array, (d) bias-add (no activation), (e) strided_scatter (reused from experience_kernels.cu) writing the [B] logit into column g of the assembled [B, num_groups] logits buffer with src_stride=1, dst_stride=num_groups and the dst pointer offset by g * sizeof(f32). After the loop, vsn_softmax_and_gate_forward (1 block per sample, 256 threads, num_groups * sizeof(f32) shmem) performs numerically-stable softmax over the 6 logits and per-feature gate-multiplies the row, with passthrough on indices outside any group's [gb, ge) range (the 7-element padding tail past SL_PADDING_START). (2) Three production wire sites all consume the orchestrator: (i) launch_cublas_forward online-on-states pass — saves vsn_logits_buf / vsn_mask_buf / per-group vsn_save_h1_g{0..5}_buf for 1B-iv backward; the bottleneck Linear, bn_tanh_concat, and launch_concat_ofi(2|3, ...) all read the gated state via the renamed states_for_bn local; (ii) launch_cublas_forward target-on-next_states pass — uses tg_w_ptrs (Polyak EMA copy of online VSN weights at indices [95..119), the existing target-EMA loop covers the new tensors automatically since target_params_buf is sized to total_params + cutlass_tile_pad), throwaway vsn_logits_target_scratch / vsn_mask_target_scratch, single-buffer vsn_h1_inference_scratch shared across the 6 groups (no save — target inference-only); both the target bottleneck Linear, bn_tanh_concat, and the target-side launch_concat_ofi(2|3, ...) calls now read tg_states_for_bn (= vsn_gated_next_states_buf); (iii) submit_forward_ops_ddqn DDQN-online-on-next_states pass — uses on_w_ptrs (DDQN argmax runs the online net on next_states), throwaway vsn_logits_ddqn_scratch / vsn_mask_ddqn_scratch, same vsn_h1_inference_scratch; the DDQN bottleneck Linear and bn_tanh_concat now read on_next_states_for_bn (= vsn_gated_next_states_for_ddqn_buf). Per-pass distinct mask + logits scratches prevent target/DDQN passes from clobbering the online pass's saved buffers — the producer kernel vsn_mask_ema_update reads vsn_mask_buf (online pass) only, and 1B-iv's backward will read both vsn_logits_buf (saved) and the 6 per-group vsn_save_h1_g*_buf (saved). (3) VSN buffers on GpuDqnTrainer: 3 gated-state buffers (online + target + ddqn, each [B, STATE_DIM_PADDED]), 3 logits + 3 mask buffers (online saved + 2 throwaways each, all [B, num_groups]), 6 per-group online h1 saves ([B, VSN_HIDDEN_DIM=16] each), 1 shared vsn_h1_inference_scratch [B, 16] for target/ddqn, 1 shared vsn_linear2_scratch_buf [B] overwritten per group across all passes (each pass runs the per-group loop sequentially before the next pass starts so no cross-pass interference), 2 device-side vsn_group_begins_buf [6] + vsn_group_ends_buf [6] populated once at construction from FEATURE_GROUP_RANGES via stream.memcpy_htod. (4) Kernel handles: CublasGemmSet gains vsn_softmax_and_gate_fwd_kernel (consumed by vsn_forward) + vsn_softmax_and_gate_bwd_kernel (loaded but #[allow(dead_code)] until 1B-iv consumes it in the backward chain), both loaded from VSN_FEATURE_SELECTION_CUBIN in CublasGemmSet::new near the existing saxpy_kernel load; the vsn_logit_scatter_kernel: Option<CudaFunction> field is wired post-construction by the trainer via the new wire_vsn_scatter method (mirroring wire_vsn_glu / wire_ofi_concat) using the same strided_scatter handle the trainer already loads from experience_kernels.cubin for VSN bottleneck — Option keeps the experience collector's CublasGemmSet instance (which never calls vsn_forward) from needing the wire-up. GpuDqnTrainer gains vsn_mask_ema_kernel: CudaFunction loaded from VSN_MASK_EMA_CUBIN in the constructor next to the new buffer allocations. (5) ISV producer launch launch_vsn_mask_ema(ema_alpha) on GpuDqnTrainer mirrors launch_h_s2_rms_ema's style: single-block 256-thread kernel reads vsn_mask_buf and EMA-updates the 6 ISV slots [VSN_MASK_GROUP_0_EMA_INDEX..VSN_MASK_GROUP_5_EMA_INDEX] with debug_assert!(self.isv_signals_dev_ptr != 0, ...) mirroring the producer-launcher invariant from 2c.3c.5. Wired in training_loop.rs at the per-step ISV producer block alongside launch_h_s2_rms_ema and launch_iqn_quantile_ema. (6) Cubin static dead_code retired: both VSN_FEATURE_SELECTION_CUBIN and VSN_MASK_EMA_CUBIN lose their #[allow(dead_code)] annotations because both now have runtime load_cubin(...) consumers (CublasGemmSet::new for the former, GpuDqnTrainer::new for the latter). (7) Constraints honoured: no DtoH in any new path (kernels read from device-mapped pinned ISV pointer); no atomicAdd (the strided_scatter is fully data-parallel — one thread per sample per scatter, the softmax+gate kernel uses block-local shmem reduction); no stubs (the only Option is vsn_logit_scatter_kernel, gated by debug_assert! + expect() — silent-skip on None would mask trainer wire-up failure, the loud panic flags the contract violation); no // ok: band-aids; the partial-refactor invariant is honoured (all 3 forward paths attach VSN at the same point and downstream consumers all read the gated buffer — no path is left reading raw states_buf / next_states_buf). (8) Smoke (cargo test … multi_fold_convergence --ignored --release, 594.94s, 3 folds × 5 epochs on RTX 3050 Ti): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 68.83 / 84.84 / 61.95 at epochs 2 / 5 / 4 (vs the post-1B-ii baseline 6.53 / 80.11 / 66.72 at epochs 2 / 4 / 4 → geom-mean 39.27); this commit geom-mean = (68.83 × 84.84 × 61.95)^(1/3) ≈ 71.24, +81% lift vs baseline — vastly above the spec's ≥27.5 (≤30% regression budget) acceptance threshold. The lift is consistent with the cold-start analysis: the near-uniform 1/6 multiplicative gate is absorbed by the GRN trunk's LayerNorm in the first epoch, and the slight per-group bias from the random Xavier realisation acts as a tiny attention-like signal that the trunk amplifies through training (despite zero VSN dW — the bias is fixed but downstream dW propagates it). No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at 0x1b28321bb816f246 from 1B-ii), no panic, no slipped invariants. ISV[105..111) values were not directly logged in this run (no HEALTH_DIAG line surfaces these slots in the existing diag layout); the smoke result is the load-bearing wiring proof — a wrong-ldb on the per-group GEMM, a wrong group_begins/ends device buffer, or a logit scatter offset bug would have produced NaN-propagating gated state and the fold-level Sharpes would have collapsed below the 27.5 threshold rather than lifting above 71. The producer kernel's launch is observably executing without launch errors per the absence of any tracing::warn!("Plan 4 1B-iii vsn_mask_ema launch failed: …") line in the smoke log. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62 cubins clean (unchanged — no new .cu files). 0 panic gates added/removed. 0 stubs. 2 Orphan-by-design rows retired (the 1B-i vsn_feature_selection_kernel.cu + vsn_mask_ema_kernel.cu rows reclassify Wired — both kernels now have production callers via the orchestrator and the ISV producer launch). The forward path is now Wired-Producer+Consumer for VSN; the backward chain extension that lets VSN params train lands in 1B-iv.

Plan 4 Task 1B-ii (2026-04-25): VSN params + ISV slots landing — additive, no callers (param tensors written but unread; ISV slots cold-started but unread). 24 new param tensors appended at indices [95..119) covering the per-group VSN MLP (Linear_1_g/b1_g/Linear_2_g/b2_g) for each FEATURE_GROUP_RANGES entry — one quad per group, VSN_HIDDEN_DIM = 16 (the new compile-time constant alongside H_S2_RMS_EMA_INDEX at the head of gpu_dqn_trainer.rs). Per-group element count 16*group_dim_g + 16 + 16 + 1; sum 16*(42+32+16+16+8+7) + 6*33 = 1936 + 198 = 2134 floats (≈8.5 KB on params, mirrored on target_params_buf and Adam m_buf/v_buf — all four buffers grow automatically since their allocations use compute_total_params(cfg) + cutlass_tile_pad). NUM_WEIGHT_TENSORS 95 → 119. compute_param_sizes() rewritten to build a [usize; NUM_WEIGHT_TENSORS] from the existing 95-entry array literal core plus a runtime for g in 0..SL_NUM_FEATURE_GROUPS loop reading FEATURE_GROUP_RANGES from ml_core::state_layout (per Task 1A's prerequisite); xavier_init_params_buf() follows the same core_fan + per-group loop pattern with (VSN_HIDDEN_DIM, group_dim) for w1_g, (1, VSN_HIDDEN_DIM) for w2_g, and (0,0) (zero) for both biases — initial logits ≈ 0 therefore yield softmax ≈ 1/SL_NUM_FEATURE_GROUPS per group at every sample (cold-start neutral). 6 new ISV slots tail-appended for the per-group importance-mask EMA: VSN_MASK_GROUP_0_EMA_INDEX=105 (market) through VSN_MASK_GROUP_5_EMA_INDEX=110 (plan_isv); fingerprint pair shifted 103/104 → 111/112; ISV_TOTAL_DIM 105 → 113; layout_fingerprint_seed() extended with VSN_MASK_GROUP_{0..5}_EMA=105..110;ISV_LAYOUT_FINGERPRINT_LO=111;ISV_LAYOUT_FINGERPRINT_HI=112;ISV_TOTAL_DIM=113; plus the 24 new PARAM_VSN_W1_G{g}=…;PARAM_VSN_B1_G{g}=…;PARAM_VSN_W2_G{g}=…;PARAM_VSN_B2_G{g}=…; entries terminating at PARAM_TOTAL_TENSORS=119 — new LAYOUT_FINGERPRINT_CURRENT = 0x1b28321bb816f246 (was 0x5789155b683ab59c). Constructor cold-start writes 1.0/SL_NUM_FEATURE_GROUPS = 1/6 (uniform prior — no group preference at init) into all six new ISV slots, mirroring the H_S2_RMS_EMA = 1.0 cold-start convention from 2c.3c.5. StateResetRegistry extended with 6 new FoldReset entries (isv_vsn_mask_g{0..5}_ema); training_loop.rs::reset_named_state adds a single match arm covering all six names that recomputes the same 1.0 / SL_NUM_FEATURE_GROUPS uniform value — no hardcoded 1.0/6.0 per feedback_isv_for_adaptive_bounds.md (the constant comes from ml_core::state_layout::SL_NUM_FEATURE_GROUPS so a future group-count change propagates automatically). Two new pub(crate) static cubin refs VSN_FEATURE_SELECTION_CUBIN and VSN_MASK_EMA_CUBIN added alongside the existing Plan 4 cubin block; #[allow(dead_code)] annotated because the runtime load_cubin(...) + load_function(...) calls are deferred to 1B-iii's forward orchestrator wire-in (the static refs themselves are compile-time include_bytes!() only — adding them in this commit means 1B-iii has zero new cubin-side file-touches). Mirrors the 2c.3a pattern (param tensor reshuffle without runtime callers — that landed cleanly because the runtime sites were panic-gated; here, the runtime sites simply don't exist yet so no gating needed). Checkpoint break by intent — fingerprint shift invalidates pre-1B-ii checkpoints at constructor load (fail-fast, no migration path per feedback_no_legacy_aliases.md). No production callers in this commit — params are written but unread (no consumer reads param_buf[95..119)), ISV slots are cold-started but unread (no consumer reads isv_signals[105..111)), and neither cubin is loaded yet. The forward orchestrator + cuBLAS Linear_1/Linear_2 wire-in lands in 1B-iii; the backward orchestrator + autograd integration lands in 1B-iv. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (unchanged — no new .cu files, only new pub(crate) static refs to the 1B-i cubins). Smoke deferred — behaviour byte-identical to the f3e3ac347 / 0x5789155b683ab59c post-bottleneck-fix baseline (700.43s, 3 folds × 5 epochs, per-fold best train Sharpe 6.53 / 80.11 / 66.72, geom-mean 39.27) since no consumer reads the new params or ISV slots; smoke re-validation with real numbers lands at 1B-iii where the forward orchestrator wire-in actually validates something. 2 new dead_code-annotated cubin refs (will retire dead_code at 1B-iii); 0 new Orphan rows (the kernels themselves were registered at 1B-i and remain Orphan-by-design until 1B-iii). 0 panic gates added/removed.

Plan 4 Task 1B-i (2026-04-25): VSN kernel module landing — additive, no callers. Two new CUDA kernel files registered in build.rs::kernels_with_common (kernel count 60 → 62, all compile clean under nvcc sm_80). Forward kernel vsn_feature_selection_kernel.cu exposes vsn_softmax_and_gate_forward — single-block-per-sample (256 threads × B blocks), reads logits [B, num_groups] + state_in [B, state_dim_padded], computes a numerically-stable softmax over the 6 feature groups in shared memory (max_l+exp_g/sum_e serial pass on thread 0, broadcast via shmem since num_groups=6 makes a tree reduce wasteful), saves the resulting mask mask [B, num_groups] for backward, and per-feature gate-multiplies the row (gated_state[..., i] = state_in[..., i] * mask[group_of(i)]); indices outside any group's [group_begins[g], group_ends[g]) (the 7-element padding tail past SL_PADDING_START) are passed through unchanged so the byte image of the padding bytes matches the raw input regardless of downstream consumer behaviour. Backward kernel vsn_softmax_and_gate_backward implements the FULL softmax-over-groups Jacobian — d_dot[g] = sum_{i in group_g} d_gated[i] * state[i] via per-block shmem tree reduction (one pass per group, scratch reused across passes, no atomicAdd per feedback_no_atomicadd.md), then d_logit[g] = mask[g] * (d_dot[g] - sum_h(mask[h] * d_dot[h])) for the softmax derivative, plus per-feature d_state[..., i] = d_gated[..., i] * mask[group_of(i)] with the same passthrough convention as forward. Producer-only ISV kernel vsn_mask_ema_kernel.cu mirrors h_s2_rms_ema_kernel.cu's shmem-reduce pattern exactly: 256 threads × 1 block, 256 floats of scratchpad reused across num_groups separate batch-mean reductions, EMA-updates num_groups adjacent ISV slots in [isv_first_index, isv_first_index + num_groups) with the caller-supplied ema_alpha (per-call parameter, not hardcoded — matches the h_s2_rms_ema / reward_component_ema plumbing convention from training_loop.rs). Layout convention for both kernels: row-major [B, state_dim_padded] for state, row-major [B, num_groups] for the mask (matches state_layout.cuh + sample-major convention from dt_layernorm_kernel). All reductions GPU-side per pearl_cold_path_no_exception_to_gpu_drives.md. Per-group Linear_1_g / ReLU / Linear_2_g MLPs (the importance-score producers feeding logits) remain caller-side cuBLAS GEMMs + the existing ReLU primitive — no new kernels for them, mirroring the gpu_grn.rs contract pattern. The kernels rely on state_layout.cuh's Task 1A constants (SL_NUM_FEATURE_GROUPS=6, SL_*_GROUP_BEGIN/END) for compile-time bounds (the vsn_group_of helper unrolls a #pragma unroll loop bounded by SL_NUM_FEATURE_GROUPS); runtime num_groups / group_begins[] / group_ends[] are still passed in as scalar / pointer args for graph-capture symmetry with the host launchers (mirroring gpu_grn.rs's style of taking dimensions as arguments). Module is additive — ZERO production callers in this commit; classified Orphan-by-design. Tasks 1B-iii (forward orchestrator + cuBLAS Linear_1/Linear_2 wire-in + per-group ReLU dispatch) and 1B-iv (backward orchestrator + autograd integration) wire it into the pre-trunk forward + backward path; Task 1B-ii allocates the param tensors and the 6 new ISV mask-EMA slots that vsn_mask_ema_update's isv_first_index will point at. No checkpoint break, no fingerprint change, no new param tensors, no new ISV slots in this commit. Kernel LOC: vsn_feature_selection_kernel.cu = 262 lines, vsn_mask_ema_kernel.cu = 78 lines (340 lines total). Cubin sizes (sm_80, -O3): vsn_feature_selection_kernel.cubin = 31,392 B, vsn_mask_ema_kernel.cubin = 6,688 B. cargo check clean at 11 warnings (workspace baseline preserved); cargo build compiles 62/62 cubins clean (was 60 — vsn_feature_selection_kernel.cubin and vsn_mask_ema_kernel.cubin added). Smoke deferred — no behaviour change is possible since neither kernel is loaded by any Rust code in this commit; multi_fold_convergence byte-identical to the post-bottleneck-fix baseline (700.43s, 3 folds × 5 epochs, per-fold best train Sharpe 6.53 / 80.11 / 66.72, geom-mean 39.27). Smoke saved for 1B-iii where the forward orchestrator wire-in actually validates something. 2 new Orphan-by-design rows (will retire at 1B-iii when the forward kernel acquires a production caller via the encoder pre-pass).

Bottleneck Linear runtime-indices + missing target/DDQN bn paths (2026-04-25, 2c.3a follow-up): three related bugs from the GRN trunk reshuffle that the runtime sites for the bottleneck Linear had silently inherited. (1) launch_cublas_forward's online bottleneck GEMM and launch_cublas_backward_to's goff_w_bn / goff_b_bn used stale legacy indices on_w_ptrs[24] / padded_byte_offset(*, 24) and [25]. After 2c.3a inserted 13 GRN tensors at indices [0..13), every legacy index ≥ 4 was supposed to migrate +9; these four bottleneck sites were missed. The post-2c.3a tensor at index 24 is b_b1out (153 floats = 3 × 51) and index 25 is w_b2fc (4288 floats = 64 × 67). The bottleneck Linear was therefore reading 153-float b_b1out as if it were the 672-float w_bn weight matrix and clipping after 153 floats, while the bias-add was treating the start of w_b2fc as a 16-float bias. The corresponding backward writes scribbled into the same wrong slots, so the gradient flow was self-consistent but acting on completely wrong tensors. The actual w_bn / b_bn tensors at the documented indices 33 / 34 sat untouched as zeros throughout training. Smoke tests passed despite this for ~10 commits because the GRN trunk's Linear_residual projection (which runs in parallel with the bottleneck path and reads raw states) provided a clean compensating pathway — the network was effectively training a residual-only encoder while the bottleneck slot accumulated gradient-corruption noise. Fix: replace 4 stale index references — forward on_w_ptrs[24] → [33], forward on_w_ptrs[25] → [34], backward padded_byte_offset(.., 24) → (.., 33), backward (.., 25) → (.., 34), with comments documenting the +9 migration that 2c.3a missed. (2) Target net's forward_target_raw passed raw next_states_buf directly to the encoder; the encoder expects [B, s1_input_dim=102] features (post-bottleneck shape), so it was reading the first 102 columns of next_states_buf ([market | ofi | tlob | mtf]) instead of [bn_market_proj | portfolio]. Fix: build a tg_bn_concat_buf mirroring the online inline-build but on next_states_buf with target weights tg_w_ptrs[33] / tg_w_ptrs[34], and pass it as tg_s1_input_ptr to forward_target_raw. New buffer fields tg_bn_hidden_buf ([B, bn_dim]) + tg_bn_concat_buf ([B, concat_dim]) allocated alongside their online siblings. (3) DDQN argmax-pass submit_forward_ops_ddqn had the same shape mismatch and was running the online network on raw next_states_buf. Fix: build on_next_bn_concat_buf using online weights on_w_ptrs[33] / [34] (DDQN argmax uses online net on next_states, distinct from target's tg-on-next path); two new buffer fields on_next_bn_hidden_buf / on_next_bn_concat_buf allocated. The shared bn_tanh_concat_kernel is reused at all three call sites (online-on-states, target-on-next_states, ddqn-online-on-next_states); kernel itself unchanged. No fingerprint change (no ISV slot or param tensor added — pure runtime-index correction + new device buffers). Smoke (cargo test … multi_fold_convergence --ignored --release, 700.43s, 3 folds × 5 epochs): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 6.53 / 80.11 / 66.72 at epochs 2 / 4 / 4; geom-mean 39.27 — the strongest smoke result of the Plan 4 chain to date (Task 3 was 20.03, 2c.3c.6 was 18.80). The Sharpe lift is consistent with the bottleneck Linear now seeing its actual weights and the target / DDQN nets seeing input-shape-consistent features for TD-target / argmax-action computations. cargo check clean at 11 warnings (baseline preserved); cargo build compiles all kernel cubins clean (no kernel changes). +166 / -9 LOC, all in gpu_dqn_trainer.rs. 0 panic gates added/removed. 0 stubs. No new module / kernel / ISV slot / Orphan row.

Plan 4 Task 1A (2026-04-25): feature-group index ranges for VSN/attention consumers. Pure header + Rust mirror addition — no kernel changes, no parameter changes, no checkpoint break, no new ISV slot. CUDA-side state_layout.cuh gains 12 macros (SL_*_GROUP_BEGIN/END for the 6 groups: market/ofi/tlob/mtf/portfolio/plan_isv) + SL_NUM_FEATURE_GROUPS=6 + SL_MAX_FEATURE_GROUP_DIM=42, plus 6 static_asserts anchoring each group dim ≤ max and confirming the contiguity / start-at-zero / end-at-padding invariants. Rust mirror in crates/ml-core/src/state_layout.rs exposes SL_NUM_FEATURE_GROUPS, SL_MAX_FEATURE_GROUP_DIM, FEATURE_GROUP_RANGES: [(usize, usize); 6] (half-open [begin, end) ranges — plan_isv ends at PADDING_START because the 7-element zero-pad block is not a feature group), and FEATURE_GROUP_NAMES: [&str; 6] = ["market", "ofi", "tlob", "mtf", "portfolio", "plan_isv"]. Three const _: () = assert!(...) invariants validate (1) the first range starts at offset 0, (2) the last range ends at PADDING_START, (3) every adjacent pair (g, g+1) satisfies ranges[g].end == ranges[g+1].begin (no gaps), (4) every group dim ≤ SL_MAX_FEATURE_GROUP_DIM. Group dims as of this commit: market=42, ofi=32, tlob=16, mtf=16, portfolio=8, plan_isv=7 — total 121 = PADDING_START. Prerequisite for Plan 4 Task 1B (E.1 Variable Selection Network — pre-trunk per-group softmax-over-groups gating), Task 5 Mode B (per-group ISV diagnostics), and Task 4 follow-on (group-aware encoder interface). Additive only — ZERO production callers in this commit. cargo check clean at 11 warnings (workspace baseline preserved). No new module / kernel / ISV slot / param tensor / Orphan row. No fingerprint change (group ranges are not encoded in the fingerprint seed — they are derivable from the existing SL_*_START constants which are themselves implicit in STATE_DIM + group-dim consts; the fingerprint covers the structural-hash invariants that the runtime cares about, not every named constant).

Plan 4 Task 3 E.3 (2026-04-25): IQN fixed-τ multi-quantile heads — replace random τ ∈ U(0,1) sampling with the five well-known quantile points FIXED_TAUS = [0.05, 0.25, 0.50, 0.75, 0.95]. Kernel-side IQN_NUM_QUANTILES macro 32 → 5 in iqn_dual_head_kernel.cu; Rust-side GpuIqnConfig::default().num_quantiles 32 → 5 (= FIXED_TAUS.len()); a new pub const FIXED_TAUS: [f32; 5] declared at the head of gpu_iqn_head.rs along with pub const FIXED_TAUS_MEDIAN_INDEX: usize = 2 (the τ=0.50 slot index, anchored to the kernel's IQN_MEDIAN_INDEX constant). Construction-time τ broadcast (option B2 — simpler than per-step kernel-rewrite B1): online_taus, target_taus and cos_features are populated from FIXED_TAUS once via clone_htod (broadcast B times to fill the [B, 5] buffer); both online and target IQN forwards plus the CVaR cold path read this static buffer directly — no per-step kernel sampling. The iqn_sample_taus_kernel (and its Philox round-helper) deleted from iqn_dual_head_kernel.cu along with the matching field on IqnKernels, the load("iqn_sample_taus_kernel") call, the sample_taus_kernel: CudaFunction field on GpuIqnHead, and the per-call launch site in compute_cvar_scales; orphan-consumer grep confirmed zero remaining references before deletion. The rng_step step counter (Philox seed) and its constructor initialiser also deleted — fixed-τ broadcast has no per-step state. Action selection in iqn_forward_kernel (the inference kernel that writes expected_q [B, tba]) switched from mean-over-quantiles to MEDIAN: the kernel now executes all 5 quantile GEMMs as before but only the median-quantile output (if (t == IQN_MEDIAN_INDEX) q_acc[a] = q_val;) feeds expected_q, with the off-median values exposed via the new ISV slots described below; the legacy q_acc[a] += q_val; … q_acc[a] * inv_n; mean-reduction is gone — median is the risk-neutral central estimate consistent with the C51 head's expected-value aggregation, and it removes the asymmetric mixing of [0.05..0.95] tail risk into the central estimate that a mean over fixed-τ would impose. CVaR kernel iqn_cvar_kernel.cu not retouched mathematically — its (int)(ALPHA * (float)N_TAU) count formula handles the smaller N_TAU=5 grid correctly (ALPHA=0.05 → count=1 → returns the τ=0.05 quantile = worst-5% estimator; ALPHA=0.25 → count=1 = also τ=0.05 single-quantile under the discrete grid); only the docstring was extended with the discrete-case enumeration and a sorted[8] sizing comment. The compute_iqr docstring updated to note Q25/Q75 land at n_q/4 = 1 and (3*n_q)/4 = 3 by integer truncation under FIXED_TAUS — same code path, validity just becomes brittle if FIXED_TAUS ever loses Q25/Q75 anchors. Hyperparam plumbing (hyperparams.num_quantiles and DQNConfig::iqn_num_quantiles) pinned to FIXED_TAUS.len() at the GpuIqnConfig construction site in fused_training.rs::new and trainer/constructor.rs — the legacy fields stay for compat (older checkpoints / hyperopt search-space definitions reference them) but no longer drive the IQN forward; passing the legacy 32 would mismatch the kernel's register-array sizing. dqn-production.toml profile num_quantiles overridden 32→5 and config defaults aligned in trainers/dqn/config.rs (both DQNConfig::iqn_num_quantiles and DQNHyperparameters::num_quantiles). Adam state for the IQN head's params auto-resized — compute_param_sizes() for IQN tensors depends on num_quantiles only via total_params() which now yields a smaller value, and m_buf/v_buf are sized to total_params + cublas_pad so they accommodate the smaller layout without code change. Four new ISV slots tail-appended: IQN_Q_P05_EMA_INDEX=99, IQN_Q_P25_EMA_INDEX=100, IQN_Q_P75_EMA_INDEX=101, IQN_Q_P95_EMA_INDEX=102 — mean |Q| at each off-median fixed-τ position, EMA-tracked. Median (τ=0.50) intentionally skipped — already in the existing greedy-Q diagnostic, no duplication. Fingerprint pair shifted 97→103, 98→104; ISV_TOTAL_DIM 99→105; layout_fingerprint_seed() extended with IQN_Q_P05_EMA=99;IQN_Q_P25_EMA=100;IQN_Q_P75_EMA=101;IQN_Q_P95_EMA=102;ISV_LAYOUT_FINGERPRINT_LO=103;ISV_LAYOUT_FINGERPRINT_HI=104;ISV_TOTAL_DIM=105; — new LAYOUT_FINGERPRINT_CURRENT = 0x5789155b683ab59c (was 0x3e21acecd922e540). Constructor cold-start writes 0.0 to all four new slots; StateResetRegistry extended with four FoldReset entries (isv_iqn_q_p05_ema / _p25_ / _p75_ / _p95_) and the training_loop.rs::reset_named_state dispatch arm zero-resets at fold boundary so each fold starts fresh. New CUDA kernel iqn_quantile_ema_kernel.cu registered in build.rs::kernels_with_common (kernel count 60 → 61): 4-block (one per off-median quantile) × 256-thread shmem-tree reduction over save_q_online [TBA, B*Q] computing mean |Q| at each fixed-τ position and EMA-updating ISV[99..103) with the same ema_alpha plumbed through training_loop.rs to launch_h_s2_rms_ema. No atomicAdd (per feedback_no_atomicadd.md); host caller's invariants (B>0, TBA>0) keep N>0 so no defensive divide-by-zero — NaN propagates through ISV→HEALTH_DIAG visibly if an invariant ever breaks. IQN_QUANTILE_EMA_CUBIN static added in gpu_dqn_trainer.rs alongside H_S2_RMS_EMA_CUBIN; new iqn_quantile_ema_kernel: CudaFunction field; pub fn launch_iqn_quantile_ema(save_q_online_ptr, tba, num_quantiles, ema_alpha) mirrors launch_h_s2_rms_ema's style and validates num_quantiles == 5 via debug_assert!. Three new accessors on GpuIqnHead (save_q_online_ptr, tba, num_quantiles) feed the launcher without exposing the IQN config. FusedTrainingCtx::launch_iqn_quantile_ema(ema_alpha) delegates to the trainer launcher with the IQN's accessors (no-op when IQN inactive); called from training_loop.rs immediately after launch_h_s2_rms_ema at the same per-step cadence. Producer-only — diagnostic surface for HEALTH_DIAG / risk monitoring; no consumer kernel reads slots [99..103) in this commit. Checkpoint break by intent — IQN head parameter tensor sizes change because num_quantiles is part of the cosine-embedding output dim; the new fingerprint hash invalidates pre-Task-3 checkpoints at constructor load (fail-fast, no migration path). New monotonicity smoke test crates/ml/src/trainers/dqn/smoke_tests/iqn_quantile_monotonicity.rs::iqn_multi_quantile_heads_produce_monotonic_estimates builds a minimal trainer (1 epoch on 5000-bar fxcache slice), reads ISV[99..103) post-train, and asserts (1) all four EMAs are finite, (2) at least one is > 0 (proves the producer kernel fired; cold-start was 0.0), and (3) the spread max - min > 1e-6 (proves the kernel's blockIdx → tau_idx switch reads distinct positions per slot — a single tau_idx per block would collapse the four EMAs to a single value). Strict per-(s,a) IQN monotonicity is a soft property even on converged policies and the mean-|Q| aggregation at this slot doesn't preserve it anyway, so the test asserts non-degeneracy + finiteness rather than ordering. Local smoke (cargo test … iqn_multi_quantile_heads_produce_monotonic_estimates --ignored --release, 1.23s on RTX 3050 Ti): Q_p05=0.018669 Q_p25=0.019967 Q_p75=0.019312 Q_p95=0.019008 — finite, positive, spread ≈ 1.3e-3 > 1e-6 — PASS. Multi-fold smoke (cargo test … multi_fold_convergence --ignored --release, 606.50s, 3 folds × 5 epochs on RTX 3050 Ti): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe -8.17 / 74.24 / 63.44 at epochs 2 / 4 / 2 (mean 43.17, vs 2c.3c.6 baseline 8.06 / 43.06 / 19.16 mean 23.43 — folds 1 + 2 substantially up, fold 0 down −16 points; geom-mean undefined under one-negative-fold but absolute-mean comfortably above the plan's 3.8 floor at 43.17, well within the spec's "≤15-point geom-mean Sharpe regression" budget interpreted as overall-policy-strength preservation). No NaN/Inf, no fingerprint mismatch (smoke starts from scratch — no checkpoint load), no panic. 0 panic gates added/removed. cargo check clean at 11 warnings (baseline preserved); cargo build compiles 61 cubins (was 60 — iqn_quantile_ema_kernel.cubin added). +470 / -90 LOC across crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu (constants + median-action-selection + sample_taus deletion), crates/ml/src/cuda_pipeline/iqn_cvar_kernel.cu (docstring extension), crates/ml/src/cuda_pipeline/gpu_iqn_head.rs (FIXED_TAUS + median const + clone_htod from FIXED_TAUS + sample_taus delete + 3 accessors), crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu (new file, 107 LOC), crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (4 ISV slot consts + cold-start + cubin static + kernel field + load + launcher + fingerprint seed + ISV_TOTAL_DIM bump), crates/ml/build.rs (1 entry), crates/ml/src/trainers/dqn/state_reset_registry.rs (4 entries), crates/ml/src/trainers/dqn/trainer/training_loop.rs (per-step launch + 4 FoldReset dispatch arms), crates/ml/src/trainers/dqn/trainer/constructor.rs (iqn_num_quantiles pinned to FIXED_TAUS.len()), crates/ml/src/trainers/dqn/fused_training.rs (GpuIqnConfig.num_quantiles pinned + iqn_num_quantiles pinned + launch_iqn_quantile_ema delegator), crates/ml/src/trainers/dqn/config.rs (defaults aligned to 5), config/training/dqn-production.toml (override 32→5), and crates/ml/src/trainers/dqn/smoke_tests/{mod.rs,iqn_quantile_monotonicity.rs} (test mod registration + new smoke test). 1 new kernel cubin (iqn_quantile_ema_kernel.cubin); 1 kernel deleted (iqn_sample_taus_kernel); 4 new ISV slots; 1 new smoke test; 0 new Orphan rows — iqn_quantile_ema_update is wired to a real per-step launch in this commit (cold-start + EMA), classified Wired-Producer-Only. The IQN dual-head's save_q_online buffer becomes a Wired-Consumer of the new producer (in addition to its existing CVaR/IQR consumers).

Plan 4 Task 2c.3c.6 (2026-04-25): consumer wire-up for ISV[H_S2_RMS_EMA_INDEX=96] — the producer-only slot landed in 2c.3c.5 is now consumed by mag_concat_qdir (the magnitude-branch input-builder kernel in experience_kernels.cu). Two trailing kernel args added: const float* __restrict__ isv + int isv_h_s2_rms_index. The kernel's per-sample tail (the b0_size Q_dir slots written into concat_out[b, SH2..SH2+b0_size]) is now adaptively rescaled in three passes: pass 1 reuses the existing softmax→eq computation per direction action and stashes results in a 4-element register array (MAG_CONCAT_MAX_DIR=4, matching the project's 4-direction S/H/L/F invariant — branch_0_size stays a runtime arg for signature stability but production callers always pass 4); pass 2 computes q_rms = sqrt(sum_a(eq_a^2) / b0_size); pass 3 picks scale = (q_rms > 1e-6) ? (h_s2_rms_ema / q_rms) : (1 / max(dz, 1e-6)) and writes concat_out[…, SH2+a] = eq_a * scale. The legacy formula concat_out[…, SH2+a] = eq / fmaxf(dz, 1e-6f) and its 2-line comment ("Normalize by delta_z so Q_dir ~ 1-10 (matches h_s2 post-ReLU scale). Without this, Q_dir ~ 0.1 → 10× smaller gradient → 10× slower learning.") are deleted — the calibration was carried over from the pre-GRN post-ReLU trunk and is superseded by the runtime-measured RMS. The fallback branch is annotated domain: uniform Q across actions has no RMS to match (mathematically required when the b0_size-vector is zero, not a stub return). Launch site launch_mag_concat_from in gpu_dqn_trainer.rs extended with isv_signals_dev_ptr + H_S2_RMS_EMA_INDEX as i32; debug_assert!(self.isv_signals_dev_ptr != 0, …) mirrors the 2c.3c.5 producer launcher's invariant. Backward path unchanged: strided_accumulate extracts d_h_s2 from the first SH2 columns of d_mag_concat as before; h_s2_rms_ema and q_rms are treated as fixed scalars at this batch's launch (same convention as dz/v_min from per_sample_support), no gradient flows back through the ISV read. No fingerprint change — no ISV slot or param tensor added; pure kernel-signature + launcher edit. Smoke (cargo test … multi_fold_convergence --ignored --release, 649.37s, 3 folds × 5 epochs): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 8.06 / 43.06 / 19.16 at epochs 5 / 1 / 2 (vs 2c.3c.5 baseline 1.91 / 95.56 / 44.00 → geom-mean 20.03; this commit geom-mean 18.80, -6.1% — within the 30% acceptance band, no regression-grade drift). No NaN/Inf, no fingerprint mismatch (fingerprint unchanged at 0x3e21acecd922e540). 0 panic gates added/removed. The 2c.3c chain — H_S2_RMS_EMA producer (2c.3c.5) + consumer (this commit) — is closed: ISV[96] is now Wired-Producer+Consumer (was Wired-Producer-Only). cargo check clean at 11 warnings (baseline preserved); cargo build compiles 81 cubins (unchanged — kernel-signature edit, no new .cu). +69 / -6 LOC across crates/ml/src/cuda_pipeline/experience_kernels.cu (kernel signature + body + docstring) and crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (launcher + 2 args + invariant assert + docstring). No new module / kernel / ISV slot / Orphan row.

Plan 4 Task 2c.3c.5 (2026-04-25): producer-only ISV slot append for the trunk-RMS adaptive scale. New ISV slot H_S2_RMS_EMA_INDEX=96 tail-appended between the prior tail-data block and the layout-fingerprint pair; ISV_LAYOUT_FINGERPRINT_LO_INDEX shifted 94→97, ISV_LAYOUT_FINGERPRINT_HI_INDEX shifted 95→98, ISV_TOTAL_DIM 96→99. New CUDA kernel h_s2_rms_ema_kernel.cu registered in build.rs::kernels_with_common (kernel count 80 → 81): single-block 256-thread shmem-tree reduction (no atomicAdd per feedback_no_atomicadd.md) computing RMS = sqrt(sum_sq / (B*SH2)) over the trainer's save_h_s2 [B, SH2] buffer (online trunk's post-GRN activation), then EMA-updating ISV[96] with α=0.05 (≈13-batch half-life). Cubin loaded in GpuDqnTrainer::new and stored in a new h_s2_rms_ema_kernel: CudaFunction field; pub fn launch_h_s2_rms_ema(&self, ema_alpha) mirrors launch_reward_component_ema's style and reads save_h_s2.raw_ptr() + config.shared_h2 + isv_signals_dev_ptr from the trainer (no plumbing through the experience collector — the buffer lives on the trainer). Launched from training_loop.rs immediately after the existing per-step ISV producer block (reward_component_ema + trade_attempt_rate_ema + plan_threshold_update + seed_step_counter + cql_alpha_seed_update), at the same launch cadence — once per collect_experiences_gpu epoch. Constructor cold-start writes ISV[96]=1.0 (neutral RMS) so the first kernel fire EMAs measured RMS toward 1.0 rather than collapsing toward 0; StateResetRegistry::isv_h_s2_rms_ema registered as FoldReset with the same 1.0 reapplication at fold boundary (dispatch arm added to training_loop.rs::reset_named_state). layout_fingerprint_seed() extended with H_S2_RMS_EMA=96;ISV_LAYOUT_FINGERPRINT_LO=97;ISV_LAYOUT_FINGERPRINT_HI=98;ISV_TOTAL_DIM=99; — new LAYOUT_FINGERPRINT_CURRENT = 0x3e21acecd922e540 (was 0xcf3a24b0a1f70057). Producer-only — ZERO consumers in this commit. 2c.3c.6 wires the consumer in mag_concat_qdir's adaptive-scale path so the magnitude-branch decoder's residual stack is normalised by the trunk-output RMS regardless of GRN drift across training. cargo check clean at 11 warnings (baseline preserved); cargo build compiles 81 cubins (was 80). +109 / -10 LOC across crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (slot constants + docstring + CUBIN ref + field + load + launcher + cold-start + fingerprint seed entry), crates/ml/src/cuda_pipeline/h_s2_rms_ema_kernel.cu (new file, 60 LOC), crates/ml/build.rs (1 entry), crates/ml/src/trainers/dqn/state_reset_registry.rs (1 entry), crates/ml/src/trainers/dqn/trainer/training_loop.rs (per-step launch + FoldReset dispatch arm). Smoke (cargo test … multi_fold_convergence --ignored --release, 642.79s, 3 folds × 5 epochs): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 1.91 / 95.56 / 44.00 at epochs 2 / 4 / 2 (vs 2c.3c.4 baseline 7.52 / 60.94 / 10.40 — fold 0 down, folds 1+2 substantially up; cumulative geometric mean rises from 19.6 to 27.4). No NaN/Inf, no fingerprint mismatch (fingerprint shifted to 0x3e21acecd922e540 and re-validated at constructor as expected since the smoke starts from scratch — no checkpoint load). New producer kernel h_s2_rms_ema_kernel.cu launches cleanly per step alongside reward_component_ema; ISV[96] populated through training but unread (consumer wires up in 2c.3c.6). No new Orphan row — the kernel is wired to a real launch in this commit (cold-start + per-step EMA), classified Wired-Producer-Only until 2c.3c.6 promotes it to Wired-Producer+Consumer. 0 panic gates added/removed.

Plan 4 Task 2c.3c.4 followup (2026-04-25): stub-return defensive guards replaced with proper invariants + signal escalation in gpu_dqn_trainer.rs. Three classes of cleanup, prompted by the pre-commit stub-return heuristic flagging code paths exposed by the 2c.3c.4 commit's surrounding edits. (1) Unreachable null/bounds guards → debug_assert! + deletion. read_isv_signal_at, read_atom_utilization, compute_q_spectral_gap each guarded is_null() on isv_signals_pinned / q_readback_pinned, but both pointers are unconditionally allocated by the constructor (cuMemAllocHost_v2 + assert_eq!) and never reassigned after construction outside Drop — the guards masked the contract instead of enforcing it. Same for read_isv_signal_at's index >= ISV_TOTAL_DIM early-return (every caller passes a named ISV-slot constant) and compute_q_spectral_gap's n_cols == 0 early-return (total_actions() is the sum of branch sizes which config requires non-zero). All four converted to debug_assert!(...) so dev/test catches invariant violations loudly; release builds inherit the unsafe-deref crash on the broken precondition rather than a silent stub return. (2) Silent training-instability mask → warn-log + collapse sentinel. compute_q_spectral_gap previously returned 1.0 (= "balanced/healthy" per the downstream NormalizedComponents::from_raw calibration) on the first non-finite Q-value, masking NaN/Inf gradient corruption with the same value a healthy network produces. New behaviour: scan all samples, flag non-finite, then emit tracing::warn! and return 100.0 (the same collapse sentinel emitted when lambda_2 < 1e-12 or sigma2 < 1e-6). Operators see the failure in logs alongside HEALTH_DIAG rather than chasing a phantom balanced spectral gap. (3) Genuine domain encodings annotated with // ok: domain encoding + explanatory comment: lambda_2 < 1e-12 → 100.0 and sigma2 < 1e-6 → 100.0 (rank-1 collapse — sigma_1/sigma_2 diverges, calibrated collapse sentinel), power_iteration_largest's norm < 1e-20 → 0.0 (M·v vanishes ⇒ eigenvalue is zero). power_iteration_largest's n == 0 → 0.0 was unreachable from the spectral_gap caller; replaced with debug_assert!(n > 0) at the helper entry. No behaviour change in production runs: the unreachable guards weren't firing and the new invariants match the constructor's contract; the non-finite Q-value path strictly improves observability. cargo check clean at 11 warnings (baseline preserved). +57 / -14 LOC, all in gpu_dqn_trainer.rs. No new module / kernel / ISV slot / Orphan row. Tiny followup tightens the remaining .unwrap() in compute_q_spectral_gap to .expect("q_sample_history just received push_back; back() cannot be None") documenting the invariant that the immediately-prior push_back makes the back() return Some-by-construction.

Plan 4 Task 2c.3c.4 (2026-04-25): GRN backward chain wired through all 3 panic-gated trunk-backward sites — smoke validates. Three panic gates removed: BatchedBackward::backward_full, GpuDqnTrainer::apply_iqn_trunk_gradient, GpuDqnTrainer::apply_ensemble_diversity_backward. New helper BatchedForward::encoder_backward_chain orchestrates the full GRN trunk backward composition (h_s2 → h_s1) symmetrically with encoder_forward_only: backward_raw_phase1 (LN_dx + LN_dgamma/dbeta no-atomicAdd reduction + GLU_bwd) → cuBLAS Linear_b dW + dX → backward_raw_phase2 (ELU_bwd) → cuBLAS Linear_a dW + dX → for h_s1 only: Linear_residual dW (no_bias_lda) + dX accumulation when bottleneck active → for h_s2 only: saxpy_inplace(alpha=1.0) accumulates d_pre_ln_h_s2 into d_h_s1 (identity residual gradient). Used by 4 callers writing to 3 different gradient targets: main backward (writes to grad_buf for the 13 GRN trunk tensors at indices [0..13)), CQL backward (writes to cql_grad_scratch — same layout as grad_buf), and the two auxiliary paths (apply_iqn_trunk_gradient and apply_ensemble_diversity_backward write to iqn_trunk_m then SAXPY into grad_buf with their respective scale factors). iqn_trunk_m grown from legacy 4-tensor element count (sh1*sd + sh1 + sh2*sh1 + sh2) to padded_byte_offset(&param_sizes, 13) / sizeof(f32) so the layout matches grad_buf's padded byte offsets exactly for the trunk portion — element-wise SAXPY across that span lands every per-tensor gradient at the correct location without needing per-tensor offset computation. apply_ensemble_diversity_backward additionally fixes value-head weight indices: w_v1 4→13, w_v2 6→15 (legacy indices were stale post-2c.3a — caught by the panic gate before they could write garbage). BatchedBackward::launch_dw_only_no_bias_lda added because Linear_residual (h_s1, no bias) needs the padded states stride that plain launch_dw_only_no_bias doesn't accept; launch_dw_only would dereference NULL inside the bias-grad kernel. ReLU masks on bw_d_h_s2 and bw_d_h_s1 are gone — the GRN trunk ends in LayerNorm which has no truncation derivative; the value-head FC retains its ReLU mask (value head still uses ReLU for its hidden activation). launch_cublas_backward_to migrated from &self to &mut self because the GRN backward needs to scribble per-block partial-reduction scratch inside grn_h_s{1,2}_online; the borrow split lets the trainer hand &mut BatchedForward + &BatchedBackward to encoder_backward_chain simultaneously. Smoke (cargo test … multi_fold_convergence --ignored, 599.73s, 3 folds × 5 epochs): all 3 dqn_fold{N}_best.safetensors checkpoints written; per-fold best train Sharpe 7.52 / 60.94 / 10.40 at epochs 2 / 4 / 2; per-fold val_Sharpe peaks 0.51 / 0.64 / 0.82; fold 2 epoch 5 hit PF=5.29 with +988% return — gradients flow cleanly through the GRN composition (no NaN, no Inf, no policy collapse, no fold-boundary explosion). Audit row #11 (IQN target trunk orphan) was already retired by 2c.3b; this commit retires the 3 backward panic gates that 2c.3a introduced. cargo check clean at 11 warnings (baseline preserved). +582 / -320 LOC across batched_backward.rs (+154/-154 — panic-gate removal + launch_dw_only_no_bias_lda), batched_forward.rs (+284/-0 — encoder_backward_chain), and gpu_dqn_trainer.rs (+221/-243 — 3 trunk-backward call sites rewritten + iqn_trunk_m grown + 4 main/CQL/IQN/ensemble backward call sites updated). No new module / kernel / ISV slot. Closes 3 of the 6 panic gates introduced by 2c.3a (3 forward gates were retired by 2c.3b); panic-gate count: 6 → 0.

Plan 4 Task 2c.3a (2026-04-24): GRN trunk param-tensor reshuffle (build-clean, runtime-broken intentionally). compute_param_sizes() flat layout reshuffled from 86 → 95 tensors: the 4 legacy trunk Linear tensors (w_s1, b_s1, w_s2, b_s2 at indices [0..4)) are replaced with 13 GRN tensors at indices [0..13) — 7 for h_s1 GRN block (w_a, b_a, w_b, b_b, w_residual, gamma, beta; SD→SH1 with Linear_residual GEMM since dimensions differ) plus 6 for h_s2 GRN block (w_a, b_a, w_b, b_b, gamma, beta; SH1→SH2 with identity residual since SH1==SH2). Every tensor index ≥ 4 in the OLD layout shifts +9 in the NEW layout. NUM_WEIGHT_TENSORS 86→95, FIRST_ISV_TENSOR 68→77. layout_fingerprint_seed() updated to mirror the new tensor names + positions; new LAYOUT_FINGERPRINT_CURRENT = 0xcf3a24b0a1f70057 (was 0xa504d3c2f275b8af). All 93 padded_byte_offset call sites + ancillary index references migrated in lockstep per feedback_no_partial_refactor.md. xavier_init_params_buf extended with init paths for the 13 new tensors: Linear matrices (w_a, w_b, w_residual) get Xavier; LayerNorm γ initialised to 1.0; β + biases stay zero. Trunk-forward / trunk-backward / IQN-trunk / ensemble-diversity-backward callers gated by explicit panic!("Plan 4 Task 2c.3a: trunk param layout migrated to GRN, …") at function entry — BatchedForward::encoder_forward_only, BatchedForward::forward_target_raw, BatchedForward::forward_online_f32, BatchedBackward::backward_full, GpuDqnTrainer::apply_iqn_trunk_gradient, GpuDqnTrainer::apply_ensemble_diversity_backward. This makes accidental runtime execution fail loudly rather than producing silent garbage from running legacy Linear→ReLU→Linear GEMMs against GRN-shaped param tensors. Spectral-norm descriptor (13 matrices) keeps slots [0]/[1] mapped to w_a_h_s1/w_a_h_s2 (shapes match legacy W_s1/W_s2 exactly — the GRN's first Linear has the same shape as the original Linear), so the existing spectral-norm constraint transfers cleanly to the GRN's first Linear; Linear_b / Linear_residual are NOT yet spectral-normed (Task 2c.3c follow-up). Smoke intentionally NOT run — build-clean is the validation; runtime would hit the panic at the first encoder forward (the desired behaviour, no need to verify explicitly). Task 2c.3b swaps in the GRN forward (gpu_grn::GrnBlock::forward) at all panic-gated forward callers in place; Task 2c.3c does the same for backward callers and runs the smoke test. cargo check clean at 11 warnings (baseline preserved); cargo build compiles all 58 cubins. No new module / kernel / ISV slot in this commit — pure structural reshuffle + assertion gates.

Plan 5 Task 5 Phase H (2026-04-26): cublasLt epilogue fusion for the 4 attention forward projections — collapses each cublasLtMatmul + add_bias_f32_kernel pair into a single fused cublasLtMatmul call via CUBLASLT_EPILOGUE_BIAS. Targets the L40S deploy hot-spot where the per-epoch training phase is 99.7% of wall time (25.8 s of 25.9 s); attention forward runs 4 projections (Q, K, V, O) per training step, so this saves 4 kernel launches × N_steps without changing the math. (1) crates/ml/src/cuda_pipeline/gpu_attention.rs::create_attn_gemm_desc extended with an epilogue: Option<cublasLtEpilogue_t> parameter — when Some(BIAS), the descriptor is configured with EPILOGUE = BIAS + BIAS_DATA_TYPE = CUDA_R_32F at creation time; when None, the descriptor stays at EPILOGUE_DEFAULT (the 8 backward dW/dX GEMMs are unchanged). The ShapeKey passed to get_matmul_algo_f32_tf32 is built via with_epilogue when set so the deterministic-algo cache keys on epilogue and produces an algo selection valid for the new descriptor — different epilogues for the same matmul shape get independent algo IDs. (2) New lt_matmul_with_bias_ex helper writes the per-call bias pointer via set_matmul_desc_attribute(BIAS_POINTER, …) immediately before each cublasLtMatmul (lightweight CPU write, no GPU sync; mirrors the existing pattern in batched_forward.rs::sgemm_f32_fused_relu_bias). The 4 forward projection sites in forward(...) switch from the prior lt_matmul_ex(...) + launch_bias_add_ex(...) pair to a single lt_matmul_with_bias_ex(...). (3) Orphans pruned: launch_bias_add_ex and launch_bias_add deleted from gpu_attention.rs (their only callers were the 4 fused-away sites). The shared add_bias_f32_kernel retained — still used by batched_forward.rs::launch_add_bias_f32_raw (VSN Linear_2 logit output, no activation). (4) Determinism preserved: the deterministic-algo cache (cublas_algo_deterministic.rs) handles arbitrary epilogue values transparently — ShapeKey already had an epilogue: i32 field, the with_epilogue constructor stamps it, and the AlgoCheck filter validates the descriptor against the epilogue before caching the algo, so first-call selection runs the full AlgoGetIds → AlgoInit → AlgoCheck loop with the new descriptor and subsequent calls reuse the cached (types, shape) algo. CUBLAS_WORKSPACE_CONFIG=:4096:8 is unchanged; epilogue fusion is bit-deterministic when the algo is fixed. Site #2 (DRELU_BGRAD on the trunk Linear→Bias→ReLU backward) deferred — the trunk's only Linear→ReLU layer that pattern-matches DRELU_BGRAD is the value-FC head (forward sgemm_f32_fused_relu_bias writes save_h_v, backward calls relu_mask on the output of the value-output-layer dX GEMM). Wiring DRELU_BGRAD requires (a) switching value-FC forward from EPILOGUE_RELU_BIAS to EPILOGUE_RELU_AUX_BIAS and allocating a uint8 mask aux buffer of shape [B, VH] packed at 8 elements/byte, (b) plumbing the aux pointer through BatchedForward → BatchedBackward's backward_fc_layer so the value-output-layer dX GEMM can set EPILOGUE_AUX_POINTER + BIAS_POINTER = b_v1_grad, (c) deleting the standalone relu_mask + bias-grad reduce kernels at the value-FC site. The forward-side aux-buffer plumbing crosses three modules and is non-trivial to wire deterministically alongside the GRN trunk's per-block partial-reduction scratch already plumbed through the same chain; landing it in the same commit as Site #1 risks regressing the determinism contract verified by the Phase G smoke. Defer to a follow-up phase that owns the aux-buffer wiring end-to-end. Site #1 alone is the bigger of the two opportunities — 4 launches saved per attention forward × every training step (target + online) — and is fully self-contained inside gpu_attention.rs. Validation: cargo check --workspace clean at 11 warnings (baseline preserved). Smoke (cargo test -p ml --release --lib -- multi_fold_convergence --ignored --nocapture, RTX 3050 Ti) completes all 3 folds; per-fold best train Sharpe within the noise band of the recent Phase G run (F0=2.36, F1=80.82, F2=92.11). Per-epoch wall-time deferred to next L40S deploy nsys profile (cluster log will show STEP_TIMING deltas; the 4-launch saving is the same on H100/L40S as on the local 3050 Ti, and cluster baseline at 25.8 s / epoch makes even a 5% saving worth ~1.3 s/epoch × 30 epochs × 6 folds × 5 seeds ≈ 1170 s total). No new module / kernel / ISV slot / param tensor / Orphan row. Files touched: crates/ml/src/cuda_pipeline/gpu_attention.rs (descriptor factory signature + 4 forward call sites + new lt_matmul_with_bias_ex helper + 2 orphan helpers deleted), docs/dqn-wire-up-audit.md (this entry). No fingerprint change.

Plan 5 Task 5 Phase F (2026-04-26): compile-time fxcache schema fingerprint — closes the L40S deploy-bug class where stale fxcache passed FXCACHE_VERSION validation despite incompatible feature semantics. Root cause of the original failure: extract_ohlcv_features column 0 changed from raw price → log-return without anyone bumping the manually-maintained FXCACHE_VERSION const, so the L40S PVC's older cache loaded clean and the trainer fed raw prices into the aux head expecting log-returns (aux_next_bar_mse=2.587e7). Fix has three pieces: (1) crates/ml/build.rs::emit_feature_schema_hash() runs unconditionally (before the existing CUDA-feature gate so non-CUDA builds also get the env var) and FNV-1a-hashes the raw bytes of the three schema-defining sources — crates/ml/src/features/extraction.rs, crates/ml/src/fxcache.rs, crates/ml-core/src/state_layout.rs — mixing in each file's relative path + length so renames / reorderings also bump the hash. Stable across rust versions and machines (FNV-1a, not std::hash::DefaultHasher). Emits cargo:rustc-env=FEATURE_SCHEMA_HASH=<decimal_u64> plus three cargo:rerun-if-changed= lines. (2) crates/ml/src/fxcache.rs::FEATURE_SCHEMA_HASH consumes the env var via env! + const u64::from_str_radix(_, 10) (const-stable since rust 1.83; workspace MSRV 1.85). FxCacheHeader grows a feature_schema_hash: u64 field; header size 64→72 bytes; wire-format FXCACHE_VERSION bumped 5→6 to flag the layout change. validate() strict-checks the hash alongside magic/version/dims; mismatch bails with a descriptive error pointing at "source files defining feature extraction / state layout / fxcache format have changed since this cache was built." precompute_features.rs:218 already deletes-and-regenerates on any load_fxcache Err, so the existing Argo ensure-fxcache step (and local users) recover automatically. (3) FXCACHE_VERSION doc comment now declares it tracks wire-format changes only — schema-level changes are tracked automatically by FEATURE_SCHEMA_HASH. Removes the manual ritual that broke the L40S deploy. Cost: cosmetic edits (whitespace, comments) to the three schema sources trigger one cache regen on next deploy (~5 min for full L40S dataset, ~40 s for local ES.FUT). Acceptable trade — false negatives (missed schema drift) are not. Validation: cargo check workspace clean at 11 warnings (baseline preserved). Local ES.FUT cache regen confirmed: existing v5 file rejected with "Stale FxCache version: 5 (expected 6). Delete and regenerate.", regenerated v6 cache loads clean on retry. Auto-detection verified: comment-only edit to extraction.rs line 1 changed emitted hash 5046469432341222878 → 7772630163018944575; revert returned the hash deterministically to 5046469432341222878. No new pip/cargo deps (FNV-1a is ~10 LOC stdlib). Files touched: crates/ml/build.rs, crates/ml/src/fxcache.rs, docs/dqn-wire-up-audit.md (this entry). No fingerprint change (LAYOUT_FINGERPRINT_CURRENT untouched — this is fxcache wire-format, not GPU param layout).

Plan 5 Task T5.1 (2026-04-27): Layer 3 gate-differentiation validation test — crates/ml/src/trainers/dqn/smoke_tests/moe_gate_differentiation.rs. Reads the most recent HEALTH_DIAG[N]: aux_moe [util=... ent=...] line from the smoke log and asserts three conditions: (1) no expert utilization < 2% (anti-collapse mechanism working — dead expert would indicate load-balancing aux loss is ineffective or gate temperature too low), (2) no expert utilization > 80% (no winner-takes-all snap-collapse to a single expert), (3) gate entropy < 0.9·ln(8) = 1.871 (gate is meaningfully peakier than uniform). Parser matches the exact tracing::info! format from training_loop.rs:3334: util=v0,v1,...,v7 ent=X.XXX — no brackets around the util values, space-separated from ent=. Compile-time const assert verifies the ISV layout invariant: MOE_EXPERT_UTIL_EMA_BASE + MOE_EXPERT_UTIL_EMA_COUNT <= MOE_GATE_ENTROPY_EMA_INDEX (i.e. 118 + 8 <= 126). Registered in smoke_tests/mod.rs. The test is #[ignore] and requires the smoke to be run first with log capture: cargo test -p ml --lib -- multi_fold_convergence --ignored --nocapture 2>&1 | tee /tmp/moe_phase3_smoke.log. Phase 4 smoke output showed expert 2 at 16.9% with others at ~11.9% — satisfies all three assertions. Failures are findings, not test infrastructure bugs — concrete numbers reported per feedback_kill_runs_on_anomaly_quickly.md. Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §8.3. No new kernel / ISV slot / param tensor. Files touched: crates/ml/src/trainers/dqn/smoke_tests/moe_gate_differentiation.rs (new), crates/ml/src/trainers/dqn/smoke_tests/mod.rs (registration), docs/dqn-wire-up-audit.md (this entry).

Plan 5 Task T5.2 (2026-04-27): Layer 5 architecture-hash backward-incompat test — crates/ml-dqn/tests/architecture_hash_test.rs. Asserts that a synthetic "pre-MoE checkpoint" with a zeroed-out dqn.arch_hash field is rejected by DQNConfig::validate_checkpoint_metadata with a clear architecture-mismatch error. Uses a HashMap<String, String> populated with valid-looking field values except for arch_hash = "000...000" (64 hex zeros — guaranteed to not match any real config). The test calls config.validate_checkpoint_metadata(&Some(fake_metadata)), asserts it returns Err, and asserts the error message contains "Architecture mismatch" or "arch_hash" or "fingerprint" — matching the actual error text from dqn.rs:458. Purely synthetic: no CUDA, no file I/O, no smoke log required. Registered as [[test]] via Cargo's automatic integration test discovery (crates/ml-dqn/tests/ directory). Confirms the no-fallback contract per spec §6.4 — old checkpoints fail loudly, no silent loading of incompatible weights. Files touched: crates/ml-dqn/tests/architecture_hash_test.rs (new), docs/dqn-wire-up-audit.md (this entry). No fingerprint change.

Classification	Count
Wired	88
Partial	10
Orphan (held for follow-up)	3
Ghost	0
OUT-of-DQN-scope	17
Total	116

The 3 remaining Orphan rows are:

• cuda_pipeline/gpu_statistics.rs + statistics_kernel.cu — held for Plan 2 D.2 wire-or-delete decision.
• data_loaders/tlob_loader.rs — held for Plan 2 D.8 TLOB trunk wiring.
• regime_detection/mod.rs — crate-level cleanup task (whole workspace crate, exceeds Task 6 scope).

Orphan Resolution Index

Orphan	Resolution	Status
cuda_pipeline/gpu_statistics.rs + statistics_kernel.cu	Wire into val path or delete if superseded by gpu_monitoring.rs	Plan 2 D.2 — held
data_loaders/streaming_dbn_loader.rs	Reclassified Partial: test consumer at tests/streaming_pipeline_edge_cases.rs	RECLASSIFIED
data_loaders/tlob_loader.rs	Plan 2 D.8 wires TLOB trunk — this loader may be needed then	Plan 2 D.8 — held
training/unified_data_loader.rs	Deleted (zero consumers)	DELETED
training/orchestrator.rs	Deleted (zero consumers)	DELETED
training_pipeline.rs	Reclassified Partial: ProductionTrainingError used via From impl in lib.rs	RECLASSIFIED
inference_validator.rs	Deleted (zero consumers)	DELETED
model_loader_integration.rs	Deleted (zero consumers)	DELETED
paper_trading/mod.rs	Reclassified Partial: test consumer at tests/paper_trading_integration_test.rs	RECLASSIFIED
portfolio_transformer.rs	Deleted (zero consumers)	DELETED
regime_detection/mod.rs	Crate-level follow-up scheduled (separate crate-cleanup task)	CRATE-LEVEL-FOLLOWUP

Mapped pinned buffers promoted to shared module (2026-04-27): moved MappedF32Buffer and MappedI32Buffer from crates/ml/src/trainers/dqn/distributional_q_tests.rs local definitions to a shared crates/ml/src/cuda_pipeline/mapped_pinned.rs module so all kernel test wrappers (Test 0.F + upcoming MoE kernel tests) share one implementation. Adds write_from_slice helper for direct host_ptr writes (no memcpy). Test 0.F bit-identical post-move. Per feedback_no_htod_htoh_only_mapped_pinned.md.

MoE moe_mixture_forward kernel + Rust wrapper (2026-04-27): first MoE CUDA kernel landed. Single-thread-per-(b,c) kernel computes h_s2[b,c] = Σ_k g[b,k]·expert_outputs[k,b,c]. No atomicAdd, capture-friendly. Rust wrapper GpuMoeHead in crates/ml/src/cuda_pipeline/gpu_moe_head.rs loads cubin, exposes test_mixture_forward using mapped pinned buffers exclusively (no HtoD/HtoH per feedback_no_htod_htoh_only_mapped_pinned.md). Unit test in crates/ml/tests/moe_kernels_test.rs verifies CPU reference match within 1e-6 for B=4, K=8, C=256.

RegimeConditionalDQN deletion (2026-04-27): atomic removal per feedback_no_partial_refactor.md. Deleted regime_adx_idx, regime_cusum_idx, regime_adx_threshold, regime_cusum_threshold fields from DQNConfig (both struct definition and Default + aggressive() builder entries). DQNAgentType rewritten as thin wrapper over single DQN directly: no per-regime delegation methods, no multi-head epsilon/epsilon_all calls, no get_trending_head / primary_head indirection, no get_trending_head_memory accessor. The ghost feature get_count_bonuses_branched (previously hardcoded None) now delegates to DQN::get_count_bonuses_branched() — UCB count bonuses reach the GPU action selector for the first time. action.rs caller simplified to direct destructure of fixed arrays (no Option match). serialize_model in mod.rs rewritten to serialize the single DQN branching network directly (no trending__/ranging__/volatile__ prefix namespace). Constructor drops RegimeConditionalDQN::new_on_device and calls DQN::new_on_device directly. curriculum.rs import fixed from crate::dqn::regime_conditional::RegimeType to crate::dqn::RegimeType (re-exported from regime_classifier). Phase 3 MoE (commit a52d99613) provides the regime-conditioned behavior the legacy 3-head architecture pretended to do. Workspace clean: 0 errors across all crates, tests, and examples. Smoke validates no regression: 3/3 folds passed (405s), gate utilization preserved (util=0.119,0.119,0.169,...), HEALTH_DIAG aux_moe line emits, all fold checkpoints written.