diff --git a/config/training/dqn-smoketest.toml b/config/training/dqn-smoketest.toml index f343a6a49..48eb48301 100644 --- a/config/training/dqn-smoketest.toml +++ b/config/training/dqn-smoketest.toml @@ -49,6 +49,12 @@ noise_sigma = 0.1 buffer_size = 256 min_replay_size = 64 regime_replay_decay = 0.3 +# Plan 3 Task 8 B.3: smoke override. Production default 100_000; smoke runs +# only ~200 samples per `collect_experiences_gpu` (4 episodes × 50 timesteps) +# × 5 epochs × 3 folds ≈ 3000 total. Setting to 1_000 lets the seed phase +# complete mid-fold so the seed→network transition is observable in the smoke +# (HEALTH_DIAG cql_alpha.eff visibly ramps from 0 toward config.cql_alpha). +replay_seed_steps = 1000 [early_stopping] enabled = true diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f29f53fe4..438aff1ee 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -103,6 +103,21 @@ fn main() { // escape amplifier (producer for ISV[78]+ISV[79]). Adaptive amp // multiplies B.1 opp_cost and B.2 bonus consumers. "state_kl_divergence_kernel.cu", + // B.3 Plan 3 Task 8: GPU-only seeded warm-start scripted policies. + // Replaces network-Q action source during the seed phase. 4 policies + // (uniform/momentum/mean-rev/vwap-deviation) deterministically mixed + // 40/20/20/20 by `i % 5`. Driven by ISV[SEED_STEPS_DONE/TARGET] + // dispatch at the launcher boundary. + "scripted_policy_kernel.cu", + // B.3 Plan 3 Task 8: per-collect_experiences seed-phase progress + // counter. Increments ISV[SEED_STEPS_DONE], EMAs the derived + // seed-fraction into ISV[SEED_FRAC_EMA] (consumed by Task 9). + "seed_step_counter_update_kernel.cu", + // C.5 Plan 3 Task 9: CQL α ramp coupled to ISV[SEED_FRAC_EMA]. + // EMAs ISV[CQL_ALPHA_INDEX=48] toward `final × (1 - seed_frac)`. + // Producer-only upgrade for slot 48; CQL gradient kernel consumer + // (already reading ISV[48] per Plan 1 Task 12) unchanged. + "cql_alpha_seed_update_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/cql_alpha_seed_update_kernel.cu b/crates/ml/src/cuda_pipeline/cql_alpha_seed_update_kernel.cu new file mode 100644 index 000000000..b592e508f --- /dev/null +++ b/crates/ml/src/cuda_pipeline/cql_alpha_seed_update_kernel.cu @@ -0,0 +1,50 @@ +/* cql_alpha_seed_update — GPU-driven CQL α ramp coupled to seed-phase decay. + * + * Plan 3 Task 9 (C.5 CQL α seed-coupled ramp). + * + * Single-thread cold-path kernel launched per epoch boundary alongside the + * other Plan 3 ISV producers. Reads ISV[SEED_FRAC_EMA_INDEX] (Task 8) and + * computes + * + * target = cql_alpha_final × max(0, 1 - seed_frac) + * + * EMAs target into ISV[CQL_ALPHA_INDEX=48] using the adaptive-rate + * convention (α = α_base × (1 + 0.5 × |clamp(sharpe, ±2)|)). + * + * Semantics: + * - During seed phase (seed_frac ≈ 1): target → 0, so CQL α decays toward + * 0 — no pessimism on exploration data populated by scripted policies. + * - As seed_frac → 0: target → cql_alpha_final, CQL α ramps to its + * config value, and the loss kernel sees full pessimism on + * network-driven trajectories. + * + * Constructor still cold-starts ISV[CQL_ALPHA_INDEX=48] to config.cql_alpha; + * this kernel begins overwriting that value at the first epoch boundary. + * The CQL gradient kernel reads ISV[48] directly (no consumer change + * needed — Plan 1 Task 12 / Bug #7 fix `90e1e3dbb` already wired the read). + */ + +#include + +extern "C" __global__ void cql_alpha_seed_update( + float* __restrict__ isv_out, /* ISV bus [ISV_TOTAL_DIM] */ + int cql_alpha_idx, /* = CQL_ALPHA_INDEX = 48 */ + int seed_frac_idx, /* = SEED_FRAC_EMA_INDEX = 84 */ + int sharpe_ema_idx, /* = SHARPE_EMA_INDEX = 22 */ + float alpha_base, /* base EMA rate (e.g. 0.05) */ + float cql_alpha_final /* config.cql_alpha — full pessimism target */ +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float seed_frac = isv_out[seed_frac_idx]; + /* During seed phase (frac=1) → target=0 (no pessimism on exploration + * data). As frac → 0 → target = cql_alpha_final (full pessimism on + * network-driven trajectories). */ + const float target = cql_alpha_final * fmaxf(0.0f, 1.0f - seed_frac); + + /* Adaptive α: higher-Sharpe epochs ramp faster. */ + const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx])); + const float alpha = alpha_base * (1.0f + 0.5f * fabsf(sharpe)); + const float prev = isv_out[cql_alpha_idx]; + isv_out[cql_alpha_idx] = (1.0f - alpha) * prev + alpha * target; +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 1ba2342bc..50fa84325 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -97,6 +97,16 @@ pub(crate) static PLAN_THRESHOLD_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(en /// Plan 3 Task 7 C.3: per-epoch train-vs-val state-distribution KL EMA + /// adaptive Flat-trap escape amplifier (consumed by B.1/B.2 multipliers). pub(crate) static STATE_KL_DIVERGENCE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/state_kl_divergence_kernel.cubin")); +/// Plan 3 Task 8 B.3: GPU-only seeded warm-start scripted policies. +/// Replaces the network-Q action source during the seed phase. 4 policies +/// (uniform/momentum/mean-rev/vwap-deviation) deterministically mixed. +pub(crate) static SCRIPTED_POLICY_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/scripted_policy_kernel.cubin")); +/// Plan 3 Task 8 B.3: per-collect_experiences seed-phase progress counter +/// (ISV[83]) + adaptive seed-fraction EMA (ISV[84]). +pub(crate) static SEED_STEP_COUNTER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/seed_step_counter_update_kernel.cubin")); +/// Plan 3 Task 9 C.5: CQL α ramp coupled to ISV[SEED_FRAC_EMA]. +/// EMAs ISV[CQL_ALPHA_INDEX=48] toward `final × (1 - seed_frac)`. +pub(crate) static CQL_ALPHA_SEED_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/cql_alpha_seed_update_kernel.cubin")); /// Mamba2 temporal scan configuration. const MAMBA2_HISTORY_K: usize = 8; // Rolling history length @@ -241,11 +251,19 @@ const ISV_NETWORK_DIM: usize = 23; /// [79] STATE_KL_AMPLIFICATION_INDEX — Flat-trap escape multiplier ∈ [1, 2] /// smoothly tracking the trailing-EMA-of-self KL ratio. Multiplied into /// B.1 opp_cost and B.2 bonus consumers in `experience_kernels.cu`. -/// [80] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint (shifted from 76). -/// [81] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint (shifted from 77). +/// [82] SEED_STEPS_TARGET_INDEX — Plan 3 Task 8 B.3: config replay_seed_steps target. +/// CPU constructor write from `config.replay_seed_steps`; constant for the run. +/// [83] SEED_STEPS_DONE_INDEX — Plan 3 Task 8 B.3: cumulative seed-phase steps +/// completed. GPU-incremented by `seed_step_counter_update` kernel after every +/// `collect_experiences_gpu`. Capped at TARGET to avoid overflow. +/// [84] SEED_FRAC_EMA_INDEX — Plan 3 Task 8 B.3: adaptive EMA of +/// `max(0, 1 - DONE/TARGET)`. Smoothly decays 1→0 as the seed phase completes. +/// Cold-start 1.0 (fully in seed phase). Consumed by Task 9 CQL α ramp kernel. +/// [85] ISV_LAYOUT_FINGERPRINT_LO_INDEX — low 32 bits of layout fingerprint (shifted from 80). +/// [86] ISV_LAYOUT_FINGERPRINT_HI_INDEX — high 32 bits of layout fingerprint (shifted from 81). /// 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 structural-hash rationale. -const ISV_TOTAL_DIM: usize = 82; +const ISV_TOTAL_DIM: usize = 87; /// 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). @@ -459,7 +477,23 @@ pub const STATE_KL_TRAIN_VAL_EMA_INDEX: usize = 78; /// `fmaxf(1.0, …)` to no-op against cold-start 0. pub const STATE_KL_AMPLIFICATION_INDEX: usize = 79; -/// ISV slot [80] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +/// ISV slot [82] — Plan 3 Task 8 B.3: target number of seed-phase +/// experience-collection steps. CPU-written at constructor from +/// `config.replay_seed_steps` (default 100k). SchemaContract — never reset +/// after construction. +pub const SEED_STEPS_TARGET_INDEX: usize = 82; +/// ISV slot [83] — Plan 3 Task 8 B.3: cumulative seed-phase steps completed. +/// GPU-incremented by `seed_step_counter_update` per `collect_experiences_gpu` +/// call while DONE < TARGET. Capped at TARGET. Reset to 0 at fold boundary +/// (FoldReset). +pub const SEED_STEPS_DONE_INDEX: usize = 83; +/// ISV slot [84] — Plan 3 Task 8 B.3: seed-fraction adaptive EMA of +/// `max(0, 1 - DONE/TARGET)`. Cold-start 1.0 (fully in seed phase), +/// smoothly decays toward 0 as DONE approaches TARGET. Consumed by Task 9 +/// CQL α ramp kernel. FoldReset (re-bootstrap to 1.0 each fold). +pub const SEED_FRAC_EMA_INDEX: usize = 84; + +/// ISV slot [85] — 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 @@ -471,12 +505,13 @@ pub const STATE_KL_AMPLIFICATION_INDEX: usize = 79; /// 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). /// Shifted 61→69 in Plan 3 Task 1, 69→73 in Plan 3 Task 3, 73→76 in -/// Plan 3 Task 4 B.4, then 76→80 in Plan 3 Task 7 C.3. -pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 80; -/// ISV slot [81] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). +/// Plan 3 Task 4 B.4, 76→80 in Plan 3 Task 7 C.3, then 80→85 in +/// Plan 3 Task 8 B.3 (3 new seed-phase slots tail-appended). +pub const ISV_LAYOUT_FINGERPRINT_LO_INDEX: usize = 85; +/// ISV slot [86] — high 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// Shifted 62→70 in Plan 3 Task 1, 70→74 in Plan 3 Task 3, 74→77 in Plan 3 Task 4 B.4, -/// then 77→81 in Plan 3 Task 7 C.3. -pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 81; +/// 77→81 in Plan 3 Task 7 C.3, then 81→86 in Plan 3 Task 8 B.3. +pub const ISV_LAYOUT_FINGERPRINT_HI_INDEX: usize = 86; /// Canonical alias for the fingerprint slot (the low half). pub const ISV_LAYOUT_FINGERPRINT_INDEX: usize = ISV_LAYOUT_FINGERPRINT_LO_INDEX; @@ -539,8 +574,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] { TRADE_ATTEMPT_RATE_EMA=71;TRADE_TARGET_RATE=72;\ READINESS_EMA=75;\ STATE_KL_EMA=78;STATE_KL_AMP=79;\ - ISV_LAYOUT_FINGERPRINT_LO=80;ISV_LAYOUT_FINGERPRINT_HI=81;\ - ISV_TOTAL_DIM=82" + SEED_STEPS_TARGET=82;SEED_STEPS_DONE=83;SEED_FRAC_EMA=84;\ + ISV_LAYOUT_FINGERPRINT_LO=85;ISV_LAYOUT_FINGERPRINT_HI=86;\ + ISV_TOTAL_DIM=87" } /// Compile-time layout fingerprint. Burned into the binary at build time. @@ -781,6 +817,17 @@ pub struct GpuDqnTrainConfig { /// a CPU→GPU round-trip. Default 0 means "unknown / single fold"; set this from /// `DQNHyperparameters::epochs` at construction. pub total_epochs: usize, + + /// Plan 3 Task 8 B.3: number of experience-collection samples that should be + /// drawn from the GPU scripted-policy mixture before handing action selection + /// back to the network. Written to ISV[SEED_STEPS_TARGET_INDEX=82] at + /// constructor time. Default 100_000 (≈ first ~5 epochs at production + /// batch_size×timesteps); set to 0 to disable the seed phase entirely + /// (the dispatch branch in the launcher will see DONE >= TARGET on the + /// very first call and route to the network kernel). Smoke configs may + /// override to a smaller value (e.g. 10_000) so the seed→network transition + /// is observable inside the smoke run. + pub replay_seed_steps: u32, } impl Default for GpuDqnTrainConfig { @@ -828,6 +875,7 @@ impl Default for GpuDqnTrainConfig { bottleneck_dim: 16, market_dim: 42, // Default: 42 base features. Overridden to 50 when OFI (MBP-10) enabled. total_epochs: 0, + replay_seed_steps: 100_000, } } } @@ -8521,6 +8569,18 @@ impl GpuDqnTrainer { * neutral element until the kernel populates a real value. */ *sig_ptr.add(STATE_KL_TRAIN_VAL_EMA_INDEX) = 0.0_f32; *sig_ptr.add(STATE_KL_AMPLIFICATION_INDEX) = 1.0_f32; + /* B.3 Plan 3 Task 8: GPU-only replay seed warm-start. + * Cold-start the seed-phase progress slots. TARGET is the + * static config (SchemaContract); DONE starts at 0 and the + * GPU `seed_step_counter_update` kernel increments it after + * every `collect_experiences_gpu`. SEED_FRAC_EMA cold-starts + * at 1.0 (fully in seed phase) so Task 9's CQL ramp sees + * `target = final × max(0, 1 - 1.0) = 0` until the seed + * phase actually decays — preserving the existing CQL=0 + * pre-warmup behaviour byte-for-byte. */ + *sig_ptr.add(SEED_STEPS_TARGET_INDEX) = config.replay_seed_steps as f32; + *sig_ptr.add(SEED_STEPS_DONE_INDEX) = 0.0_f32; + *sig_ptr.add(SEED_FRAC_EMA_INDEX) = 1.0_f32; // Plan 2 C.1 Q-quantile bootstrap (2026-04-24): // q_p05 = v_min, q_p95 = v_max — matches the cold-start atom range // so the first update_eval_v_range call reads meaningful bounds. diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 4a463c56f..7d93a4692 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -775,6 +775,27 @@ pub struct GpuExperienceCollector { /// `launch_state_kl_inplace`. state_kl_kernel: CudaFunction, + /// B.3 Plan 3 Task 8: scripted_policy_select GPU kernel. + /// Replaces the network-Q action source during the seed phase. 4 policies + /// (uniform/momentum/mean-rev/vwap-deviation) deterministically mixed + /// 40/20/20/20 by `i % 5`. Per-thread one sample. Launched from + /// `launch_scripted_policy_select` only when ISV[SEED_STEPS_DONE] < + /// ISV[SEED_STEPS_TARGET]. + scripted_policy_kernel: CudaFunction, + + /// B.3 Plan 3 Task 8: seed_step_counter_update GPU kernel. + /// Single-block single-thread cold-path. Increments + /// ISV[SEED_STEPS_DONE_INDEX] by `n_samples`, EMAs the derived + /// `max(0, 1 - DONE/TARGET)` into ISV[SEED_FRAC_EMA_INDEX]. Launched + /// after every `collect_experiences_gpu`. + seed_step_counter_kernel: CudaFunction, + + /// C.5 Plan 3 Task 9: cql_alpha_seed_update GPU kernel. + /// Single-block single-thread cold-path. EMAs ISV[CQL_ALPHA_INDEX=48] + /// toward `cql_alpha_final × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX])`. + /// Launched alongside `seed_step_counter_update`. + cql_alpha_seed_update_kernel: CudaFunction, + /// ISV signals dev_ptr from trainer — for adaptive hold enforcement. /// Set via set_isv_signals_ptr(). 0 = NULL (static hold). isv_signals_dev_ptr: u64, @@ -818,6 +839,17 @@ pub struct GpuExperienceCollector { /// argmax-favoring sampling to argmin-favoring. Used briefly to escape Q-uniform /// attractors when the policy is systematically anti-correlated with market. contrarian_active_cache: u8, + + /// B.3 Plan 3 Task 8: cached "in seed phase" flag set per-epoch by the + /// trainer. When `true`, `launch_timestep_loop` dispatches the + /// `scripted_policy_select` kernel BEFORE the network forward pass and + /// SKIPS `experience_action_select` for that step (the scripted action + /// is already in `batch_actions`). When `false`, the network's + /// `experience_action_select` runs normally. Set via + /// `set_seed_phase_active`. Cold-path CPU read of ISV slot at the + /// per-epoch boundary is the trigger; the hot-path action computation + /// itself is 100% GPU. + seed_phase_active_cache: bool, } impl Drop for GpuExperienceCollector { @@ -1355,6 +1387,36 @@ impl GpuExperienceCollector { .map_err(|e| MLError::ModelError(format!("state_kl_moment_match load: {e}")))? }; + // B.3 Plan 3 Task 8: load scripted_policy_select kernel. + let scripted_policy_kernel = { + use super::gpu_dqn_trainer::SCRIPTED_POLICY_CUBIN; + let sp_module = stream.context() + .load_cubin(SCRIPTED_POLICY_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("scripted_policy cubin load: {e}")))?; + sp_module.load_function("scripted_policy_select") + .map_err(|e| MLError::ModelError(format!("scripted_policy_select load: {e}")))? + }; + + // B.3 Plan 3 Task 8: load seed_step_counter_update kernel. + let seed_step_counter_kernel = { + use super::gpu_dqn_trainer::SEED_STEP_COUNTER_CUBIN; + let sc_module = stream.context() + .load_cubin(SEED_STEP_COUNTER_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("seed_step_counter cubin load: {e}")))?; + sc_module.load_function("seed_step_counter_update") + .map_err(|e| MLError::ModelError(format!("seed_step_counter_update load: {e}")))? + }; + + // C.5 Plan 3 Task 9: load cql_alpha_seed_update kernel. + let cql_alpha_seed_update_kernel = { + use super::gpu_dqn_trainer::CQL_ALPHA_SEED_UPDATE_CUBIN; + let cs_module = stream.context() + .load_cubin(CQL_ALPHA_SEED_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("cql_alpha_seed_update cubin load: {e}")))?; + cs_module.load_function("cql_alpha_seed_update") + .map_err(|e| MLError::ModelError(format!("cql_alpha_seed_update load: {e}")))? + }; + // Sort buffers: sized to next power-of-2 of max possible N (episodes × timesteps × cf_mult). // Bitonic sort requires power-of-2 sized arrays. let max_reward_n = alloc_episodes * alloc_timesteps * 2; // cf_mult = 2 @@ -1535,6 +1597,9 @@ impl GpuExperienceCollector { trade_attempt_rate_ema_kernel, plan_threshold_update_kernel, state_kl_kernel, + scripted_policy_kernel, + seed_step_counter_kernel, + cql_alpha_seed_update_kernel, step_counter_gpu, step_counter_kernel, noise_kernel, @@ -1557,6 +1622,7 @@ impl GpuExperienceCollector { learning_health_cache: 1.0, // D4/N4: assume healthy at construction last_cf_ratio_eff: 0.5, // D4/N4: standard cf_ratio at healthy state contrarian_active_cache: 0, // D7/N7: off by default + seed_phase_active_cache: false, // B.3 Plan 3 Task 8: cold-start in network mode; trainer flips on at epoch boundary if SEED_STEPS_DONE < TARGET reward_rank_kernel, reward_compute_abs_sharpe_kernel, bitonic_sort_step_kernel, @@ -1621,6 +1687,23 @@ impl GpuExperienceCollector { self.learning_health_cache = health.clamp(0.0, 1.0); } + /// B.3 Plan 3 Task 8: propagate the seed-phase flag from the trainer + /// (called once per epoch boundary). The trainer reads + /// ISV[SEED_STEPS_DONE_INDEX] vs ISV[SEED_STEPS_TARGET_INDEX] from its + /// pinned ISV mirror; if DONE < TARGET it sets `active=true` and the + /// next `launch_timestep_loop` will dispatch the scripted-policy + /// kernel instead of the network's `experience_action_select`. + /// Cold-path CPU decision; the action computation itself stays + /// 100% GPU. + pub fn set_seed_phase_active(&mut self, active: bool) { + self.seed_phase_active_cache = active; + } + + /// Read-only accessor for diagnostics / smoke tests. + pub fn seed_phase_active(&self) -> bool { + self.seed_phase_active_cache + } + /// D4/N4: return the last effective cf_ratio (set during collect_experiences_gpu). pub fn last_cf_ratio_eff(&self) -> f32 { self.last_cf_ratio_eff @@ -2288,6 +2371,82 @@ impl GpuExperienceCollector { Ok(()) } + /// B.3 Plan 3 Task 8 (spec §4.B.3): GPU-driven seed-phase progress + /// tracker. Increments ISV[SEED_STEPS_DONE_INDEX] by `n_samples` and EMAs + /// the derived `max(0, 1 - DONE/TARGET)` into ISV[SEED_FRAC_EMA_INDEX]. + /// + /// Single-block single-thread cold-path. Launched alongside the other + /// per-epoch ISV producers (`launch_reward_component_ema_inplace`, + /// `launch_trade_attempt_rate_ema_inplace`, …). + pub fn launch_seed_step_counter_update_inplace( + &mut self, + n_samples: usize, + alpha_base: f32, + ) -> Result<(), crate::MLError> { + use crate::cuda_pipeline::gpu_dqn_trainer::{ + SEED_STEPS_DONE_INDEX, SEED_STEPS_TARGET_INDEX, + SEED_FRAC_EMA_INDEX, SHARPE_EMA_INDEX, + }; + if self.isv_signals_dev_ptr == 0 { + return Ok(()); + } + let done_idx = SEED_STEPS_DONE_INDEX as i32; + let target_idx = SEED_STEPS_TARGET_INDEX as i32; + let frac_idx = SEED_FRAC_EMA_INDEX as i32; + let sharpe_idx = SHARPE_EMA_INDEX as i32; + let n_step_i32 = n_samples as i32; + let isv_ptr = self.isv_signals_dev_ptr; + unsafe { + self.stream.launch_builder(&self.seed_step_counter_kernel) + .arg(&isv_ptr) + .arg(&done_idx) + .arg(&target_idx) + .arg(&frac_idx) + .arg(&sharpe_idx) + .arg(&alpha_base) + .arg(&n_step_i32) + .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), shared_mem_bytes: 0 }) + .map_err(|e| crate::MLError::ModelError(format!("seed_step_counter_update: {e}")))?; + } + Ok(()) + } + + /// C.5 Plan 3 Task 9 (spec §4.C.5): GPU-driven CQL α ramp coupled to + /// ISV[SEED_FRAC_EMA]. EMAs ISV[CQL_ALPHA_INDEX=48] toward + /// `cql_alpha_final × max(0, 1 - seed_frac)`. + /// + /// Single-block single-thread cold-path. Launched alongside + /// `launch_seed_step_counter_update_inplace` after every + /// `collect_experiences_gpu`. + pub fn launch_cql_alpha_seed_update_inplace( + &mut self, + cql_alpha_final: f32, + alpha_base: f32, + ) -> Result<(), crate::MLError> { + use crate::cuda_pipeline::gpu_dqn_trainer::{ + CQL_ALPHA_INDEX, SEED_FRAC_EMA_INDEX, SHARPE_EMA_INDEX, + }; + if self.isv_signals_dev_ptr == 0 { + return Ok(()); + } + let cql_idx = CQL_ALPHA_INDEX as i32; + let frac_idx = SEED_FRAC_EMA_INDEX as i32; + let sharpe_idx = SHARPE_EMA_INDEX as i32; + let isv_ptr = self.isv_signals_dev_ptr; + unsafe { + self.stream.launch_builder(&self.cql_alpha_seed_update_kernel) + .arg(&isv_ptr) + .arg(&cql_idx) + .arg(&frac_idx) + .arg(&sharpe_idx) + .arg(&alpha_base) + .arg(&cql_alpha_final) + .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), shared_mem_bytes: 0 }) + .map_err(|e| crate::MLError::ModelError(format!("cql_alpha_seed_update: {e}")))?; + } + Ok(()) + } + /// Task 2.X prerequisite (policy-quality scoping): per-magnitude win rate /// + realized step-return variance, computed host-side from the per-sample /// buffers written by `experience_env_step`. @@ -3175,42 +3334,80 @@ impl GpuExperienceCollector { // ── 4. Action selection: epsilon-greedy + Q-gap conviction filter // Uses q_select_buf (quantile-blended Q-values) instead of q_values (E[Q]). // E[Q] in q_values is preserved for monitoring, Q-gap computation, and loss. - let _epsilon = config.epsilon; // v8: epsilon now computed GPU-side via cosine schedule - let q_gap = config.q_gap_threshold; + // + // B.3 Plan 3 Task 8: when `seed_phase_active_cache` is true, dispatch + // the scripted-policy kernel INSTEAD of the network's action_select. + // The scripted kernel writes the same `batch_actions` + `conviction_buf` + // outputs the network kernel would have written, so all downstream + // env_step/cf paths behave identically. Per-epoch cold-path CPU + // decision; the action computation itself stays 100% GPU. + // + // `max_pos` is declared here (outer scope) because the env_step + // launcher below also reads it; previously it was nested inside + // the action_select block and the env_step site relied on + // shadowing through the if-else. let max_pos = config.max_position; - let min_hold_bars_i32: i32 = 0; // hold enforcement removed — cost-driven - unsafe { - self.stream - .launch_builder(&self.action_select_kernel) - .arg(&self.q_select_buf) - .arg(&mut self.batch_actions) - .arg(&mut self.q_gaps_buf) - .arg(&config.eps_start) - .arg(&config.eps_end) - .arg(&config.current_epoch) - .arg(&config.total_epochs) - .arg(&n_i32) - .arg(&b0) - .arg(&b1) - .arg(&b2) - .arg(&b3) - .arg(&q_gap) - .arg(&self.portfolio_states) // ABI compat — unused by action_select - .arg(&min_hold_bars_i32) // ABI compat — unused (pass 0) - .arg(&max_pos) // max_position for direction index - .arg(&config.eps_exp_mult) // per-branch epsilon multiplier: exposure - .arg(&config.eps_ord_mult) // per-branch epsilon multiplier: order - .arg(&config.eps_urg_mult) // per-branch epsilon multiplier: urgency - .arg(&(t as i32)) // timestep for stateless Philox RNG - .arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule) - .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static) - .arg(&(self.contrarian_active_cache as i32)) // D7/N7: Q-negation flag - .arg(&0u64) // out_intent_mag: NULL — training path does not collect - .arg(&mut self.conviction_buf) // out_conviction: per-sample Kelly warmup signal - .launch(launch_cfg) - .map_err(|e| MLError::ModelError(format!( - "experience_action_select t={t}: {e}" - )))?; + if self.seed_phase_active_cache { + let sd_i32 = ml_core::state_layout::STATE_DIM_PADDED as i32; + let bar_idx_i32 = t as i32; + unsafe { + self.stream + .launch_builder(&self.scripted_policy_kernel) + .arg(&self.batch_states) // assembled state [N, state_dim] + .arg(&self.portfolio_states) // [N, PORTFOLIO_STRIDE] + .arg(&n_i32) + .arg(&sd_i32) + .arg(&bar_idx_i32) + .arg(&mut self.batch_actions) // [N] factored action ∈ [0, 108) + .arg(&mut self.conviction_buf) // [N] conviction ∈ [0, 1] + .launch(launch_cfg) + .map_err(|e| MLError::ModelError(format!( + "scripted_policy_select t={t}: {e}" + )))?; + } + // q_gaps_buf is normally written by experience_action_select. + // The scripted policy doesn't compute Q-gaps, so leave the buffer + // at whatever the previous step wrote (or zero on the first step + // post-reset). Downstream env_step does NOT read q_gaps for the + // physics path — it's a diagnostic-only signal — so the seed + // phase produces no q_gap signal until the network takes over. + } else { + let _epsilon = config.epsilon; // v8: epsilon now computed GPU-side via cosine schedule + let q_gap = config.q_gap_threshold; + let min_hold_bars_i32: i32 = 0; // hold enforcement removed — cost-driven + unsafe { + self.stream + .launch_builder(&self.action_select_kernel) + .arg(&self.q_select_buf) + .arg(&mut self.batch_actions) + .arg(&mut self.q_gaps_buf) + .arg(&config.eps_start) + .arg(&config.eps_end) + .arg(&config.current_epoch) + .arg(&config.total_epochs) + .arg(&n_i32) + .arg(&b0) + .arg(&b1) + .arg(&b2) + .arg(&b3) + .arg(&q_gap) + .arg(&self.portfolio_states) // ABI compat — unused by action_select + .arg(&min_hold_bars_i32) // ABI compat — unused (pass 0) + .arg(&max_pos) // max_position for direction index + .arg(&config.eps_exp_mult) // per-branch epsilon multiplier: exposure + .arg(&config.eps_ord_mult) // per-branch epsilon multiplier: order + .arg(&config.eps_urg_mult) // per-branch epsilon multiplier: urgency + .arg(&(t as i32)) // timestep for stateless Philox RNG + .arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule) + .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static) + .arg(&(self.contrarian_active_cache as i32)) // D7/N7: Q-negation flag + .arg(&0u64) // out_intent_mag: NULL — training path does not collect + .arg(&mut self.conviction_buf) // out_conviction: per-sample Kelly warmup signal + .launch(launch_cfg) + .map_err(|e| MLError::ModelError(format!( + "experience_action_select t={t}: {e}" + )))?; + } } // ── 4b. Expert action override (GPU-native MA crossover + ADX) diff --git a/crates/ml/src/cuda_pipeline/scripted_policy_kernel.cu b/crates/ml/src/cuda_pipeline/scripted_policy_kernel.cu new file mode 100644 index 000000000..d1eeb326d --- /dev/null +++ b/crates/ml/src/cuda_pipeline/scripted_policy_kernel.cu @@ -0,0 +1,111 @@ +/* scripted_policy_select — GPU-driven seed-phase action selection. + * + * Plan 3 Task 8 (B.3 GPU-only seeded warm-start). + * + * Replaces `experience_action_select`'s network-Q argmax during the seed + * phase. Per-sample: deterministically assign one of 4 policies (40/20/20/20 + * mix) based on `i % 5`, compute the scripted action, write to actions_out + * and conviction_out. Driven from the launcher in `gpu_experience_collector.rs` + * which reads ISV[SEED_STEPS_DONE_INDEX] vs ISV[SEED_STEPS_TARGET_INDEX] at + * the per-epoch boundary (cold path) to decide whether to dispatch this + * kernel or `experience_action_select`. The action computation itself is + * 100% GPU; CPU only orchestrates the dispatch. + * + * Policy types (factored action layout: dir(4) × mag(3) × ord(3) × urg(3)): + * 0 (UNIFORM, 40%): LCG random direction + half magnitude + market + + * normal urgency. Conviction floored at 0.5 + * (uniform doesn't express confidence). + * 1 (MOMENTUM, 20%): sign(recent_return) → direction; conviction = + * |return| × 10 capped at 1. + * 2 (MEAN_REV, 20%): −sign(recent_return) → direction; same conviction. + * 3 (VWAP_DEV, 20%): Approximate VWAP as one-bar lag (PS_PREV_CLOSE); + * deviation outside ±5bp band → mean-revert direction; + * conviction by relative band breach magnitude. + * + * "Recent return" comes from `state[MARKET_START + 0] = close_now` and + * `portfolio[PS_PREV_CLOSE] = prev_close`. We do NOT pass historical bars — + * one-bar return is sufficient for a seed-phase exploration policy. + * + * Seed-policy direction-flip thresholds (`±0.0001f`) and VWAP band width + * (`±0.0005f × prev_close`) are noise-floor cutoffs below which we hold — + * NOT scaling coefficients (per feedback_adaptive_not_tuned.md carve-out). + * + * GPU-only: no CPU physics mirror. Existing `experience_env_step` runs + * unchanged; we only swap the action source kernel. + */ + +#include +#include "state_layout.cuh" + +extern "C" __global__ void scripted_policy_select( + const float* __restrict__ states, /* [N, state_dim] assembled state vector */ + const float* __restrict__ portfolio_states, /* [N, PS_STRIDE] for prev_close access */ + int n_samples, + int state_dim, + int bar_idx, /* current timestep, mixes into LCG seed */ + int* __restrict__ actions_out, /* [N] factored action index ∈ [0, 108) */ + float* __restrict__ conviction_out /* [N] conviction ∈ [0, 1] */ +) { + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n_samples) return; + + /* Policy assignment: deterministic mix. + * i % 5 ∈ {0,1,2,3,4}; map {0,1} → UNIFORM (2/5 = 40%), 2 → MOMENTUM, + * 3 → MEAN_REV, 4 → VWAP_DEV. Each of the latter three gets 20%. */ + const int mod5 = i % 5; + const int policy_id = (mod5 < 2) ? 0 : (mod5 - 1); /* 0=UNIFORM, 1=MOM, 2=REV, 3=VWAP */ + + /* Read raw_close (state[MARKET_START]) and prev_close (PS slot) for + * return calc. SL_MARKET_START = 0; PS_PREV_CLOSE = 6. */ + const float close_now = states[(long long)i * state_dim + SL_MARKET_START]; + const float prev_close = portfolio_states[(long long)i * PS_STRIDE + PS_PREV_CLOSE]; + const float recent_ret = (prev_close > 1e-3f) + ? (close_now - prev_close) / prev_close + : 0.0f; + + int dir = DIR_HOLD; /* default Hold */ + float conviction = 0.5f; + + if (policy_id == 0) { + /* UNIFORM: LCG random direction across all 4 values */ + unsigned int seed = ((unsigned int)i * 2654435761u) + ^ ((unsigned int)bar_idx * 16777619u); + seed = seed * 1103515245u + 12345u; + dir = (int)((seed >> 16) & 0x3); /* 0..3 */ + conviction = 0.5f; + } else if (policy_id == 1) { + /* MOMENTUM: trade with the move */ + if (recent_ret > 0.0001f) { dir = DIR_LONG; } + else if (recent_ret < -0.0001f) { dir = DIR_SHORT; } + else { dir = DIR_HOLD; } + conviction = fminf(1.0f, fabsf(recent_ret) * 10.0f); + } else if (policy_id == 2) { + /* MEAN_REV: trade against the move */ + if (recent_ret > 0.0001f) { dir = DIR_SHORT; } + else if (recent_ret < -0.0001f) { dir = DIR_LONG; } + else { dir = DIR_HOLD; } + conviction = fminf(1.0f, fabsf(recent_ret) * 10.0f); + } else { + /* VWAP_DEV: approximate VWAP as one-bar lag (PS_PREV_CLOSE). + * The actual VWAP isn't stored in PS; the one-bar-lag proxy is + * acceptable for a seed-phase exploration policy. */ + const float dev = close_now - prev_close; + const float band_width = fmaxf(prev_close * 0.0005f, 1e-3f); + if (dev > prev_close * 0.0005f) { dir = DIR_SHORT; } + else if (dev < -prev_close * 0.0005f) { dir = DIR_LONG; } + else { dir = DIR_HOLD; } + conviction = fminf(1.0f, fabsf(dev) / band_width); + } + + /* Magnitude = Half (idx 1), order = Market (idx 0), urgency = Normal (idx 1). + * Factored layout: action = dir*27 + mag*9 + ord*3 + urg. + * For DIR_HOLD/DIR_FLAT, mag is forced to MAG_QUARTER (0) elsewhere by + * the env-step kernel's deterministic encoding (see CLAUDE.md project + * memory note on direction encoding); we still emit MAG_HALF here + * because the env-step kernel ignores magnitude for Hold/Flat. */ + const int mag = MAG_HALF; /* 1 */ + const int ord = 0; /* Market */ + const int urg = 1; /* Normal */ + actions_out[i] = dir * 27 + mag * 9 + ord * 3 + urg; + conviction_out[i] = conviction; +} diff --git a/crates/ml/src/cuda_pipeline/seed_step_counter_update_kernel.cu b/crates/ml/src/cuda_pipeline/seed_step_counter_update_kernel.cu new file mode 100644 index 000000000..78cbf6868 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/seed_step_counter_update_kernel.cu @@ -0,0 +1,49 @@ +/* seed_step_counter_update — GPU-driven seed-phase progress tracker. + * + * Plan 3 Task 8 (B.3 GPU-only seeded warm-start). + * + * Single-thread cold-path kernel launched after every + * `collect_experiences_gpu` call, alongside the other Plan 3 EMA producers. + * Increments ISV[SEED_STEPS_DONE_INDEX] by `n_samples_this_step`, capped at + * ISV[SEED_STEPS_TARGET_INDEX]. Computes the seed fraction + * `frac = max(0, 1 - done/target)` (smoothly decays 1→0 as the seed phase + * completes) and EMA-tracks it into ISV[SEED_FRAC_EMA_INDEX] with the + * adaptive-rate convention shared with the other Plan 3 producers + * (α = α_base × (1 + 0.5 × |clamp(sharpe, ±2)|)). + * + * SEED_FRAC_EMA is the consumer signal for Task 9 (CQL α ramp). When + * frac ≈ 1 (fully in seed phase), CQL α decays toward 0 (no pessimism on + * exploration data); as frac → 0, CQL α ramps to its config-final value. + */ + +#include + +extern "C" __global__ void seed_step_counter_update( + float* __restrict__ isv_out, /* ISV bus [ISV_TOTAL_DIM] */ + int seed_done_idx, /* = SEED_STEPS_DONE_INDEX = 83 */ + int seed_target_idx, /* = SEED_STEPS_TARGET_INDEX = 82 */ + int seed_frac_idx, /* = SEED_FRAC_EMA_INDEX = 84 */ + int sharpe_ema_idx, /* = SHARPE_EMA_INDEX = 22 */ + float alpha_base, /* base EMA rate (e.g. 0.05) */ + int n_samples_this_step /* number of (i,t) experience samples this collect */ +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float target = isv_out[seed_target_idx]; + const float prev_done = isv_out[seed_done_idx]; + const float new_done = fminf(target, prev_done + (float)n_samples_this_step); + isv_out[seed_done_idx] = new_done; + + /* Smoothly decay seed fraction from 1.0 (cold-start) toward 0 as DONE + * approaches TARGET. Floor 1.0 in denominator avoids div-by-zero when + * target=0 (smoke configs may set this — whole seed phase becomes + * inert, frac→0 immediately). */ + const float new_frac = fmaxf(0.0f, 1.0f - new_done / fmaxf(1.0f, target)); + + /* Adaptive α matches Task 3 (trade_rate_ema) / Task 4 (plan_threshold) + * convention. Higher-Sharpe epochs update the estimate faster. */ + const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx])); + const float alpha = alpha_base * (1.0f + 0.5f * fabsf(sharpe)); + const float prev_frac_ema = isv_out[seed_frac_idx]; + isv_out[seed_frac_idx] = (1.0f - alpha) * prev_frac_ema + alpha * new_frac; +} diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 421d3eea3..8d27cf567 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -146,6 +146,10 @@ #define ISV_PLAN_THRESHOLD_IDX 49 // == PLAN_THRESHOLD_INDEX — plan activation threshold (shifted +3 by D.2 per-branch gamma, Plan 2 Task 3) #define ISV_SHARPE_EMA_IDX 22 // == SHARPE_EMA_INDEX — EMA of training Sharpe ratio (Plan 3 D.4b adaptive threshold) #define ISV_STATE_KL_AMP_IDX 79 // == STATE_KL_AMPLIFICATION_INDEX — Flat-trap escape amplifier ∈ [1, 2] (Plan 3 C.3) +#define ISV_CQL_ALPHA_IDX 48 // == CQL_ALPHA_INDEX — CQL pessimism base coefficient (Plan 1 Task 12; producer upgraded to GPU `cql_alpha_seed_update` by Plan 3 Task 9 C.5) +#define ISV_SEED_STEPS_TARGET_IDX 82 // == SEED_STEPS_TARGET_INDEX — config replay_seed_steps target (Plan 3 Task 8 B.3) +#define ISV_SEED_STEPS_DONE_IDX 83 // == SEED_STEPS_DONE_INDEX — cumulative seed-phase steps completed (Plan 3 Task 8 B.3) +#define ISV_SEED_FRAC_EMA_IDX 84 // == SEED_FRAC_EMA_INDEX — adaptive EMA of (1 - done/target) ∈ [0, 1] (Plan 3 Task 8 B.3; consumed by Task 9 CQL ramp) // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs index 560266f4e..539a33b71 100644 --- a/crates/ml/src/trainers/dqn/config.rs +++ b/crates/ml/src/trainers/dqn/config.rs @@ -1147,6 +1147,15 @@ pub struct DQNHyperparameters { /// Breaks the dueling symmetry trap where advantage heads produce zero Q-gap. /// 0.0 = disabled. Default: 0.1. pub noise_sigma: f32, + + /// Plan 3 Task 8 B.3: number of experience-collection samples drawn from + /// the GPU scripted-policy mixture (40% uniform / 20% momentum / 20% + /// mean-reversion / 20% vwap-deviation) before action selection switches + /// back to the network. Default 100_000. Set to 0 to disable. Smoke configs + /// override to a smaller value (e.g. 10_000) so the transition is + /// observable inside the smoke run. Cast to `u32` at the + /// `GpuDqnTrainConfig` boundary. + pub replay_seed_steps: u32, } impl Default for DQNHyperparameters { @@ -1611,6 +1620,10 @@ impl DQNHyperparameters { // Advantage noise for dueling symmetry breaking (replaces dead NoisyLinear) noise_sigma: 0.1, + + // Plan 3 Task 8 B.3: GPU-only seeded warm-start. + // 100k samples ≈ first 5 production-scale collection epochs. + replay_seed_steps: 100_000, } } } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index a8f17432f..fbba8253c 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -409,6 +409,7 @@ impl FusedTrainingCtx { bottleneck_dim: hyperparams.bottleneck_dim, market_dim: 42, // Always 42 base market features — OFI features bypass bottleneck via portfolio_dim total_epochs: hyperparams.epochs, + replay_seed_steps: hyperparams.replay_seed_steps, }; // Create weight set pointer views AFTER GpuDqnTrainer is constructed below. diff --git a/crates/ml/src/trainers/dqn/monitors/cql_alpha_monitor.rs b/crates/ml/src/trainers/dqn/monitors/cql_alpha_monitor.rs new file mode 100644 index 000000000..632ecdee0 --- /dev/null +++ b/crates/ml/src/trainers/dqn/monitors/cql_alpha_monitor.rs @@ -0,0 +1,114 @@ +//! CPU-side observer for the GPU-driven CQL α ramp. +//! +//! Plan 3 Task 9 C.5. The `cql_alpha_seed_update` GPU kernel reads +//! ISV[SEED_FRAC_EMA_INDEX=84] (Task 8) and EMAs ISV[CQL_ALPHA_INDEX=48] +//! toward `config.cql_alpha × max(0, 1 - seed_frac)`. During the seed +//! phase (frac ≈ 1) the target is 0 so CQL α decays toward 0 — no +//! pessimism on exploration data populated by the scripted-policy +//! mixture. As frac → 0 the target ramps to `config.cql_alpha` and the +//! CQL gradient kernel sees full pessimism on network-driven +//! trajectories. +//! +//! This monitor is a pure read-only observer: it surfaces both the +//! current CQL α effective value and the underlying seed-fraction +//! signal in HEALTH_DIAG and tracks per-epoch fire-rate so the +//! `controller_activity` smoke test can flag thrash. + +use crate::cuda_pipeline::gpu_dqn_trainer::{ + CQL_ALPHA_INDEX, SEED_FRAC_EMA_INDEX, +}; +use crate::trainers::dqn::adaptive_monitor::{ + AdaptiveMonitor, DiagSnapshot, FireRateStats, IsvBus, +}; + +pub(crate) struct CqlAlphaMonitor { + last_alpha: f32, + fire_rate: FireRateStats, +} + +impl CqlAlphaMonitor { + pub(crate) fn new() -> Self { + Self { last_alpha: 0.0, fire_rate: FireRateStats::default() } + } +} + +impl AdaptiveMonitor for CqlAlphaMonitor { + fn read(&self, isv: &IsvBus<'_>) -> f32 { + isv.read(CQL_ALPHA_INDEX) + } + + fn diagnose(&self, isv: &IsvBus<'_>) -> DiagSnapshot { + let alpha = isv.read(CQL_ALPHA_INDEX); + let frac = isv.read(SEED_FRAC_EMA_INDEX); + DiagSnapshot::new() + .with("cql_alpha.eff", alpha) + .with("cql_alpha.seed_frac", frac) + .with("fire_rate", self.fire_rate.fire_rate()) + } + + fn observe(&mut self, current: f32) { + let fired = (current - self.last_alpha).abs() > 1e-7; + self.fire_rate.record_fire(fired); + self.last_alpha = current; + } + + fn fire_rate(&self) -> &FireRateStats { + &self.fire_rate + } + + fn name(&self) -> &'static str { + "cql_alpha" + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn make_isv(alpha: f32, frac: f32) -> Vec { + let n = SEED_FRAC_EMA_INDEX.max(CQL_ALPHA_INDEX) + 1; + let mut v = vec![0.0f32; n]; + v[CQL_ALPHA_INDEX] = alpha; + v[SEED_FRAC_EMA_INDEX] = frac; + v + } + + #[test] + fn read_returns_alpha_slot() { + let mut buf = make_isv(0.014, 0.50); + let bus = IsvBus::new(&mut buf); + let mon = CqlAlphaMonitor::new(); + let v = mon.read(&bus); + assert!((v - 0.014).abs() < 1e-7, "expected alpha=0.014, got {v}"); + } + + #[test] + fn diagnose_exposes_alpha_and_seed_frac() { + let mut buf = make_isv(0.007, 0.50); + let bus = IsvBus::new(&mut buf); + let mon = CqlAlphaMonitor::new(); + let snap = mon.diagnose(&bus); + assert!(snap.fields.contains_key("cql_alpha.eff")); + assert!(snap.fields.contains_key("cql_alpha.seed_frac")); + assert!(snap.fields.contains_key("fire_rate")); + let got_alpha = snap.fields["cql_alpha.eff"]; + let got_frac = snap.fields["cql_alpha.seed_frac"]; + assert!((got_alpha - 0.007).abs() < 1e-6); + assert!((got_frac - 0.50).abs() < 1e-6); + } + + #[test] + fn observe_tracks_fire_rate_on_change() { + // Mirrors the PlanThresholdMonitor convention: `last_alpha` starts + // at 0.0 (sentinel), so the first observe(0.005) is a fire (0→0.005), + // the second observe(0.005) is not (no change), and observe(0.010) + // fires again. Expected 2/3. + let mut mon = CqlAlphaMonitor::new(); + mon.observe(0.005); + mon.observe(0.005); + mon.observe(0.010); + let rate = mon.fire_rate().fire_rate(); + assert!((rate - 2.0 / 3.0).abs() < 1e-6, + "expected 2/3 fire rate, got {rate}"); + } +} diff --git a/crates/ml/src/trainers/dqn/monitors/mod.rs b/crates/ml/src/trainers/dqn/monitors/mod.rs index 3f780461f..20eb7eac3 100644 --- a/crates/ml/src/trainers/dqn/monitors/mod.rs +++ b/crates/ml/src/trainers/dqn/monitors/mod.rs @@ -6,11 +6,13 @@ //! smoke-test fire-rate tracking. No CPU-side adaptive computation. pub(crate) mod atoms_monitor; +pub(crate) mod cql_alpha_monitor; pub(crate) mod epsilon_monitor; pub(crate) mod gamma_monitor; pub(crate) mod grad_balancer_monitor; pub(crate) mod kelly_cap_monitor; pub(crate) mod plan_threshold_monitor; pub(crate) mod reward_component_monitor; +pub(crate) mod seed_monitor; pub(crate) mod state_kl_monitor; pub(crate) mod tau_monitor; diff --git a/crates/ml/src/trainers/dqn/monitors/seed_monitor.rs b/crates/ml/src/trainers/dqn/monitors/seed_monitor.rs new file mode 100644 index 000000000..2f5cfc2f2 --- /dev/null +++ b/crates/ml/src/trainers/dqn/monitors/seed_monitor.rs @@ -0,0 +1,119 @@ +//! CPU-side observer for the GPU-driven seed-phase progress. +//! +//! Plan 3 Task 8 B.3. The `seed_step_counter_update` GPU kernel reads +//! ISV[SEED_STEPS_TARGET_INDEX=82] (constructor-written from +//! `config.replay_seed_steps`) and increments ISV[SEED_STEPS_DONE_INDEX=83] +//! by the number of (i,t) samples in each `collect_experiences_gpu` call. +//! It then EMAs the derived `max(0, 1 - DONE/TARGET)` into +//! ISV[SEED_FRAC_EMA_INDEX=84]. The seed-fraction signal is consumed by +//! the Task 9 `cql_alpha_seed_update` kernel to ramp CQL α from 0 +//! (during the seed phase) toward `config.cql_alpha` (post-seed). +//! +//! This monitor is a pure read-only observer: it surfaces both the +//! cumulative steps and the seed-fraction EMA in HEALTH_DIAG and tracks +//! per-epoch fire-rate against the seed-fraction so the +//! `controller_activity` smoke test can flag thrash. + +use crate::cuda_pipeline::gpu_dqn_trainer::{ + SEED_FRAC_EMA_INDEX, SEED_STEPS_DONE_INDEX, SEED_STEPS_TARGET_INDEX, +}; +use crate::trainers::dqn::adaptive_monitor::{ + AdaptiveMonitor, DiagSnapshot, FireRateStats, IsvBus, +}; + +pub(crate) struct SeedMonitor { + last_frac: f32, + fire_rate: FireRateStats, +} + +impl SeedMonitor { + pub(crate) fn new() -> Self { + // Sentinel `0.0` matches the convention used by `PlanThresholdMonitor` + // and `StateKlMonitor`: the first `observe` call always fires (the + // GPU-driven seed-fraction starts at cold-start 1.0, so the + // 0.0 → 1.0 first transition is a real "first observation" event). + Self { last_frac: 0.0, fire_rate: FireRateStats::default() } + } +} + +impl AdaptiveMonitor for SeedMonitor { + fn read(&self, isv: &IsvBus<'_>) -> f32 { + isv.read(SEED_FRAC_EMA_INDEX) + } + + fn diagnose(&self, isv: &IsvBus<'_>) -> DiagSnapshot { + let target = isv.read(SEED_STEPS_TARGET_INDEX); + let done = isv.read(SEED_STEPS_DONE_INDEX); + let frac = isv.read(SEED_FRAC_EMA_INDEX); + DiagSnapshot::new() + .with("seed.steps_target", target) + .with("seed.steps_done", done) + .with("seed.frac_ema", frac) + .with("fire_rate", self.fire_rate.fire_rate()) + } + + fn observe(&mut self, current: f32) { + let fired = (current - self.last_frac).abs() > 1e-7; + self.fire_rate.record_fire(fired); + self.last_frac = current; + } + + fn fire_rate(&self) -> &FireRateStats { + &self.fire_rate + } + + fn name(&self) -> &'static str { + "seed" + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn make_isv(target: f32, done: f32, frac: f32) -> Vec { + let mut v = vec![0.0f32; SEED_FRAC_EMA_INDEX + 1]; + v[SEED_STEPS_TARGET_INDEX] = target; + v[SEED_STEPS_DONE_INDEX] = done; + v[SEED_FRAC_EMA_INDEX] = frac; + v + } + + #[test] + fn read_returns_frac_slot() { + let mut buf = make_isv(100_000.0, 5_000.0, 0.95); + let bus = IsvBus::new(&mut buf); + let mon = SeedMonitor::new(); + let v = mon.read(&bus); + assert!((v - 0.95).abs() < 1e-7, "expected frac=0.95, got {v}"); + } + + #[test] + fn diagnose_exposes_target_done_frac() { + let mut buf = make_isv(100_000.0, 25_000.0, 0.75); + let bus = IsvBus::new(&mut buf); + let mon = SeedMonitor::new(); + let snap = mon.diagnose(&bus); + assert!(snap.fields.contains_key("seed.steps_target")); + assert!(snap.fields.contains_key("seed.steps_done")); + assert!(snap.fields.contains_key("seed.frac_ema")); + assert!(snap.fields.contains_key("fire_rate")); + let got_target = snap.fields["seed.steps_target"]; + let got_done = snap.fields["seed.steps_done"]; + let got_frac = snap.fields["seed.frac_ema"]; + assert!((got_target - 100_000.0).abs() < 1e-3); + assert!((got_done - 25_000.0).abs() < 1e-3); + assert!((got_frac - 0.75).abs() < 1e-6); + } + + #[test] + fn observe_tracks_fire_rate_on_change() { + let mut mon = SeedMonitor::new(); + mon.observe(1.00); + mon.observe(1.00); + mon.observe(0.80); + let rate = mon.fire_rate().fire_rate(); + assert!((rate - 2.0 / 3.0).abs() < 1e-6, + "expected 2/3 fire rate, got {rate}"); + } +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 7f74fda36..b10c9e492 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -153,8 +153,8 @@ impl StateResetRegistry { }, RegistryEntry { name: "isv_cql_alpha", - category: ResetCategory::SchemaContract, - description: "ISV[CQL_ALPHA_INDEX=48] — CQL pessimism base coefficient; Plan 3 B.3 may make reactive", + category: ResetCategory::FoldReset, + description: "ISV[CQL_ALPHA_INDEX=48] — CQL pessimism base coefficient; producer upgraded SchemaContract → FoldReset by Plan 3 Task 9 C.5 (GPU `cql_alpha_seed_update` kernel ramps slot 48 toward `config.cql_alpha × (1 - seed_frac)`); cold-start config.cql_alpha reapplied at fold boundary", }, RegistryEntry { name: "isv_plan_threshold", @@ -241,6 +241,22 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[79] — Flat-trap escape amplifier ∈ [1, 2] driven by KL ratio; consumed by B.1 opp_cost + B.2 bonus (Plan 3 C.3)", }, + // ───── B.3 GPU-only seeded warm-start (Plan 3 Task 8) ─── + RegistryEntry { + name: "isv_seed_steps_target", + category: ResetCategory::SchemaContract, + description: "ISV[SEED_STEPS_TARGET_INDEX=82] — config replay_seed_steps; constant for the run, written once at constructor (Plan 3 B.3)", + }, + RegistryEntry { + name: "isv_seed_steps_done", + category: ResetCategory::FoldReset, + description: "ISV[SEED_STEPS_DONE_INDEX=83] — cumulative seed-phase steps; GPU-incremented each collect_experiences (Plan 3 B.3)", + }, + RegistryEntry { + name: "isv_seed_frac_ema", + category: ResetCategory::FoldReset, + description: "ISV[SEED_FRAC_EMA_INDEX=84] — seed-fraction EMA, Task 9 CQL ramp consumer; cold-start 1.0 (Plan 3 B.3)", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 7e370d9c9..46025fd85 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -648,6 +648,32 @@ impl DQNTrainer { } } + // B.3 Plan 3 Task 8: bootstrap the seed-phase flag from the + // pinned ISV BEFORE the first `collect_experiences_gpu` call. + // The trainer's constructor cold-started ISV[SEED_STEPS_DONE]=0 + // and ISV[SEED_STEPS_TARGET]=config.replay_seed_steps; if the + // target is non-zero, the first epoch starts in seed phase. + // Subsequent epochs re-evaluate this flag at the per-epoch + // boundary (see metric-emit block ~3000 lines below). + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::gpu_dqn_trainer::{ + SEED_STEPS_DONE_INDEX, SEED_STEPS_TARGET_INDEX, + }; + let seed_done = fused.read_isv_signal_at(SEED_STEPS_DONE_INDEX); + let seed_target = fused.read_isv_signal_at(SEED_STEPS_TARGET_INDEX); + let in_seed_phase = seed_target > 0.0 && seed_done < seed_target; + if let Some(ref mut collector) = self.gpu_experience_collector { + collector.set_seed_phase_active(in_seed_phase); + } + if in_seed_phase { + tracing::info!( + target: "dqn::B3", + seed_target = seed_target, + "GPU-only replay seed warm-start: scripted-policy phase begins" + ); + } + } + // Trade plan params + readiness for plan activation/enforcement if let Some(ref fused) = self.fused_ctx { let plan_ptr = fused.plan_params_buf_ptr(); @@ -2714,6 +2740,25 @@ impl DQNTrainer { if let Err(e) = c.launch_plan_threshold_update_inplace(n_main, ema_alpha) { tracing::warn!("B.4 plan_threshold_update launch failed: {e}"); } + // B.3 Plan 3 Task 8 (spec §4.B.3): GPU-driven seed-phase + // progress tracker. Increments + // ISV[SEED_STEPS_DONE_INDEX=83] by `n_main` and EMAs + // `max(0, 1 - DONE/TARGET)` into + // ISV[SEED_FRAC_EMA_INDEX=84]. Cold-path single-thread + // kernel; same α_base as the other Plan 3 producers. + if let Err(e) = c.launch_seed_step_counter_update_inplace(n_main, ema_alpha) { + tracing::warn!("B.3 seed_step_counter_update launch failed: {e}"); + } + // C.5 Plan 3 Task 9 (spec §4.C.5): GPU-driven CQL α ramp + // coupled to ISV[SEED_FRAC_EMA]. EMAs + // ISV[CQL_ALPHA_INDEX=48] toward + // `cql_alpha_final × max(0, 1 - seed_frac)`. Producer- + // only upgrade — CQL gradient kernel consumer (already + // reading slot 48 per Plan 1 Task 12) unchanged. + let cql_alpha_final = self.hyperparams.cql_alpha as f32; + if let Err(e) = c.launch_cql_alpha_seed_update_inplace(cql_alpha_final, ema_alpha) { + tracing::warn!("C.5 cql_alpha_seed_update launch failed: {e}"); + } } } @@ -2995,6 +3040,32 @@ impl DQNTrainer { collector.set_learning_health(health_value); self.last_cf_ratio_eff = Some(collector.last_cf_ratio_eff()); } + + // B.3 Plan 3 Task 8: per-epoch seed-phase orchestration. + // Cold-path CPU read of ISV[SEED_STEPS_DONE/TARGET] to decide + // whether the next `collect_experiences_gpu` should dispatch the + // scripted-policy kernel (DONE < TARGET) or the network's + // `experience_action_select`. The action computation itself stays + // 100% GPU; only the per-epoch dispatch decision is CPU. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::gpu_dqn_trainer::{ + SEED_STEPS_DONE_INDEX, SEED_STEPS_TARGET_INDEX, + }; + let seed_done = fused.read_isv_signal_at(SEED_STEPS_DONE_INDEX); + let seed_target = fused.read_isv_signal_at(SEED_STEPS_TARGET_INDEX); + let in_seed_phase = seed_target > 0.0 && seed_done < seed_target; + if let Some(ref mut collector) = self.gpu_experience_collector { + collector.set_seed_phase_active(in_seed_phase); + } + tracing::debug!( + target: "dqn::B3", + epoch = epoch, + seed_done = seed_done, + seed_target = seed_target, + in_seed_phase = in_seed_phase, + "seed-phase orchestration" + ); + } } // Feed per-step Q-value extremes into monitor (real min/max, not single-sample). @@ -4298,6 +4369,44 @@ impl DQNTrainer { ); } } + // B.3 Plan 3 Task 8: seed-phase progress slots reset at fold + // boundary. SEED_STEPS_DONE → 0 and SEED_FRAC_EMA → 1.0 mirror + // the constructor cold-start so the new fold begins fully in + // seed phase (the GPU `seed_step_counter_update` kernel + // re-decays them as the fold's experience collection proceeds). + // SEED_STEPS_TARGET is SchemaContract — never reset here. + "isv_seed_steps_done" => { + if let Some(ref fused) = self.fused_ctx { + fused.trainer().write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::SEED_STEPS_DONE_INDEX, + 0.0, + ); + } + } + "isv_seed_frac_ema" => { + if let Some(ref fused) = self.fused_ctx { + fused.trainer().write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::SEED_FRAC_EMA_INDEX, + 1.0, + ); + } + } + // C.5 Plan 3 Task 9: CQL α slot promoted to FoldReset. + // Constructor cold-starts to config.cql_alpha; the GPU + // `cql_alpha_seed_update` kernel ramps it toward + // `config.cql_alpha × (1 - seed_frac)` over training. At fold + // boundary we re-apply the cold-start so the new fold's seed + // phase starts from the same configured base. Producer-only + // upgrade — CQL gradient kernel consumer (already reading + // slot 48 per Plan 1 Task 12) unchanged. + "isv_cql_alpha" => { + if let Some(ref fused) = self.fused_ctx { + fused.trainer().write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::CQL_ALPHA_INDEX, + self.hyperparams.cql_alpha as f32, + ); + } + } "isv_gamma_dir_eff" | "isv_gamma_mag_eff" | "isv_gamma_ord_eff" | "isv_gamma_urg_eff" => { // D.2 per-branch gamma slots. Reset to 0.0 at fold boundary; // per_branch_gamma_update GPU kernel re-populates on next epoch-boundary launch. diff --git a/crates/ml/src/training_profile.rs b/crates/ml/src/training_profile.rs index 06a64191b..337030f0d 100644 --- a/crates/ml/src/training_profile.rs +++ b/crates/ml/src/training_profile.rs @@ -124,6 +124,11 @@ pub struct ReplayBufferSection { pub per_alpha: Option, pub per_beta_start: Option, pub regime_replay_decay: Option, + /// Plan 3 Task 8 B.3: number of GPU scripted-policy seed-phase samples + /// before action selection switches to the network. Default 100_000; + /// smoke configs override smaller (e.g. 1_000) so the seed→network + /// transition is observable inside the smoke run. + pub replay_seed_steps: Option, } /// C51 distributional RL parameters. @@ -844,6 +849,9 @@ impl DqnTrainingProfile { if let Some(v) = rb.regime_replay_decay { hp.regime_replay_decay = v as f32; } + if let Some(v) = rb.replay_seed_steps { + hp.replay_seed_steps = v; + } } // [distributional] diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3218d5fe3..f78aa1c7b 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -70,6 +70,11 @@ | `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. | — | +| `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 — Rust Wrappers diff --git a/docs/isv-slots.md b/docs/isv-slots.md index 37e2a8e68..7ddc5ca2c 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -15,7 +15,7 @@ at the head would displace those live signals and require shifting every upstrea 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`:** 82 (Plan 1 + Plan 2 Task 1 C.1 + Plan 2 Task 3 D.2 per-branch gamma + Plan 2 Task 6C D.8 TLOB + Plan 3 Task 1 C.2 reward-component EMAs + Plan 3 Task 3 B.2 trade-attempt novelty + Plan 3 Task 4 B.4 readiness-EMA + Plan 3 Task 7 C.3 state-distribution KL). Post-full DQN v2 rollout: 82+. +**Current `ISV_TOTAL_DIM`:** 87 (Plan 1 + Plan 2 Task 1 C.1 + Plan 2 Task 3 D.2 per-branch gamma + Plan 2 Task 6C D.8 TLOB + Plan 3 Task 1 C.2 reward-component EMAs + Plan 3 Task 3 B.2 trade-attempt novelty + Plan 3 Task 4 B.4 readiness-EMA + Plan 3 Task 7 C.3 state-distribution KL + Plan 3 Task 8 B.3 GPU-only seed warm-start [82..85)). Post-full DQN v2 rollout: 87+. | Index | Name constant | Type | Producer | Consumers | Reset-category | Notes | |---|---|---|---|---|---|---| @@ -75,6 +75,10 @@ new slots are appended to the bus. | [75] | `READINESS_EMA_INDEX` | f32 | GPU `plan_threshold_update` kernel (Plan 3 B.4) | derived → ISV[PLAN_THRESHOLD_INDEX=49]; `PlanThresholdMonitor` | FoldReset | Per-batch mean readiness EMA (α=α_base × (1+0.5×\|clamp(sharpe,−2,2)\|), α_base=0.05). Cold-start 1.0 so derived threshold = 0.5 matches the previous static default. Drives slot 49 producer-only — consumers unchanged. | | [78] | `STATE_KL_TRAIN_VAL_EMA_INDEX` | f32 | GPU `state_kl_moment_match` kernel (Plan 3 C.3) | `StateKLMonitor` (HEALTH_DIAG) | FoldReset | Per-validation-epoch moment-match KL EMA between train-state sample and val-state batch over the OFI block (32 dims). Adaptive α=α_base × (1+0.5×\|clamp(sharpe,−2,2)\|), α_base=0.05. Cold-start 0.0 — first kernel fire EMAs the actual measured KL toward this. | | [79] | `STATE_KL_AMPLIFICATION_INDEX` | f32 | GPU `state_kl_moment_match` kernel (Plan 3 C.3) | `experience_kernels.cu` B.1 opp_cost + B.2 bonus consumers (via `ISV_STATE_KL_AMP_IDX` macro) | FoldReset | Bounded amplifier ∈ [1, 2] tracking trailing-EMA-of-self KL ratio: `1 + clamp(new_ema/prev_ema − 1, 0, 1)`. Cold-start 1.0 — neutral element behind consumers' `fmaxf(1.0, …)` guard. When train/val distributions diverge, both Flat opp-cost and trade-attempt bonus scale up together. Per `pearl_one_unbounded_signal_per_reward.md`: amp is bounded so it stacks safely with B.1's `q_abs_ref` unbounded multiplicand. | -| [80] | `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. Shifted 69→73 by Plan 3 Task 3 B.2, 73→76 by Plan 3 Task 4 B.4, 76→80 by Plan 3 Task 7 C.3. | -| [81] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. Shifted 70→74 by Plan 3 Task 3 B.2, 74→77 by Plan 3 Task 4 B.4, 77→81 by Plan 3 Task 7 C.3. | -| [82) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 3-5 | +| [48] | `CQL_ALPHA_INDEX` (post-Plan-3-T9 producer upgrade) | f32 | GPU `cql_alpha_seed_update` kernel (Plan 3 C.5) | `compute_cql_logit_gradients` (Rust ISV read) | FoldReset (was SchemaContract) | CQL pessimism base coefficient. Producer upgraded SchemaContract → FoldReset by Plan 3 Task 9 C.5: per-epoch GPU kernel EMAs slot 48 toward `config.cql_alpha × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX=84])`. During seed phase (frac=1), target=0 → CQL α decays to 0 (no pessimism on exploration data); as frac → 0, α ramps to `config.cql_alpha` (full pessimism on network-driven trajectories). Constructor cold-starts to `config.cql_alpha`; fold-boundary reset re-applies. | +| [82] | `SEED_STEPS_TARGET_INDEX` | f32 | construct (CPU static, from `config.replay_seed_steps`) | `seed_step_counter_update` GPU kernel (Plan 3 B.3) | SchemaContract | Plan 3 Task 8 B.3 — target number of seed-phase experience-collection samples. Default 100_000; smoke configs may override to 10_000 so the seed→network transition is observable inside the smoke run. | +| [83] | `SEED_STEPS_DONE_INDEX` | f32 | GPU `seed_step_counter_update` kernel (Plan 3 B.3) | training_loop CPU dispatch decision (per-epoch read) + `seed_step_counter_update` self | FoldReset | Plan 3 Task 8 B.3 — cumulative seed-phase steps completed. GPU-incremented by `n_samples_this_step` per `collect_experiences_gpu` call, capped at TARGET. Cold-start 0; fold-boundary reset re-applies. CPU per-epoch read of (DONE, TARGET) decides whether next collect dispatches scripted-policy kernel (DONE < TARGET) or `experience_action_select`. | +| [84] | `SEED_FRAC_EMA_INDEX` | f32 | GPU `seed_step_counter_update` kernel (Plan 3 B.3) | GPU `cql_alpha_seed_update` kernel (Plan 3 C.5) + `SeedMonitor` | FoldReset | Plan 3 Task 8 B.3 — adaptive EMA of `max(0, 1 - DONE/TARGET) ∈ [0, 1]`. Cold-start 1.0 (fully in seed phase) so Task 9's CQL ramp sees `target = final × max(0, 1 - 1.0) = 0` until the seed phase actually decays. Adaptive α = α_base × (1 + 0.5 × \|clamp(sharpe, ±2)\|), α_base = 0.05. | +| [85] | `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. Shifted 69→73 by Plan 3 Task 3 B.2, 73→76 by Plan 3 Task 4 B.4, 76→80 by Plan 3 Task 7 C.3, 80→85 by Plan 3 Task 8 B.3. | +| [86] | `ISV_LAYOUT_FINGERPRINT_HI_INDEX` | u32 bits (in f32) | construct | check_layout_fingerprint | SchemaContract | High 32 bits of u64 FNV-1a structural hash. Shifted 70→74 by Plan 3 Task 3 B.2, 74→77 by Plan 3 Task 4 B.4, 77→81 by Plan 3 Task 7 C.3, 81→86 by Plan 3 Task 8 B.3. | +| [87) | (reserved for DQN v2) | | | | | Allocated incrementally by Plans 3-5 |