diff --git a/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu index 9f177becb..62324d24d 100644 --- a/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu @@ -115,8 +115,9 @@ extern "C" __global__ void backtest_plan_state_isv( * prior direction. In practice trade flips are rare and the plan * gets re-populated on the NEXT Flat→Position transition after * the flip closes. ── */ + float plan_thr = (isv_signals != nullptr) ? isv_signals[ISV_PLAN_THRESHOLD_IDX] : 0.5f; float active_tgt = plan_state[w * 7 + 0]; - if (in_position && active_tgt < 0.5f) { + if (in_position && active_tgt < plan_thr) { const float* pp = plan_params + w * 6; plan_state[w * 7 + 0] = pp[PLAN_PARAM_TARGET_BARS]; /* target_bars */ plan_state[w * 7 + 1] = pp[PLAN_PARAM_PROFIT_TARGET]; /* profit_target */ @@ -146,7 +147,7 @@ extern "C" __global__ void backtest_plan_state_isv( float pisv[6]; /* [0] plan progress [0, 2]: hold_time/target_bars */ - pisv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > 0.5f) + pisv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > plan_thr) ? fminf(hold_time / plan_tgt_bars, 2.0f) : 0.0f; /* [1] PnL vs profit target [bounded 2.0] */ @@ -161,11 +162,11 @@ extern "C" __global__ void backtest_plan_state_isv( pisv[PLAN_ISV_ENTRY_CONVICTION] = plan_conv_ent; /* [4] Conviction drift: current MLP conviction - entry conviction */ - pisv[PLAN_ISV_CONVICTION_DRIFT] = (plan_tgt_bars > 0.5f) + pisv[PLAN_ISV_CONVICTION_DRIFT] = (plan_tgt_bars > plan_thr) ? (plan_params[w * 6 + PLAN_PARAM_CONVICTION] - plan_conv_ent) : 0.0f; /* [5] Regime shift magnitude since entry */ - pisv[PLAN_ISV_REGIME_SHIFT] = (isv_signals != nullptr && plan_tgt_bars > 0.5f) + pisv[PLAN_ISV_REGIME_SHIFT] = (isv_signals != nullptr && plan_tgt_bars > plan_thr) ? fabsf(isv_signals[11] - entry_stab) : 0.0f; #pragma unroll diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 7beadba05..f7673d9c5 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -593,7 +593,11 @@ extern "C" __global__ void experience_state_gather( float plan_stop = ps[PS_PLAN_STOP_LOSS]; float plan_conv = ps[PS_PLAN_CONVICTION]; - plan_isv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > 0.5f) + /* Plan activation threshold from ISV[ISV_PLAN_THRESHOLD_IDX] (Plan 1 Task 16). + * Default 0.5 when ISV not yet available (smoke-scale runs). */ + float plan_thr_sg = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f; + + plan_isv[PLAN_ISV_PROGRESS] = (plan_tgt_bars > plan_thr_sg) ? fminf(f_hold_time / plan_tgt_bars, 2.0f) /* plan progress [0, 2] */ : 0.0f; plan_isv[PLAN_ISV_PNL_VS_TARGET] = (plan_profit > 1e-6f) @@ -606,15 +610,15 @@ extern "C" __global__ void experience_state_gather( /* Conviction drift — current conviction minus entry conviction. * Negative drift = model's thesis is weakening → exit signal. - * Only meaningful when a plan is active (ps[PS_PLAN_TARGET_BARS] > 0.5). */ - plan_isv[PLAN_ISV_CONVICTION_DRIFT] = (plan_params_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f) + * Only meaningful when a plan is active (ps[PS_PLAN_TARGET_BARS] > plan_thr_sg). */ + plan_isv[PLAN_ISV_CONVICTION_DRIFT] = (plan_params_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > plan_thr_sg) ? (plan_params_ptr[i * 6 + PLAN_PARAM_CONVICTION] - ps[PS_PLAN_CONVICTION]) /* conviction drift */ : 0.0f; /* Regime shift since entry — |regime_stability_now - regime_stability_at_entry|. * Large shift = regime changed mid-trade, plan may be invalid. * Entry regime stored in ps[PS_PLAN_ENTRY_REGIME] during plan activation. */ - plan_isv[PLAN_ISV_REGIME_SHIFT] = (isv_signals_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f) + plan_isv[PLAN_ISV_REGIME_SHIFT] = (isv_signals_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > plan_thr_sg) ? fabsf(isv_signals_ptr[11] - ps[PS_PLAN_ENTRY_REGIME]) /* regime shift */ : 0.0f; @@ -1145,7 +1149,8 @@ extern "C" __global__ void experience_action_select( /* Plan direction lock: during active plan, force current direction */ if (portfolio_states != NULL) { int ps_base_plan = i * PORTFOLIO_STRIDE; - int has_plan_active = (portfolio_states[ps_base_plan + PS_PLAN_TARGET_BARS] > 0.5f); + float plan_thr_as = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f; + int has_plan_active = (portfolio_states[ps_base_plan + PS_PLAN_TARGET_BARS] > plan_thr_as); if (has_plan_active) { float cur_pos_plan = portfolio_states[ps_base_plan + PS_POSITION]; if (cur_pos_plan > 0.001f) dir_idx = DIR_LONG; @@ -1771,16 +1776,20 @@ extern "C" __global__ void experience_env_step( /* ── Trade Plan: readiness-gated activation + enforcement ── * Plan head is randomly initialized → garbage plans early in training. - * Gate by model readiness: readiness < 0.5 → plan disabled (model trades freely). - * readiness ≥ 0.5 → plan activates and enforces. Smooth transition. + * Gate by model readiness: readiness < plan_thr → plan disabled (model trades freely). + * readiness >= plan_thr → plan activates and enforces. Smooth transition. * - * Phase 3 unified-env: also gated by exploration_scale > 0.5. Validation + * Phase 3 unified-env: also gated by exploration_scale > plan_thr. Validation * runs with exploration_scale=0 → plan disabled → position is the raw - * Kelly/margin-capped target, matching the backtest kernel. */ + * Kelly/margin-capped target, matching the backtest kernel. + * + * plan_thr reads from ISV[ISV_PLAN_THRESHOLD_IDX] (Plan 1 Task 16 / spec §4.C.6). + * Default 0.5 when ISV not yet available (smoke-scale runs). */ + float plan_thr_env = (isv_signals_ptr != NULL) ? isv_signals_ptr[ISV_PLAN_THRESHOLD_IDX] : 0.5f; float readiness = (readiness_ptr != NULL) ? readiness_ptr[0] : 1.0f; - int plan_enabled = (readiness >= 0.5f + int plan_enabled = (readiness >= plan_thr_env && plan_params_ptr != NULL - && exploration_scale >= 0.5f); + && exploration_scale >= plan_thr_env); /* Plan activation on Flat → Positioned (only when model is ready) */ int was_flat_plan = (fabsf(ps0_f) < 0.001f); @@ -1801,7 +1810,7 @@ extern "C" __global__ void experience_env_step( position *= fmaxf(ps[PS_PLAN_CONVICTION] * readiness, 0.1f); } - int has_plan = (ps[PS_PLAN_TARGET_BARS] > 0.5f); + int has_plan = (ps[PS_PLAN_TARGET_BARS] > plan_thr_env); /* Recursive plan revision (N17): living plan document. * When model is highly confident (readiness ≥ 0.7) and conviction changes @@ -2391,12 +2400,12 @@ extern "C" __global__ void experience_env_step( } /* Plan conviction as reward scaling — ONLY when plan head is mature. - * Before readiness ≥ 0.5, conviction defaults to 1.0 (identity) so the + * Before readiness >= plan_thr_env, conviction defaults to 1.0 (identity) so the * reward signal flows undampened. Without this gate, random init conviction * (~0.1) kills the gradient signal and prevents the model from learning. */ if (plan_params_ptr != NULL && fabsf(reward) > 1e-8f) { float readiness_val = (readiness_ptr != NULL) ? readiness_ptr[0] : 0.0f; - float conviction = (readiness_val >= 0.5f) + float conviction = (readiness_val >= plan_thr_env) ? plan_params_ptr[i * 6 + PLAN_PARAM_CONVICTION] /* [0, 1] from plan head */ : 1.0f; /* identity until plan matures */ reward *= conviction; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index d0e1aede4..e1d3f64b3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -164,17 +164,25 @@ const ISV_NETWORK_DIM: usize = 23; /// to ensure no (sample, branch) pair collapses to zero gradient /// contribution. /// -/// Slots [37..39) carry the layout fingerprint (ISV_LAYOUT_FINGERPRINT_LO_INDEX +/// Slots [39..47) are the DQN v2 Plan 1 expansion (spec §4.C.6, 2026-04-24): +/// [39] EPOCH_IDX_INDEX — CPU writes current epoch index at epoch boundary. +/// [40] TOTAL_EPOCHS_INDEX — CPU writes total epoch count at constructor (static). +/// [41] EPSILON_EFF_INDEX — GPU-written by epsilon adaptive kernel (follow-up task). +/// [42] TAU_EFF_INDEX — GPU-written by tau adaptive kernel (follow-up task). +/// [43] GAMMA_EFF_INDEX — GPU-written by gamma adaptive kernel (follow-up task). +/// [44] KELLY_CAP_EFF_INDEX — GPU-written by kelly_cap adaptive kernel (follow-up task). +/// [45] CQL_ALPHA_INDEX — constructor writes `config.cql_alpha`; CQL formula reads +/// base from here instead of config field (Plan 1 Task 12). +/// [46] PLAN_THRESHOLD_INDEX — constructor writes 0.5f; plan kernels read instead of +/// hardcoded literal (Plan 1 Task 16). +/// +/// Slots [47..49) carry the layout fingerprint (ISV_LAYOUT_FINGERPRINT_LO_INDEX /// and ISV_LAYOUT_FINGERPRINT_HI_INDEX). These are placed at the tail because /// slots [0] and [1] are actively written by `isv_signal_update` (Q-drift and /// gradient-norm EMA respectively). Written by the constructor; checked at /// checkpoint load. Fail-fast only — no migration path exists. See spec §4.A.2 /// and `LAYOUT_FINGERPRINT_CURRENT` for the structural-hash design rationale. -/// -/// Extending past 39 requires updating the pinned allocation in the -/// constructor and any kernel that takes an ISV pointer expecting a -/// specific length. -const ISV_TOTAL_DIM: usize = 39; +const ISV_TOTAL_DIM: usize = 49; /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -263,7 +271,36 @@ pub const GRAD_SCALE_LIMIT_INDEX: usize = 35; /// `feedback_isv_for_adaptive_bounds.md` carve-out. pub const IQL_BRANCH_SCALE_FLOOR_INDEX: usize = 36; -/// ISV slot [37] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +/// ISV slot [39] — current epoch index. CPU writes at each epoch boundary. +/// Zero at construction; updated by the epoch loop in a follow-up task. +pub const EPOCH_IDX_INDEX: usize = 39; +/// ISV slot [40] — total epoch count for this training run. CPU writes at +/// constructor from `config`; constant for the run's lifetime. +pub const TOTAL_EPOCHS_INDEX: usize = 40; +/// ISV slot [41] — effective epsilon written by the GPU epsilon adaptive +/// kernel (follow-up task). Starts at 0; kernel fills it each step. +pub const EPSILON_EFF_INDEX: usize = 41; +/// ISV slot [42] — effective tau written by the GPU tau adaptive kernel +/// (follow-up task). Starts at 0; kernel fills it each step. +pub const TAU_EFF_INDEX: usize = 42; +/// ISV slot [43] — effective gamma written by the GPU gamma adaptive kernel +/// (follow-up task). Starts at 0; kernel fills it each step. +pub const GAMMA_EFF_INDEX: usize = 43; +/// ISV slot [44] — effective Kelly cap written by the GPU kelly_cap adaptive +/// kernel (follow-up task). Starts at 0; kernel fills it each step. +pub const KELLY_CAP_EFF_INDEX: usize = 44; +/// ISV slot [45] — CQL pessimism base coefficient. Constructor writes +/// `config.cql_alpha`; the CQL launch formula reads from here instead of +/// the config field (Plan 1 Task 12 / spec §4.C.6). Plan 3 B.3 may later +/// make this reactive by writing an adaptive value each epoch. +pub const CQL_ALPHA_INDEX: usize = 45; +/// ISV slot [46] — plan-MLP activation threshold. Constructor writes 0.5f; +/// plan kernels (`experience_kernels.cu`, `backtest_plan_kernel.cu`) read +/// from here instead of the hardcoded literal (Plan 1 Task 16 / spec §4.C.6). +/// Plan 3 B.4 may later make this reactive. +pub const PLAN_THRESHOLD_INDEX: usize = 46; + +/// ISV slot [47] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// /// Design note: this is NOT a version number. There is no ordered version space. /// The fingerprint is a structural hash that automatically changes whenever any @@ -274,9 +311,9 @@ pub const IQL_BRANCH_SCALE_FLOOR_INDEX: usize = 36; /// /// Placement at the tail (not head) is mandated by slots [0] and [1] being /// actively written by `isv_signal_update` (Q-drift and gradient-norm EMA). -pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 37; -/// ISV slot [38] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). -pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 38; +pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 47; +/// ISV slot [48] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 48; /// Canonical alias for the fingerprint slot (the low half). pub const ISV_LAYOUT_FINGERPRINT_INDEX: usize = ISV_LAYOUT_FINGERPRINT_LO_INDEX; @@ -327,8 +364,11 @@ const fn layout_fingerprint_seed() -> &'static [u8] { V_CENTER_ORD=27;V_HALF_ORD=28;V_CENTER_URG=29;V_HALF_URG=30;\ GRAD_NORM_TARGET_DIR=31;GRAD_NORM_TARGET_MAG=32;GRAD_NORM_TARGET_ORD=33;GRAD_NORM_TARGET_URG=34;\ GRAD_SCALE_LIMIT=35;IQL_BRANCH_SCALE_FLOOR=36;\ - ISV_LAYOUT_FINGERPRINT_LO=37;ISV_LAYOUT_FINGERPRINT_HI=38;\ - ISV_TOTAL_DIM=39" + EPOCH_IDX=39;TOTAL_EPOCHS=40;\ + EPSILON_EFF=41;TAU_EFF=42;GAMMA_EFF=43;KELLY_CAP_EFF=44;\ + CQL_ALPHA=45;PLAN_THRESHOLD=46;\ + ISV_LAYOUT_FINGERPRINT_LO=47;ISV_LAYOUT_FINGERPRINT_HI=48;\ + ISV_TOTAL_DIM=49" } /// Compile-time layout fingerprint. Burned into the binary at build time. @@ -564,6 +604,11 @@ pub struct GpuDqnTrainConfig { /// market features from portfolio+MTF+OFI features in the state vector. /// OFI features (8 dims, when enabled) bypass the bottleneck via portfolio_dim. pub market_dim: usize, + /// Total number of training epochs for this run. Written to ISV[TOTAL_EPOCHS_INDEX] + /// at constructor time so GPU kernels can compute epoch-fraction progress without + /// a CPU→GPU round-trip. Default 0 means "unknown / single fold"; set this from + /// `DQNHyperparameters::epochs` at construction. + pub total_epochs: usize, } impl Default for GpuDqnTrainConfig { @@ -610,6 +655,7 @@ impl Default for GpuDqnTrainConfig { causal_intervention_interval: 50, bottleneck_dim: 16, market_dim: 42, // Default: 42 base features. Overridden to 50 when OFI (MBP-10) enabled. + total_epochs: 0, } } } @@ -2024,11 +2070,12 @@ pub struct GpuDqnTrainer { q_dir_bin_means_reduce_kernel: CudaFunction, // ── ISV core buffers (pinned device-mapped, GPU read/write) ── - // isv_signals_pinned is sized for ISV_TOTAL_DIM (39) — the extra slots - // slots [23..31] carry per-branch Q-support centres/half-widths. The - // network (w_isv_fc1) and history rotation still consume only the first + // isv_signals_pinned is sized for ISV_TOTAL_DIM (49) — the extra slots + // slots [23..31] carry per-branch Q-support centres/half-widths; [39..49) + // are the Plan 1 C.6 expansion (static configs + pre-allocated GPU slots). + // The network (w_isv_fc1) and history rotation still consume only the first // ISV_NETWORK_DIM (23) slots; the tail is broadcast-bus scratchpad only. - isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM = 39] + isv_signals_pinned: *mut f32, // [ISV_TOTAL_DIM = 49] isv_signals_dev_ptr: u64, isv_history_pinned: *mut f32, // [ISV_K * 12 = 48] — history rotates slots [0..11] only isv_history_dev_ptr: u64, @@ -5858,7 +5905,15 @@ impl GpuDqnTrainer { Some(k) => k.clone(), None => return Ok(false), }; - if self.config.cql_alpha <= 0.0 { + // Read cql_alpha base from ISV[CQL_ALPHA_INDEX] (Plan 1 Task 12 / spec §4.C.6). + // When ISV is null (smoke scale without ISV warmup), fall back to the config + // field so smoke tests that skip ISV bootstrap still work. + let cql_alpha_base = if self.isv_signals_pinned.is_null() { + self.config.cql_alpha + } else { + unsafe { *self.isv_signals_pinned.add(CQL_ALPHA_INDEX) } + }; + if cql_alpha_base <= 0.0 { return Ok(false); } @@ -5908,14 +5963,14 @@ impl GpuDqnTrainer { should never hit this path." ); }); - self.config.cql_alpha + cql_alpha_base } else { let (health, regime_stability) = unsafe { let h = (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0); let s = (*self.isv_signals_pinned.add(11)).clamp(0.0, 1.0); (h, s) }; - self.config.cql_alpha * (1.0 - regime_stability) * health + cql_alpha_base * (1.0 - regime_stability) * health }; // Cache for logging via FusedTrainingCtx / DQNTrainer. @@ -6160,7 +6215,16 @@ impl GpuDqnTrainer { /// Whether CQL is enabled and the kernel was compiled. pub fn has_cql(&self) -> bool { - self.cql_logit_grad_kernel.is_some() && self.config.cql_alpha > 0.0 + if !self.cql_logit_grad_kernel.is_some() { + return false; + } + // Read base from ISV[CQL_ALPHA_INDEX] when available; fall back to config. + let base = if self.isv_signals_pinned.is_null() { + self.config.cql_alpha + } else { + unsafe { *self.isv_signals_pinned.add(CQL_ALPHA_INDEX) } + }; + base > 0.0 } /// F8/G5: Scale the C51-contributed portion of `grad_buf` by the adaptive c51_budget. @@ -8167,10 +8231,10 @@ impl GpuDqnTrainer { }; // ── ISV core buffers (pinned device-mapped) ────────────────────── - // Allocation uses `ISV_TOTAL_DIM` so all scratchpad slots [23..39) have + // Allocation uses `ISV_TOTAL_DIM` so all scratchpad slots [23..49) have // backing storage. The network-facing path still consumes only the first // `ISV_NETWORK_DIM` slots — the scratchpad tail is invisible to `w_isv_fc1`. - // Slots [37..39) hold the layout fingerprint (tail placement because [0..2) + // Slots [47..49) hold the layout fingerprint (tail placement because [0..2) // are actively written by `isv_signal_update`). let (isv_signals_pinned, isv_signals_dev_ptr) = { let num_bytes = ISV_TOTAL_DIM * std::mem::size_of::(); @@ -8212,7 +8276,14 @@ impl GpuDqnTrainer { // IQL readiness→1 and the losing branch has near-zero // advantage. Safety bound, not tuning. *sig_ptr.add(IQL_BRANCH_SCALE_FLOOR_INDEX) = 0.1_f32; - // Layout fingerprint (ISV[37..39)). Compile-time structural hash of + // Plan 1 C.6 static-config slots (2026-04-24): + // EPOCH_IDX starts at 0; epoch loop updates in a follow-up task. + // EPSILON_EFF / TAU_EFF / GAMMA_EFF / KELLY_CAP_EFF start at 0; + // GPU kernels in follow-up tasks fill them. + *sig_ptr.add(TOTAL_EPOCHS_INDEX) = config.total_epochs as f32; + *sig_ptr.add(CQL_ALPHA_INDEX) = config.cql_alpha; + *sig_ptr.add(PLAN_THRESHOLD_INDEX) = 0.5_f32; + // Layout fingerprint (ISV[47..49)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so // no precision is lost. See LAYOUT_FINGERPRINT_CURRENT docs for why @@ -9880,10 +9951,12 @@ impl GpuDqnTrainer { // ── A3: LearningHealth signal writers / readers ─────────────────────── /// Write a scalar into the ISV pinned buffer at the given index. - /// No-op if the pinned pointer is null or index >= ISV_DIM. + /// No-op if the pinned pointer is null or index >= ISV_TOTAL_DIM. + /// Accepts any slot in the full bus [0..ISV_TOTAL_DIM), including the + /// scratchpad tail [ISV_NETWORK_DIM..ISV_TOTAL_DIM). pub fn write_isv_signal_at(&self, index: usize, value: f32) { if self.isv_signals_pinned.is_null() { return; } - if index >= ISV_DIM { return; } + if index >= ISV_TOTAL_DIM { return; } unsafe { *self.isv_signals_pinned.add(index) = value; } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index a7a2ce12b..ed3f62cad 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -124,6 +124,16 @@ #define MAG_FULL 2 // 1.00× max_position #define NUM_MAGNITUDES 3 +// ──────────────────────────────────────────────────────────────────────────── +// ISV bus slot indices referenced by kernels that receive the ISV pointer. +// Kept in state_layout.cuh because both experience_kernels.cu and +// backtest_plan_kernel.cu include this header and need these indices. +// The authoritative named constants live in gpu_dqn_trainer.rs; the values +// here must match exactly. Any renumber requires updating both files in the +// same commit (per feedback_no_partial_refactor.md). +// ──────────────────────────────────────────────────────────────────────────── +#define ISV_PLAN_THRESHOLD_IDX 46 // == PLAN_THRESHOLD_INDEX — plan activation threshold + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 4d4a33e21..9c6274893 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -400,6 +400,7 @@ impl FusedTrainingCtx { causal_intervention_interval: hyperparams.causal_intervention_interval, bottleneck_dim: hyperparams.bottleneck_dim, market_dim: 42, // Always 42 base market features — OFI features bypass bottleneck via portfolio_dim + total_epochs: hyperparams.epochs, }; // Create weight set pointer views AFTER GpuDqnTrainer is constructed below. diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 61463673e..6de186b28 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -144,13 +144,54 @@ impl StateResetRegistry { RegistryEntry { name: "ISV_LAYOUT_FINGERPRINT", category: ResetCategory::SchemaContract, - description: "ISV[37..39) — compile-time structural hash of slot layout; fail-fast only, no migration path", + description: "ISV[47..49) — compile-time structural hash of slot layout; fail-fast only, no migration path", }, RegistryEntry { name: "isv_iql_branch_scale_floor", category: ResetCategory::SchemaContract, description: "ISV[36] — IQL branch_scales floor (safety bound, bootstrap-only)", }, + RegistryEntry { + name: "isv_total_epochs", + category: ResetCategory::SchemaContract, + description: "ISV[TOTAL_EPOCHS_INDEX=40] — total epoch count for this run; constant after construction", + }, + RegistryEntry { + name: "isv_cql_alpha", + category: ResetCategory::SchemaContract, + description: "ISV[CQL_ALPHA_INDEX=45] — CQL pessimism base coefficient; Plan 3 B.3 may make reactive", + }, + RegistryEntry { + name: "isv_plan_threshold", + category: ResetCategory::SchemaContract, + description: "ISV[PLAN_THRESHOLD_INDEX=46] — plan-MLP activation threshold; Plan 3 B.4 may make reactive", + }, + // ───── Fold-reset state (pre-allocated GPU-written slots) ─── + RegistryEntry { + name: "isv_epoch_idx", + category: ResetCategory::FoldReset, + description: "ISV[EPOCH_IDX_INDEX=39] — current epoch index; CPU writes at epoch boundary (follow-up task)", + }, + RegistryEntry { + name: "isv_epsilon_eff", + category: ResetCategory::FoldReset, + description: "ISV[EPSILON_EFF_INDEX=41] — effective epsilon; GPU epsilon-adaptive kernel fills (follow-up task)", + }, + RegistryEntry { + name: "isv_tau_eff", + category: ResetCategory::FoldReset, + description: "ISV[TAU_EFF_INDEX=42] — effective tau; GPU tau-adaptive kernel fills (follow-up task)", + }, + RegistryEntry { + name: "isv_gamma_eff", + category: ResetCategory::FoldReset, + description: "ISV[GAMMA_EFF_INDEX=43] — effective gamma; GPU gamma-adaptive kernel fills (follow-up task)", + }, + RegistryEntry { + name: "isv_kelly_cap_eff", + category: ResetCategory::FoldReset, + description: "ISV[KELLY_CAP_EFF_INDEX=44] — effective Kelly cap; GPU kelly-cap-adaptive kernel fills (follow-up task)", + }, ]; Self { entries } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 4316364c3..78310d6cd 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -219,6 +219,8 @@ 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 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. + | Classification | Count | |---|---| | Wired | 75 | diff --git a/docs/isv-slots.md b/docs/isv-slots.md index 24a8b3c11..d65760ef1 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -8,13 +8,14 @@ fail-fast; no migration functions exist. Backward compat is structurally unwritable (there is no ordered version space to pair migrations against). See spec §4.A.2. -**Tail placement rationale:** The fingerprint occupies the tail (`[37..39)`) rather +**Tail placement rationale:** The fingerprint occupies the tail (`[47..49)`) rather than the head because `isv_signals[0]` and `isv_signals[1]` are actively written by `isv_signal_update` (Q-drift EMA and gradient-norm EMA respectively). Inserting at the head would displace those live signals and require shifting every upstream -literal in `experience_kernels.cu`. +literal in `experience_kernels.cu`. The fingerprint moves to the new tail each time +new slots are appended to the bus. -**Current `ISV_TOTAL_DIM`:** 39. Post-full DQN v2 rollout: 72. +**Current `ISV_TOTAL_DIM`:** 49 (Plan 1 Tasks 12/15/16 + pre-allocation for Tasks 9-11/13/14). Post-full DQN v2 rollout: 72. | Index | Name constant | Type | Producer | Consumers | Reset-category | Notes | |---|---|---|---|---|---|---| @@ -38,6 +39,15 @@ literal in `experience_kernels.cu`. | [31..35) | `GRAD_NORM_TARGET_*_INDEX` | f32 | grad_balance_isv_update | branch_grad_rescale | SoftReset(decay_bars=500) | Per-branch grad-norm target | | [35] | `GRAD_SCALE_LIMIT_INDEX` | f32 | grad_balance_isv_update | branch_grad_rescale | SoftReset(decay_bars=500) | Scale clamp limit | | [36] | `IQL_BRANCH_SCALE_FLOOR_INDEX` | f32 | construct (static) | iql_per_branch_advantage | SchemaContract | Per-sample branch_scales floor; safety bound | -| [37] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. | -| [38] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. | -| [39..72) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 2-5 | +| [37..39) | (gap — fingerprint moved to tail) | — | — | — | — | Previously [37..39); fingerprint promoted to [47..49) by Plan 1 C.6 expansion. Slots unused; zero-filled. | +| [39] | `EPOCH_IDX_INDEX` | f32 (int cast) | CPU epoch-loop (follow-up task) | GPU adaptive kernels | FoldReset | Current epoch index; 0 at construction, CPU writes at each epoch boundary | +| [40] | `TOTAL_EPOCHS_INDEX` | f32 (int cast) | construct (CPU static) | GPU adaptive kernels | SchemaContract | Total epoch count for this run; written at construction from `config.total_epochs` | +| [41] | `EPSILON_EFF_INDEX` | f32 | GPU epsilon-adaptive kernel (follow-up) | epsilon-greedy action select | FoldReset | Effective epsilon; 0 at construction; GPU fills each step | +| [42] | `TAU_EFF_INDEX` | f32 | GPU tau-adaptive kernel (follow-up) | target-net Polyak update | FoldReset | Effective tau; 0 at construction; GPU fills each step | +| [43] | `GAMMA_EFF_INDEX` | f32 | GPU gamma-adaptive kernel (follow-up) | Bellman target kernel | FoldReset | Effective gamma; 0 at construction; GPU fills each step | +| [44] | `KELLY_CAP_EFF_INDEX` | f32 | GPU kelly-cap-adaptive kernel (follow-up) | experience_env_step | FoldReset | Effective Kelly cap; 0 at construction; GPU fills each step | +| [45] | `CQL_ALPHA_INDEX` | f32 | construct (CPU static) | CQL adaptive formula in Rust | SchemaContract | CQL pessimism base coefficient; written from `config.cql_alpha`; read in `compute_cql_logit_gradients` | +| [46] | `PLAN_THRESHOLD_INDEX` | f32 | construct (CPU static) | `experience_kernels.cu`, `backtest_plan_kernel.cu` | SchemaContract | Plan-MLP activation threshold; 0.5f at construction; read via `ISV_PLAN_THRESHOLD_IDX` in both kernels | +| [47] | `ISV_LAYOUT_FINGERPRINT_LO_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | Low 32 bits of u64 FNV-1a structural hash. Fail-fast on mismatch — NOT a version number, no migration path. | +| [48] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. | +| [49..72) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 2-5 |