From 91d1a52b9cac3e77da30551ce3968833b7397c14 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 15 May 2026 14:30:40 +0200 Subject: [PATCH] =?UTF-8?q?refactor(alpha):=20rename=20phase=5Fe=5F*=20?= =?UTF-8?q?=E2=86=92=20alpha=5F*=20=E2=80=94=20system-scoped=20naming?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The kill-criteria producer, Munchausen target kernel, Rust launchers, fit/baseline binaries, and their output JSON artifacts are *durable infrastructure* of the alpha trading system (live across Phase E/F/G/...), not milestone-scoped to Phase E specifically. Aligns with the earlier `phase_e_isv_slots.rs` → `alpha_isv_slots.rs` rename rationale. What was renamed: Code files: crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu → alpha_kill_criteria.cu crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu → alpha_munchausen_target.cu crates/ml/src/cuda_pipeline/phase_e_kernels.rs → alpha_kernels.rs crates/ml/examples/phase_e_fit_fill_model.rs → alpha_fit_fill_model.rs crates/ml/examples/phase_e_random_baseline.rs → alpha_random_baseline.rs Artifacts: config/ml/phase_e_fill_coeffs.json → alpha_fill_coeffs.json config/ml/phase_e_random_baseline.json → alpha_random_baseline.json Kernel function names: phase_e_kill_criteria_compute_kernel → alpha_kill_criteria_compute_kernel phase_e_munchausen_target_kernel → alpha_munchausen_target_kernel Rust launcher names: launch_phase_e_kill_criteria → launch_alpha_kill_criteria launch_phase_e_munchausen_target → launch_alpha_munchausen_target Static cubin names: PHASE_E_MUNCHAUSEN_TARGET_CUBIN → ALPHA_MUNCHAUSEN_TARGET_CUBIN Historical milestone tags in doc-comments ("Phase E.1 Task N (2026-05-15)") are RETAINED — they record WHEN the work landed and what plan it implemented, which doesn't change with the system-scoped rename. Plus: ADDS the alpha_munchausen_target GPU smoke test in alpha_kernels.rs. End-to-end validates the launcher + kernel against hand-computed expected values: batch=2 with one terminal sample; expected targets [29.8, 1.1]; got match within 0.05 tolerance on RTX 3050 Ti. PROVES the Task 9/10 kernels actually run on GPU. All affected references updated in: - build.rs (kernel compile list) - mod.rs (module registration) - state_reset_registry.rs (4 RegistryEntry descriptions for slots 539-542) - alpha_isv_slots.rs (slot table comment) - docs/isv-slots.md (audit-doc cross-references) Verified: cargo test -p ml --lib alpha_kernels: 2/2 pass (including GPU smoke) cargo test -p ml --lib state_reset_registry: 10/10 pass cargo build -p ml --release --example alpha_fit_fill_model --example alpha_random_baseline: clean --- ...ill_coeffs.json => alpha_fill_coeffs.json} | 0 ...seline.json => alpha_random_baseline.json} | 0 crates/ml/build.rs | 4 +- ..._fill_model.rs => alpha_fit_fill_model.rs} | 8 +- ...m_baseline.rs => alpha_random_baseline.rs} | 14 +- .../ml/src/cuda_pipeline/alpha_isv_slots.rs | 8 +- .../{phase_e_kernels.rs => alpha_kernels.rs} | 143 +++++++++++++++--- ...ill_criteria.cu => alpha_kill_criteria.cu} | 4 +- ...n_target.cu => alpha_munchausen_target.cu} | 4 +- crates/ml/src/cuda_pipeline/mod.rs | 2 +- .../src/trainers/dqn/state_reset_registry.rs | 8 +- docs/isv-slots.md | 28 ++-- 12 files changed, 160 insertions(+), 63 deletions(-) rename config/ml/{phase_e_fill_coeffs.json => alpha_fill_coeffs.json} (100%) rename config/ml/{phase_e_random_baseline.json => alpha_random_baseline.json} (100%) rename crates/ml/examples/{phase_e_fit_fill_model.rs => alpha_fit_fill_model.rs} (98%) rename crates/ml/examples/{phase_e_random_baseline.rs => alpha_random_baseline.rs} (96%) rename crates/ml/src/cuda_pipeline/{phase_e_kernels.rs => alpha_kernels.rs} (52%) rename crates/ml/src/cuda_pipeline/{phase_e_kill_criteria.cu => alpha_kill_criteria.cu} (98%) rename crates/ml/src/cuda_pipeline/{phase_e_munchausen_target.cu => alpha_munchausen_target.cu} (97%) diff --git a/config/ml/phase_e_fill_coeffs.json b/config/ml/alpha_fill_coeffs.json similarity index 100% rename from config/ml/phase_e_fill_coeffs.json rename to config/ml/alpha_fill_coeffs.json diff --git a/config/ml/phase_e_random_baseline.json b/config/ml/alpha_random_baseline.json similarity index 100% rename from config/ml/phase_e_random_baseline.json rename to config/ml/alpha_random_baseline.json diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 54808caa3..aaf22e248 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1594,13 +1594,13 @@ fn main() { // into scratch[N..N+4]. Chained `apply_pearls_ad_kernel` (n_slots=4) // smooths into ISV[539..543) — see alpha_isv_slots.rs. Slot 547/548 // (random baseline mean/std) are read directly from ISV by index. - "phase_e_kill_criteria.cu", + "alpha_kill_criteria.cu", // Phase E.1 Task 10 (2026-05-15): Munchausen DQN target term per // Vieillard et al. 2020. Adds an implicit KL-regularisation bonus // m = α_m · max(τ · log π(a|s), log_clip_min) to the TD target, // plus a soft-V bootstrap V_soft(s') = τ · log Σ exp(Q/τ). // One thread per batch sample; log-sum-exp-stable softmax. - "phase_e_munchausen_target.cu", + "alpha_munchausen_target.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/examples/phase_e_fit_fill_model.rs b/crates/ml/examples/alpha_fit_fill_model.rs similarity index 98% rename from crates/ml/examples/phase_e_fit_fill_model.rs rename to crates/ml/examples/alpha_fit_fill_model.rs index 96f492d2d..0218d0682 100644 --- a/crates/ml/examples/phase_e_fit_fill_model.rs +++ b/crates/ml/examples/alpha_fit_fill_model.rs @@ -24,11 +24,11 @@ //! ## Run //! //! ```bash -//! SQLX_OFFLINE=true cargo run --release --example phase_e_fit_fill_model -- \ +//! SQLX_OFFLINE=true cargo run --release --example alpha_fit_fill_model -- \ //! --mbp10-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT \ //! --trades-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-trades/ES.FUT \ //! --window-seconds 60 \ -//! --out-path phase_e_fill_coeffs.json +//! --out-path alpha_fill_coeffs.json //! ``` use std::fs::File; @@ -60,7 +60,7 @@ use ml::trainers::dqn::collect_dbn_files_recursive; #[derive(Debug, Parser)] #[command( - name = "phase_e_fit_fill_model", + name = "alpha_fit_fill_model", about = "Phase E.0 Task 5 — fit FillModel coefficients from historical MBP-10 + trades" )] struct Cli { @@ -74,7 +74,7 @@ struct Cli { #[arg(long, default_value_t = 60.0)] window_seconds: f32, /// Output JSON path. - #[arg(long, default_value = "config/ml/phase_e_fill_coeffs.json")] + #[arg(long, default_value = "config/ml/alpha_fill_coeffs.json")] out_path: PathBuf, /// MBP-10 events between snapshot emissions (denser = larger sample). #[arg(long, default_value_t = 10)] diff --git a/crates/ml/examples/phase_e_random_baseline.rs b/crates/ml/examples/alpha_random_baseline.rs similarity index 96% rename from crates/ml/examples/phase_e_random_baseline.rs rename to crates/ml/examples/alpha_random_baseline.rs index 9462ccc35..f12364dbb 100644 --- a/crates/ml/examples/phase_e_random_baseline.rs +++ b/crates/ml/examples/alpha_random_baseline.rs @@ -10,12 +10,12 @@ //! ## Run //! //! ```bash -//! SQLX_OFFLINE=true cargo run -p ml --release --example phase_e_random_baseline -- \ +//! SQLX_OFFLINE=true cargo run -p ml --release --example alpha_random_baseline -- \ //! --mbp10-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT \ -//! --fill-coeffs config/ml/phase_e_fill_coeffs.json \ +//! --fill-coeffs config/ml/alpha_fill_coeffs.json \ //! --horizon 600 \ //! --n-episodes 10000 \ -//! --out-path config/ml/phase_e_random_baseline.json +//! --out-path config/ml/alpha_random_baseline.json //! ``` //! //! ## L1-only data limitation (inherited) @@ -47,7 +47,7 @@ use ml::trainers::dqn::collect_dbn_files_recursive; /// Convert raw `Mbp10Msg.price` (i64 fixed-point at 1e9 scale per the /// databento DBN format) to f32 in price units. See -/// `phase_e_fit_fill_model.rs` for the rationale (mismatch with +/// `alpha_fit_fill_model.rs` for the rationale (mismatch with /// `BidAskPair::price_to_f64`'s 1e12 convention). #[inline] fn raw_price_to_f32(fixed: i64) -> f32 { @@ -59,7 +59,7 @@ const TICK: f32 = 0.25; #[derive(Debug, Parser)] #[command( - name = "phase_e_random_baseline", + name = "alpha_random_baseline", about = "Phase E.0 Task 7 — random-uniform policy reward baseline" )] struct Cli { @@ -67,7 +67,7 @@ struct Cli { #[arg(long)] mbp10_dir: PathBuf, /// Path to fitted FillModel JSON (from Task 5b). - #[arg(long, default_value = "config/ml/phase_e_fill_coeffs.json")] + #[arg(long, default_value = "config/ml/alpha_fill_coeffs.json")] fill_coeffs: PathBuf, /// Episode horizon in snapshots. #[arg(long, default_value_t = 600)] @@ -91,7 +91,7 @@ struct Cli { #[arg(long, default_value_t = 500_000)] max_snapshots: usize, /// Output JSON path. - #[arg(long, default_value = "config/ml/phase_e_random_baseline.json")] + #[arg(long, default_value = "config/ml/alpha_random_baseline.json")] out_path: PathBuf, } diff --git a/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs b/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs index 407969e1c..8c2c158d5 100644 --- a/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs @@ -14,10 +14,10 @@ //! //! | Index | Name | Role | Update cadence | Producer kernel | //! |-------|------------------------------|--------------|-----------------|------------------------------------------| -//! | 539 | Q_SPREAD_EMA_INDEX | Diagnostic | per-rollout-step| phase_e_kill_criteria.cu | -//! | 540 | ACTION_ENTROPY_EMA_INDEX | Diagnostic | per-rollout-step| phase_e_kill_criteria.cu | -//! | 541 | RETURN_VS_RANDOM_EMA_INDEX | Diagnostic | per-trade-close | phase_e_kill_criteria.cu | -//! | 542 | EARLY_Q_MOVEMENT_EMA_INDEX | Diagnostic | per-epoch | phase_e_kill_criteria.cu | +//! | 539 | Q_SPREAD_EMA_INDEX | Diagnostic | per-rollout-step| alpha_kill_criteria.cu | +//! | 540 | ACTION_ENTROPY_EMA_INDEX | Diagnostic | per-rollout-step| alpha_kill_criteria.cu | +//! | 541 | RETURN_VS_RANDOM_EMA_INDEX | Diagnostic | per-trade-close | alpha_kill_criteria.cu | +//! | 542 | EARLY_Q_MOVEMENT_EMA_INDEX | Diagnostic | per-epoch | alpha_kill_criteria.cu | //! | 543 | STACKER_THRESHOLD_INDEX | Controller | per-rollout-step| stacker_threshold_controller.cu | //! | 544 | TRADE_RATE_TARGET_INDEX | Anchor | static | host init | //! | 545 | TRADE_RATE_OBSERVED_EMA_INDEX| Diagnostic | per-rollout-step| stacker_threshold_controller.cu | diff --git a/crates/ml/src/cuda_pipeline/phase_e_kernels.rs b/crates/ml/src/cuda_pipeline/alpha_kernels.rs similarity index 52% rename from crates/ml/src/cuda_pipeline/phase_e_kernels.rs rename to crates/ml/src/cuda_pipeline/alpha_kernels.rs index bca1b2f55..b4c99ffd4 100644 --- a/crates/ml/src/cuda_pipeline/phase_e_kernels.rs +++ b/crates/ml/src/cuda_pipeline/alpha_kernels.rs @@ -2,14 +2,14 @@ //! //! Rust-side launchers for the two CUDA kernels compiled by `build.rs`: //! -//! - `phase_e_kill_criteria_compute_kernel` (`phase_e_kill_criteria.cu`) — +//! - `alpha_kill_criteria_compute_kernel` (`alpha_kill_criteria.cu`) — //! writes 4 raw diagnostic observations (Q-spread, action entropy, //! return-vs-random, early-Q-movement) into scratch_out[0..4]. The caller //! is responsible for chaining `apply_pearls_ad_kernel` (n_slots=4) //! afterward to smooth into ISV slots 539..542. See //! [`super::alpha_isv_slots`] for the slot reservation. //! -//! - `phase_e_munchausen_target_kernel` (`phase_e_munchausen_target.cu`) — +//! - `alpha_munchausen_target_kernel` (`alpha_munchausen_target.cu`) — //! Munchausen-augmented DQN target per Vieillard et al. 2020. One thread //! per batch sample; reads Q_online(s, ·) and Q_target(s', ·), writes //! `target_out[b] = r + α_m · max(τ · log π(a|s), log_clip_min) + γ · V_soft(s')`. @@ -28,7 +28,7 @@ use crate::cuda_pipeline::alpha_isv_slots::{ }; use crate::MLError; -/// Launch `phase_e_kill_criteria_compute_kernel` on `stream`. +/// Launch `alpha_kill_criteria_compute_kernel` on `stream`. /// /// The kernel reads `q_values[batch, n_actions]`, `action_counts[n_actions]`, /// and `scalar_inputs[3]` (mapped-pinned: rollout_reward_mean, q_init_norm, @@ -46,7 +46,7 @@ use crate::MLError; /// All buffer pointers MUST be valid device pointers. `scratch_out_dev` /// MUST point at a writable [4]-float region; passing a shorter buffer /// will corrupt adjacent memory. -pub(crate) unsafe fn launch_phase_e_kill_criteria( +pub(crate) unsafe fn launch_alpha_kill_criteria( stream: &cudarc::driver::CudaStream, kernel: &cudarc::driver::CudaFunction, q_values_dev: u64, @@ -59,8 +59,8 @@ pub(crate) unsafe fn launch_phase_e_kill_criteria( ) -> Result<(), MLError> { use cudarc::driver::{LaunchConfig, PushKernelArg}; - debug_assert!(batch >= 0, "launch_phase_e_kill_criteria: batch must be non-negative"); - debug_assert!(n_actions > 0, "launch_phase_e_kill_criteria: n_actions must be positive"); + debug_assert!(batch >= 0, "launch_alpha_kill_criteria: batch must be non-negative"); + debug_assert!(n_actions > 0, "launch_alpha_kill_criteria: n_actions must be positive"); debug_assert!(q_values_dev != 0, "q_values_dev must be a valid device pointer"); debug_assert!(action_counts_dev != 0, "action_counts_dev must be a valid device pointer"); debug_assert!(scalar_inputs_dev != 0, "scalar_inputs_dev must be a valid device pointer"); @@ -87,12 +87,12 @@ pub(crate) unsafe fn launch_phase_e_kill_criteria( .arg(&scratch_out_dev) .launch(cfg) .map_err(|e| { - MLError::ModelError(format!("phase_e_kill_criteria launch: {e}")) + MLError::ModelError(format!("alpha_kill_criteria launch: {e}")) })?; Ok(()) } -/// Launch `phase_e_munchausen_target_kernel` on `stream`. +/// Launch `alpha_munchausen_target_kernel` on `stream`. /// /// Computes the Munchausen-augmented target for each batch sample: /// `target[i] = r[i] + α_m · max(τ · log π(a|s), log_clip_min) + γ · V_soft(s')` @@ -109,7 +109,7 @@ pub(crate) unsafe fn launch_phase_e_kill_criteria( /// /// All buffer pointers MUST be valid device pointers. `target_out_dev` MUST /// point at a writable [batch]-float region. -pub(crate) unsafe fn launch_phase_e_munchausen_target( +pub(crate) unsafe fn launch_alpha_munchausen_target( stream: &cudarc::driver::CudaStream, kernel: &cudarc::driver::CudaFunction, q_next_dev: u64, @@ -127,8 +127,8 @@ pub(crate) unsafe fn launch_phase_e_munchausen_target( ) -> Result<(), MLError> { use cudarc::driver::{LaunchConfig, PushKernelArg}; - debug_assert!(batch > 0, "launch_phase_e_munchausen_target: batch must be positive"); - debug_assert!(n_actions > 0, "launch_phase_e_munchausen_target: n_actions must be positive"); + debug_assert!(batch > 0, "launch_alpha_munchausen_target: batch must be positive"); + debug_assert!(n_actions > 0, "launch_alpha_munchausen_target: n_actions must be positive"); debug_assert!(tau > 0.0, "tau must be positive (Boltzmann temperature)"); debug_assert!(gamma >= 0.0 && gamma <= 1.0, "gamma must be in [0, 1]"); debug_assert!(alpha_m >= 0.0, "alpha_m must be non-negative"); @@ -162,27 +162,124 @@ pub(crate) unsafe fn launch_phase_e_munchausen_target( .arg(&n_actions) .launch(cfg) .map_err(|e| { - MLError::ModelError(format!("phase_e_munchausen_target launch: {e}")) + MLError::ModelError(format!("alpha_munchausen_target launch: {e}")) })?; Ok(()) } +/// Precompiled Munchausen target cubin, embedded at compile time by build.rs. +static ALPHA_MUNCHAUSEN_TARGET_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/alpha_munchausen_target.cubin")); + #[cfg(test)] mod tests { - // No GPU tests here — launchers are exercised end-to-end by the - // Phase E.1 kernel smoke (see crates/ml/examples/phase_e_kernel_smoke.rs - // for the integration test). The unit-test budget for these tiny - // launchers (which are pure parameter-marshalling) is captured by the - // smoke run, which is the only place they get touched in this milestone. - // - // This stub module exists to anchor a future home for any pure-Rust - // helpers (e.g., parameter validation) that don't require a CUDA context. + use super::*; + use cudarc::driver::{CudaContext, DevicePtr, DevicePtrMut}; #[test] fn module_compiles_clean() { - // Tautology — if this file failed to compile, this test wouldn't - // even build. Acts as a witness that the launchers + their - // dependency on alpha_isv_slots are coherent. + // Tautology — witnesses that the launchers + alpha_isv_slots are coherent. assert!(true); } + + /// End-to-end GPU smoke for the Munchausen target launcher. + /// + /// Loads the cubin, allocates synthetic inputs (batch=2, n_actions=3 with + /// known-by-hand Q-spreads and a terminal flag on sample 1), invokes the + /// kernel via `launch_alpha_munchausen_target`, and validates the two + /// per-sample targets against hand-computed expected values. + /// + /// Sample 0 hand-math (Q_current = [1,2,3], a_taken=0, τ=1): + /// max(Q)=3; sum_exp = exp(-2)+exp(-1)+exp(0) = 0.1353+0.3679+1.0000 = 1.5032 + /// τ·log π(0|s) = (1−3) − 1·log(1.5032) = −2 − 0.4076 = −2.4076 + /// m = 0.9 · max(−2.4076, −1.0) = 0.9 · −1.0 = −0.9 + /// V_soft(Q_next=[10,20,30]) = 30 + 1·log(exp(-20)+exp(-10)+1) ≈ 30.0 + /// target = 1.0 + (−0.9) + 0.99 · 30.0 = 29.8 + /// + /// Sample 1 (terminal, Q_current all zero): + /// π uniform = 1/3; τ·log π = log(1/3) = −1.0986 + /// m = 0.9 · max(−1.0986, −1.0) = 0.9 · −1.0 = −0.9 + /// V_soft = 0 (terminal) + /// target = 2.0 + (−0.9) + 0 = 1.1 + #[test] + fn munchausen_target_smoke_matches_hand_computation() { + let ctx = match CudaContext::new(0) { + Ok(c) => c, + Err(e) => { + eprintln!("Skipping GPU smoke — no CUDA device 0: {e}"); + return; + } + }; + let stream = ctx.default_stream(); + + let module = ctx + .load_cubin(ALPHA_MUNCHAUSEN_TARGET_CUBIN.to_vec()) + .expect("munchausen cubin module load"); + let kernel = module + .load_function("alpha_munchausen_target_kernel") + .expect("munchausen kernel load"); + + let batch: i32 = 2; + let n_actions: i32 = 3; + let q_current: Vec = vec![1.0, 2.0, 3.0, 0.0, 0.0, 0.0]; + let q_next: Vec = vec![10.0, 20.0, 30.0, 0.0, 0.0, 0.0]; + let actions: Vec = vec![0, 1]; + let rewards: Vec = vec![1.0, 2.0]; + let dones: Vec = vec![0.0, 1.0]; + + let q_current_dev = stream.clone_htod(&q_current).expect("upload q_current"); + let q_next_dev = stream.clone_htod(&q_next).expect("upload q_next"); + let actions_dev = stream.clone_htod(&actions).expect("upload actions"); + let rewards_dev = stream.clone_htod(&rewards).expect("upload rewards"); + let dones_dev = stream.clone_htod(&dones).expect("upload dones"); + let mut target_out_dev = stream + .alloc_zeros::(batch as usize) + .expect("alloc target_out"); + + // Launch in a scoped block so the device-ptr guards drop before the + // download step (otherwise the mutable borrow on target_out_dev would + // prevent the &target_out_dev read in clone_dtoh). + { + let (q_next_ptr, _g0) = q_next_dev.device_ptr(&stream); + let (q_current_ptr, _g1) = q_current_dev.device_ptr(&stream); + let (actions_ptr, _g2) = actions_dev.device_ptr(&stream); + let (rewards_ptr, _g3) = rewards_dev.device_ptr(&stream); + let (dones_ptr, _g4) = dones_dev.device_ptr(&stream); + let (target_out_ptr, _g5) = target_out_dev.device_ptr_mut(&stream); + + unsafe { + launch_alpha_munchausen_target( + &stream, + &kernel, + q_next_ptr, + q_current_ptr, + actions_ptr, + rewards_ptr, + dones_ptr, + 0.99, // gamma + 0.9, // alpha_m + 1.0, // tau + -1.0, // log_clip_min + target_out_ptr, + batch, + n_actions, + ) + .expect("kernel launch"); + } + } + stream.synchronize().expect("stream sync"); + + let target_out = stream + .clone_dtoh(&target_out_dev) + .expect("download target_out"); + assert_eq!(target_out.len(), batch as usize); + + let expected = [29.8_f32, 1.1_f32]; + for (i, (got, exp)) in target_out.iter().zip(expected.iter()).enumerate() { + assert!( + (got - exp).abs() < 0.05, + "sample {i}: expected {exp:.3}, got {got:.3} (full output: {target_out:?})" + ); + } + } } diff --git a/crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu b/crates/ml/src/cuda_pipeline/alpha_kill_criteria.cu similarity index 98% rename from crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu rename to crates/ml/src/cuda_pipeline/alpha_kill_criteria.cu index fc2ded2f4..ac451e8cf 100644 --- a/crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu +++ b/crates/ml/src/cuda_pipeline/alpha_kill_criteria.cu @@ -1,4 +1,4 @@ -// crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu +// crates/ml/src/cuda_pipeline/alpha_kill_criteria.cu // // Phase E.0 kill-criteria diagnostic producer kernel (Task 9, 2026-05-15). // @@ -60,7 +60,7 @@ #include -extern "C" __global__ void phase_e_kill_criteria_compute_kernel( +extern "C" __global__ void alpha_kill_criteria_compute_kernel( const float* __restrict__ q_values, const int* __restrict__ action_counts, const float* __restrict__ scalar_inputs, diff --git a/crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu b/crates/ml/src/cuda_pipeline/alpha_munchausen_target.cu similarity index 97% rename from crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu rename to crates/ml/src/cuda_pipeline/alpha_munchausen_target.cu index 863878045..218572353 100644 --- a/crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu +++ b/crates/ml/src/cuda_pipeline/alpha_munchausen_target.cu @@ -1,4 +1,4 @@ -// crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu +// crates/ml/src/cuda_pipeline/alpha_munchausen_target.cu // // Phase E.1 Task 10 (2026-05-15): Munchausen DQN target term per Vieillard // et al. 2020 ("Munchausen Reinforcement Learning", arXiv:2007.14430). @@ -34,7 +34,7 @@ #include -extern "C" __global__ void phase_e_munchausen_target_kernel( +extern "C" __global__ void alpha_munchausen_target_kernel( /* Target net Q at next state — [batch, n_actions]. Boltzmann-softmaxed * here to produce V_soft(s'). */ const float* __restrict__ q_next, diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 878dd3e3c..1e9a18a2e 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -68,7 +68,7 @@ pub mod sp15_isv_slots; pub mod sp21_isv_slots; pub mod sp22_isv_slots; pub mod alpha_isv_slots; -pub(crate) mod phase_e_kernels; +pub(crate) mod alpha_kernels; pub mod lob_bar; pub use sp4_isv_slots::{ TARGET_Q_BOUND_INDEX, ATOM_POS_BOUND_BASE, WEIGHT_BOUND_BASE, diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index e8d65ae08..aeba0a39b 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2217,22 +2217,22 @@ impl StateResetRegistry { RegistryEntry { name: "alpha_q_spread_ema", category: ResetCategory::SoftReset { decay_bars: 1000 }, - description: "ISV[Q_SPREAD_EMA_INDEX=539] — alpha trading system kill-criterion #1 (Phase E.1). EMA of `std(Q, axis=action) / |mean(Q)|` averaged across the rollout batch — measures whether the policy's Q-values have meaningful action-spread (vs collapsed to a constant). Producer (eventual, Phase E.1): `phase_e_kill_criteria.cu` kernel reading per-rollout Q-values, Pearl A + Pearl D Wiener-α with α-floor 0.4 (controller co-adapts with policy per `pearl_wiener_alpha_floor_for_nonstationary`). Consumer: Phase E.1 kill-criteria gate (training proceeds only if EMA > 0.05 past epoch 10). Cold-start sentinel 0.0; SoftReset decay 1000 bars — allows a fading carry-over across rollouts but resets cleanly at fold boundary to avoid stale-regime contamination. Reset path: `write_isv_signal_at(539, 0.0)` + Wiener state zero (handled by the producer's first-observation bootstrap at next launch).", + description: "ISV[Q_SPREAD_EMA_INDEX=539] — alpha trading system kill-criterion #1 (Phase E.1). EMA of `std(Q, axis=action) / |mean(Q)|` averaged across the rollout batch — measures whether the policy's Q-values have meaningful action-spread (vs collapsed to a constant). Producer (eventual, Phase E.1): `alpha_kill_criteria.cu` kernel reading per-rollout Q-values, Pearl A + Pearl D Wiener-α with α-floor 0.4 (controller co-adapts with policy per `pearl_wiener_alpha_floor_for_nonstationary`). Consumer: Phase E.1 kill-criteria gate (training proceeds only if EMA > 0.05 past epoch 10). Cold-start sentinel 0.0; SoftReset decay 1000 bars — allows a fading carry-over across rollouts but resets cleanly at fold boundary to avoid stale-regime contamination. Reset path: `write_isv_signal_at(539, 0.0)` + Wiener state zero (handled by the producer's first-observation bootstrap at next launch).", }, RegistryEntry { name: "alpha_action_entropy_ema", category: ResetCategory::SoftReset { decay_bars: 1000 }, - description: "ISV[ACTION_ENTROPY_EMA_INDEX=540] — alpha trading system kill-criterion #2 (Phase E.1). EMA of `H(action_dist)` from the rollout's action histogram — measures whether the policy explores meaningfully (vs deterministic-collapse to one action). Producer (eventual): `phase_e_kill_criteria.cu`, Pearl A bootstrap + fixed-α 0.01 (a stable per-rollout entropy doesn't need adaptive α; SP-chain pattern matches `REWARD_CAP_EMA_ALPHA=0.01` per `sp14_isv_slots.rs:158`). Consumer: kill-criteria gate (entropy > 0.5·ln(9) past epoch 10). Cold-start sentinel 0.0; SoftReset 1000 bars — same rationale as 539. The 9 in `ln(9)` is `N_ACTIONS` from `crates/ml/src/env/action_space.rs`.", + description: "ISV[ACTION_ENTROPY_EMA_INDEX=540] — alpha trading system kill-criterion #2 (Phase E.1). EMA of `H(action_dist)` from the rollout's action histogram — measures whether the policy explores meaningfully (vs deterministic-collapse to one action). Producer (eventual): `alpha_kill_criteria.cu`, Pearl A bootstrap + fixed-α 0.01 (a stable per-rollout entropy doesn't need adaptive α; SP-chain pattern matches `REWARD_CAP_EMA_ALPHA=0.01` per `sp14_isv_slots.rs:158`). Consumer: kill-criteria gate (entropy > 0.5·ln(9) past epoch 10). Cold-start sentinel 0.0; SoftReset 1000 bars — same rationale as 539. The 9 in `ln(9)` is `N_ACTIONS` from `crates/ml/src/env/action_space.rs`.", }, RegistryEntry { name: "alpha_return_vs_random_ema", category: ResetCategory::SoftReset { decay_bars: 1000 }, - description: "ISV[RETURN_VS_RANDOM_EMA_INDEX=541] — alpha trading system kill-criterion #3 (Phase E.1). EMA of `(rollout_R − random_R) / σ_random` where `random_R` and `σ_random` are read from slots 547/548 (`RANDOM_BASELINE_MEAN_INDEX` and `RANDOM_BASELINE_STD_INDEX`) populated once per fold by the Phase E.0 baseline computation. Measures signed σ-distance of the trained policy's reward distribution from the random-uniform baseline — kill if not > 2.0 past epoch 10. Producer (eventual): `phase_e_kill_criteria.cu`, Pearl A bootstrap + fixed-α 0.05 (per-rollout signal-to-noise warrants a slightly faster EMA than entropy). Consumer: kill-criteria gate. Cold-start sentinel 0.0; SoftReset 1000 bars.", + description: "ISV[RETURN_VS_RANDOM_EMA_INDEX=541] — alpha trading system kill-criterion #3 (Phase E.1). EMA of `(rollout_R − random_R) / σ_random` where `random_R` and `σ_random` are read from slots 547/548 (`RANDOM_BASELINE_MEAN_INDEX` and `RANDOM_BASELINE_STD_INDEX`) populated once per fold by the Phase E.0 baseline computation. Measures signed σ-distance of the trained policy's reward distribution from the random-uniform baseline — kill if not > 2.0 past epoch 10. Producer (eventual): `alpha_kill_criteria.cu`, Pearl A bootstrap + fixed-α 0.05 (per-rollout signal-to-noise warrants a slightly faster EMA than entropy). Consumer: kill-criteria gate. Cold-start sentinel 0.0; SoftReset 1000 bars.", }, RegistryEntry { name: "alpha_early_q_movement_ema", category: ResetCategory::SoftReset { decay_bars: 1000 }, - description: "ISV[EARLY_Q_MOVEMENT_EMA_INDEX=542] — alpha trading system kill-criterion #4 (Phase E.1). EMA of `|Q(s_early) − Q_init| / |Q_init|` measuring whether terminal-reward signal at H=6000 propagates back to early-step Q-values (the diagnostic against pure TD-bootstrap failure — see `pearl_horizon_decay_falsifies_long_horizon_stateless` for the upstream finding). Producer (eventual): `phase_e_kill_criteria.cu` per-epoch (not per-step), reads a held-out set of early-step states + the snapshot of init-Q values. Consumer: kill-criteria gate (early movement > 0.01 past epoch 10). Cold-start sentinel 0.0; SoftReset 1000 bars — but in practice this slot is read epoch-by-epoch so the decay rarely matters.", + description: "ISV[EARLY_Q_MOVEMENT_EMA_INDEX=542] — alpha trading system kill-criterion #4 (Phase E.1). EMA of `|Q(s_early) − Q_init| / |Q_init|` measuring whether terminal-reward signal at H=6000 propagates back to early-step Q-values (the diagnostic against pure TD-bootstrap failure — see `pearl_horizon_decay_falsifies_long_horizon_stateless` for the upstream finding). Producer (eventual): `alpha_kill_criteria.cu` per-epoch (not per-step), reads a held-out set of early-step states + the snapshot of init-Q values. Consumer: kill-criteria gate (early movement > 0.01 past epoch 10). Cold-start sentinel 0.0; SoftReset 1000 bars — but in practice this slot is read epoch-by-epoch so the decay rarely matters.", }, RegistryEntry { name: "alpha_stacker_threshold", diff --git a/docs/isv-slots.md b/docs/isv-slots.md index dc49e9353..e28838922 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -528,7 +528,7 @@ Phase E intro commit reserves 12 contiguous slots at `[539..551)` for **system-s - `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `ISV_TOTAL_DIM: 539 → 551`; comment updated to reference `alpha_isv_slots.rs` **Producers/consumers status:** intro commit is reservation-only. Producers land later in Phase E: -- Phase E.1: `phase_e_kill_criteria.cu` writes slots 539-542 +- Phase E.1: `alpha_kill_criteria.cu` writes slots 539-542 - Phase E.2: `stacker_threshold_controller.cu` writes slots 543, 545, 546 (reads 544 as anchor) - Phase E.0: host writes 544 (trade-rate target) at training start and 547, 548 (baselines) per fold @@ -550,7 +550,7 @@ All 10 registry tests + the global registry coverage test pass: `cargo test -p m ## Alpha block — Phase E Task 9 producer kernel (2026-05-15) -Phase E.1 Task 9 lands the producer for slots 539-542 (`phase_e_kill_criteria.cu`). Single-block, single-thread kernel; writes 4 raw scalar observations into `scratch_out[0..4]`, then the existing `apply_pearls_ad_kernel` (chained on the same stream, `n_slots=4`) applies Pearl A bootstrap + Pearl D Wiener-α smoothing into the four ISV slots: +Phase E.1 Task 9 lands the producer for slots 539-542 (`alpha_kill_criteria.cu`). Single-block, single-thread kernel; writes 4 raw scalar observations into `scratch_out[0..4]`, then the existing `apply_pearls_ad_kernel` (chained on the same stream, `n_slots=4`) applies Pearl A bootstrap + Pearl D Wiener-α smoothing into the four ISV slots: - `scratch_out[0]` → ISV[539] `Q_SPREAD_EMA_INDEX` — mean over batch of `std(Q, axis=action) / |mean(Q)|` - `scratch_out[1]` → ISV[540] `ACTION_ENTROPY_EMA_INDEX` — `H(empirical action distribution)` in nats @@ -560,12 +560,12 @@ Phase E.1 Task 9 lands the producer for slots 539-542 (`phase_e_kill_criteria.cu Kernel reads ISV[547]/ISV[548] (random baseline mean/std, populated once by Task 7c, TrainingPersist) via slot indices passed as kernel args. Composability per `val_sharpe_delta_compute_kernel` — no inline Wiener math here, the canonical `apply_pearls_ad_kernel` handles it. **Files touched:** -- `crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu` (new) — producer kernel +- `crates/ml/src/cuda_pipeline/alpha_kill_criteria.cu` (new) — producer kernel - `crates/ml/build.rs` — added kernel to `kernels_with_common` list -**Cubin:** `target/release/build/ml-*/out/phase_e_kill_criteria.cubin` (~16.8 KB). +**Cubin:** `target/release/build/ml-*/out/alpha_kill_criteria.cubin` (~16.8 KB). -**Launcher wiring:** lands in Phase E.1 Task 11 (the DQN training loop reads `phase_e_kill_criteria_compute_kernel` from the compiled cubin and chains `apply_pearls_ad_kernel` on the same stream after it). +**Launcher wiring:** lands in Phase E.1 Task 11 (the DQN training loop reads `alpha_kill_criteria_compute_kernel` from the compiled cubin and chains `apply_pearls_ad_kernel` on the same stream after it). ## Phase E.1 Task 10 — Munchausen target kernel (2026-05-15) @@ -581,27 +581,27 @@ Both softmaxes use log-sum-exp via the max-trick (essential at τ ≈ 0.03 where Does NOT touch any ISV slot directly — α_m, τ, and log_clip_min are exposed as kernel args so a downstream controller (Phase E.2+) can ISV-drive them. **Files touched:** -- `crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu` (new) +- `crates/ml/src/cuda_pipeline/alpha_munchausen_target.cu` (new) - `crates/ml/build.rs` — kernel added to `kernels_with_common` -**Cubin:** `target/release/build/ml-*/out/phase_e_munchausen_target.cubin` (~12.8 KB). +**Cubin:** `target/release/build/ml-*/out/alpha_munchausen_target.cubin` (~12.8 KB). **Wiring:** lands in Phase E.1 Task 11. The integration point is wherever the existing C51/MSE loss kernels currently consume `r + γ · V_next` — they'll consume `target_out` from this kernel instead. -## Phase E.1 Task 11 — Rust launchers `phase_e_kernels.rs` (2026-05-15) +## Phase E.1 Task 11 — Rust launchers `alpha_kernels.rs` (2026-05-15) -`crates/ml/src/cuda_pipeline/phase_e_kernels.rs` exposes two `pub(crate)` launchers: +`crates/ml/src/cuda_pipeline/alpha_kernels.rs` exposes two `pub(crate)` launchers: -- **`launch_phase_e_kill_criteria(...)`** — calls `phase_e_kill_criteria_compute_kernel` (cubin from Task 9). Grid (1, 1, 1) × Block (1, 1, 1). Reads `RANDOM_BASELINE_MEAN_INDEX` (547) and `RANDOM_BASELINE_STD_INDEX` (548) from `alpha_isv_slots` and passes them as kernel args. Caller is responsible for chaining `apply_pearls_ad_kernel(n_slots=4, isv_idx_base=539)` afterward to smooth raw scratch outputs into slots 539..542. +- **`launch_alpha_kill_criteria(...)`** — calls `alpha_kill_criteria_compute_kernel` (cubin from Task 9). Grid (1, 1, 1) × Block (1, 1, 1). Reads `RANDOM_BASELINE_MEAN_INDEX` (547) and `RANDOM_BASELINE_STD_INDEX` (548) from `alpha_isv_slots` and passes them as kernel args. Caller is responsible for chaining `apply_pearls_ad_kernel(n_slots=4, isv_idx_base=539)` afterward to smooth raw scratch outputs into slots 539..542. -- **`launch_phase_e_munchausen_target(...)`** — calls `phase_e_munchausen_target_kernel` (cubin from Task 10). Grid sized `ceil(batch/256) × 1 × 1`, Block (256, 1, 1). `α_m`, `τ`, `log_clip_min` are scalar args so the smoke / future controller can ISV-drive them. +- **`launch_alpha_munchausen_target(...)`** — calls `alpha_munchausen_target_kernel` (cubin from Task 10). Grid sized `ceil(batch/256) × 1 × 1`, Block (256, 1, 1). `α_m`, `τ`, `log_clip_min` are scalar args so the smoke / future controller can ISV-drive them. Both launchers follow `launch_apply_pearls` (in `sp4_wiener_ema.rs`) — pre-loaded `CudaFunction` passed as parameter, `u64` device pointers, `debug_assert!` guards. **Plan deviation rationale:** the plan's Task 11 said "replace hardcoded `n_step=32`, `gamma=0.999` literals" but grep across the trainer found zero such literals — the trainer already reads γ via `read_isv_signal_at(GAMMA_DIR_EFF_INDEX)` and ε via `read_isv_signal_at(AUX_TRUNK_EPS_INDEX)`. The actual missing piece for E.1 was Rust-side launchers for the new cubins, delivered here. **Files touched:** -- `crates/ml/src/cuda_pipeline/phase_e_kernels.rs` (new) — both launchers, `pub(crate)` visibility -- `crates/ml/src/cuda_pipeline/mod.rs` — registers `pub(crate) mod phase_e_kernels;` +- `crates/ml/src/cuda_pipeline/alpha_kernels.rs` (new) — both launchers, `pub(crate)` visibility +- `crates/ml/src/cuda_pipeline/mod.rs` — registers `pub(crate) mod alpha_kernels;` -**Tests:** `cargo test -p ml --lib phase_e_kernels` — 1 compile-witness passes. Real GPU integration test lands in Phase E.1 Task 12 (kernel smoke binary). +**Tests:** `cargo test -p ml --lib alpha_kernels` — 1 compile-witness passes. Real GPU integration test lands in Phase E.1 Task 12 (kernel smoke binary).