refactor(alpha): rename phase_e_* → alpha_* — system-scoped naming
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
This commit is contained in:
@@ -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)
|
||||
|
||||
@@ -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)]
|
||||
@@ -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,
|
||||
}
|
||||
|
||||
@@ -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 |
|
||||
|
||||
@@ -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<f32> = vec![1.0, 2.0, 3.0, 0.0, 0.0, 0.0];
|
||||
let q_next: Vec<f32> = vec![10.0, 20.0, 30.0, 0.0, 0.0, 0.0];
|
||||
let actions: Vec<i32> = vec![0, 1];
|
||||
let rewards: Vec<f32> = vec![1.0, 2.0];
|
||||
let dones: Vec<f32> = 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::<f32>(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:?})"
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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,
|
||||
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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,
|
||||
@@ -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,
|
||||
|
||||
@@ -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",
|
||||
|
||||
Reference in New Issue
Block a user