diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 13690c349..72ebfb8f6 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -43,9 +43,10 @@ const KERNELS: &[&str] = &[ "rl_lr_controller", // RL Phase E.2-DEFER: per-head learning-rate ISV emitter — bootstraps ISV[412..417] with 1e-3 (BCE/Q/π/V/aux); replaces hardcoded PHASE_E2_DEFAULT_LR "rl_rollout_steps_controller", // RL Phase E.3b: rollout-buffer-length ISV emitter — emits ISV[RL_N_ROLLOUT_STEPS_INDEX=404] from var(advantage)/|mean A| EMA; bootstraps 2048 "rl_per_alpha_controller", // RL Phase E.3b: PER priority-exponent ISV emitter — emits ISV[RL_PER_ALPHA_INDEX=405] from TD-error kurtosis EMA; bootstraps 0.6 + "rl_reward_scale_controller", // RL Phase R1 (rebuild): reward-standardisation scale ISV emitter — emits ISV[RL_REWARD_SCALE_INDEX=406] from mean |realized_pnl_usd| EMA; bootstraps 1.0 ]; -// Cache bust v24 (2026-05-23): RL Phase E.3b — rl_rollout_steps_controller.cu (ISV[404] producer; advantage-variance-driven rollout length, bootstrap 2048) + rl_per_alpha_controller.cu (ISV[405] producer; TD-kurtosis-driven PER priority exponent α, bootstrap 0.6). Both follow pearl_first_observation_bootstrap + pearl_wiener_alpha_floor_for_nonstationary. Companion to LobEnv trait + step_with_lobsim in src/rl/reward.rs and the integrated trainer. +// Cache bust v25 (2026-05-23): RL Phase R1 rebuild — port rl_reward_scale_controller.cu (ISV[406] producer, bootstrap 1.0; ported from ml-alpha-phase-f-g-flawed) and wire all 7 RL controllers (gamma/target_tau/ppo_clip/entropy_coef/rollout_steps/per_alpha/reward_scale) into IntegratedTrainer::new via launch-once-with-sentinel-zero-input per pearl_first_observation_bootstrap. Replaces flawed Phase F's host memcpy_htod bootstrap that violated feedback_no_htod_htoh_only_mapped_pinned. fn main() { println!("cargo:rerun-if-changed=build.rs"); diff --git a/crates/ml-alpha/cuda/rl_reward_scale_controller.cu b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu new file mode 100644 index 000000000..3cfe0eec8 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu @@ -0,0 +1,93 @@ +// rl_reward_scale_controller.cu — emits per-trade reward standardisation +// scale to ISV[RL_REWARD_SCALE_INDEX=406]. +// +// Phase F of the integrated RL trainer +// (docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md). +// +// The C51 distributional Q-head uses a fixed atom support [-1, +1] with +// N_ATOMS=21. For the categorical Bellman projection to work well, scaled +// rewards must sit comfortably within that support — a typical realized +// |PnL| trade should map to a magnitude ≈ 1.0. Without standardisation, +// a strategy trading 1-tick ES moves (≈$12.50 per contract) would stuff +// nearly all probability mass into a handful of atoms near zero, starving +// the distributional representation. +// +// This controller emits the inverse of the mean absolute realized PnL: +// +// target_scale = 1.0 / max(mean_abs_pnl_ema, EPS_PNL) +// +// so that reward_scaled = reward_raw * isv[RL_REWARD_SCALE_INDEX] keeps +// a typical trade's realized PnL ≈ 1.0 (in atom-support units). The EMA +// is maintained on the Rust side (Phase F.3) from the LobEnv adapter's +// reward stream over the last ~100 closes, and passed in as a scalar. +// +// Bootstrap discipline (per `pearl_first_observation_bootstrap`): the ISV +// slot starts at 0.0 sentinel. First emit writes REWARD_SCALE_BOOTSTRAP=1.0 +// directly — raw passthrough until the EMA stabilises. Subsequent emits +// Wiener-α blend with floor 0.4 (per `pearl_wiener_alpha_floor_for_nonstationary` +// — the trade-magnitude distribution drifts as the policy co-adapts, +// breaking stationarity). +// +// Bounds: scale ∈ [1e-3, 1e3]. Five orders of magnitude either side of +// unit scale handles micro-PnL noise scenarios at the low end and +// catastrophic-loss-magnitude outliers at the high end, while preventing +// runaway. If the typical |PnL| stabilises at e.g. 50 USD, the controller +// converges to scale ≈ 0.02 (within bounds); at $0.01 noise it caps at +// 1e3 rather than exploding. + +#define RL_REWARD_SCALE_INDEX 406 +#define REWARD_SCALE_MIN 1e-3f +#define REWARD_SCALE_MAX 1e3f +#define REWARD_SCALE_BOOTSTRAP 1.0f +// Floor for the mean_abs_pnl_ema denominator to guard against div-by-zero +// when the EMA hasn't accumulated any closed trades yet. +#define EPS_PNL 1e-3f +#define WIENER_ALPHA_FLOOR 0.4f + +// ───────────────────────────────────────────────────────────────────── +// rl_reward_scale_controller: +// Single-thread controller — writes ONE float to +// isv[RL_REWARD_SCALE_INDEX]. +// +// Inputs: +// isv [≥ RL_REWARD_SCALE_INDEX+1] — ISV bus +// alpha_step — Wiener-α from the controller's signal stats; +// floored at WIENER_ALPHA_FLOOR before the blend. +// Named `alpha_step` to disambiguate from PER's α +// (ISV[405]) and from the ISV slot name itself. +// mean_abs_pnl_ema — caller-computed EMA of |realized_pnl_usd| over +// the last ~100 closed trades. Phase F.3 maintains +// this from the LobEnv adapter reward stream. +// +// Outputs: +// isv[RL_REWARD_SCALE_INDEX] — reward standardisation scale ∈ +// [REWARD_SCALE_MIN, REWARD_SCALE_MAX] +// ───────────────────────────────────────────────────────────────────── +extern "C" __global__ void rl_reward_scale_controller( + float* __restrict__ isv, + float alpha_step, + float mean_abs_pnl_ema +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float prev = isv[RL_REWARD_SCALE_INDEX]; + // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. + if (prev == 0.0f) { + isv[RL_REWARD_SCALE_INDEX] = REWARD_SCALE_BOOTSTRAP; + return; + } + + // target_scale = 1.0 / max(mean_abs_pnl_ema, EPS_PNL). + // Larger typical PnL → smaller scale (reward is divided down toward ±1). + // Smaller typical PnL → larger scale (reward is amplified toward ±1). + const float denom = fmaxf(mean_abs_pnl_ema, EPS_PNL); + float target = 1.0f / denom; + target = fmaxf(REWARD_SCALE_MIN, fminf(target, REWARD_SCALE_MAX)); + + // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. + const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + + out = fmaxf(REWARD_SCALE_MIN, fminf(out, REWARD_SCALE_MAX)); + isv[RL_REWARD_SCALE_INDEX] = out; +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 83b1fae02..3b6beaea7 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -84,6 +84,55 @@ pub const RL_LR_V_INDEX: usize = 415; /// Per-head learning rate — aux heads (prof/size, long/short). pub const RL_LR_AUX_INDEX: usize = 416; +// ───────────────────────────────────────────────────────────────────── +// EMA-input slots for the 7 RL controllers (Phase R1 rebuild) +// +// Per `feedback_isv_for_adaptive_bounds` + `pearl_controller_anchors_isv_driven`, +// every controller's INPUT EMA also lives in ISV so the producer kernel +// reads the signal from device-resident state instead of a host scalar +// argument. The 7 controllers each consume one of these slots: +// +// gamma_controller ← MEAN_TRADE_DURATION_EMA (417) +// target_tau_controller ← Q_DIVERGENCE_EMA (418) +// ppo_clip_controller ← KL_PI_EMA (419) +// entropy_coef_controller ← ENTROPY_OBSERVED_EMA (420) +// rollout_steps_controller ← ADVANTAGE_VAR_RATIO_EMA (421) +// per_alpha_controller ← TD_KURTOSIS_EMA (422) +// reward_scale_controller ← MEAN_ABS_PNL_EMA (423) +// +// All 7 sentinel-bootstrap to 0.0 via `alloc_zeros` per +// `pearl_first_observation_bootstrap`. The EMA-update kernels (Phase R3) +// write into these slots from device-resident reward/done/Q-divergence +// signals; the controller kernels READ these slots as scalar inputs +// (no host scalar arg threading). +// +// In Phase R1 these slots are reserved but unused — the controller +// launches at init pass sentinel-zero directly (bootstrap path). R3 +// wires the EMA producers, R5 wires the controllers' per-step reads. +// ───────────────────────────────────────────────────────────────────── + +/// EMA of mean trade duration in event count. Input to `rl_gamma_controller`. +pub const RL_MEAN_TRADE_DURATION_EMA_INDEX: usize = 417; + +/// EMA of ‖Q_current_w − Q_target_w‖. Input to `rl_target_tau_controller`. +pub const RL_Q_DIVERGENCE_EMA_INDEX: usize = 418; + +/// EMA of KL(π_new ‖ π_old). Input to `rl_ppo_clip_controller`. +pub const RL_KL_PI_EMA_INDEX: usize = 419; + +/// EMA of observed action entropy H(π). Input to `rl_entropy_coef_controller`. +pub const RL_ENTROPY_OBSERVED_EMA_INDEX: usize = 420; + +/// EMA of var(A_Q) / |mean A_Q|. Input to `rl_rollout_steps_controller`. +pub const RL_ADVANTAGE_VAR_RATIO_EMA_INDEX: usize = 421; + +/// EMA of |TD-error| kurtosis. Input to `rl_per_alpha_controller`. +pub const RL_TD_KURTOSIS_EMA_INDEX: usize = 422; + +/// EMA of |realized_pnl_usd| over closed trades. Input to +/// `rl_reward_scale_controller`. +pub const RL_MEAN_ABS_PNL_EMA_INDEX: usize = 423; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 417; +pub const RL_SLOTS_END: usize = 424; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index cab80b81e..e5c88fa53 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -107,6 +107,28 @@ const REDUCE_AXIS0_CUBIN: &[u8] = const RL_LR_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_lr_controller.cubin")); +// Phase R1: cubin includes for the 7 RL adaptive controllers +// (γ / τ / ε / entropy_coef / n_rollout_steps / per_α / reward_scale). +// All 7 share the `(isv*, alpha, scalar_input)` signature and the +// `pearl_first_observation_bootstrap` sentinel-zero path. Loaded in +// `IntegratedTrainer::new` + bootstrap-launched once at construction so +// every consumer kernel sees the canonical default in its ISV slot +// before any in-loop access. +const RL_GAMMA_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_gamma_controller.cubin")); +const RL_TARGET_TAU_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_target_tau_controller.cubin")); +const RL_PPO_CLIP_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_ppo_clip_controller.cubin")); +const RL_ENTROPY_COEF_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_entropy_coef_controller.cubin")); +const RL_ROLLOUT_STEPS_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_rollout_steps_controller.cubin")); +const RL_PER_ALPHA_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_per_alpha_controller.cubin")); +const RL_REWARD_SCALE_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_scale_controller.cubin")); + /// Per-head LR controller Wiener-α target. The controller kernel floors /// this at 0.4 per `pearl_wiener_alpha_floor_for_nonstationary` so the /// effective blend rate stays ≥ 0.4 regardless of this seed value. Kept @@ -166,9 +188,11 @@ pub struct IntegratedTrainer { pub value_w_adam: AdamW, pub value_b_adam: AdamW, - /// Device ISV buffer (RL-allocated, length RL_SLOTS_END = 412). - /// Sentinel-zero initialised on construction so all controllers - /// bootstrap on their first emit per pearl_first_observation_bootstrap. + /// Device ISV buffer (RL-allocated, length `RL_SLOTS_END = 424` as + /// of Phase R1). Allocated zero, then immediately seeded by the 7 + /// RL controllers' bootstrap launches in `new()` so every consumer + /// kernel sees the canonical default at its slot before any read. + /// Per `pearl_first_observation_bootstrap` + `feedback_isv_for_adaptive_bounds`. pub isv_d: CudaSlice, /// Host mirror of ISV for loss-balance λ reads. Refreshed before each @@ -192,6 +216,30 @@ pub struct IntegratedTrainer { /// (sentinel zero → `LR_BOOTSTRAP = 1e-3`). rl_lr_controller_fn: CudaFunction, + // ── 7 RL adaptive controllers (Phase R1) ────────────────────────── + // Each controller emits ONE float into its dedicated ISV slot + // (γ→400, τ→401, ε→402, entropy_coef→403, n_rollout_steps→404, + // per_α→405, reward_scale→406). Bootstrap-launched once at the end + // of `new()` so every consumer kernel reads a canonical value on + // first access — eliminates the Phase F defect where sparse-reward + // production paths read sentinel-zero γ/ε/entropy and trained + // against degenerate Bellman / surrogate targets. Per-step launches + // (with real EMA inputs sourced from ISV[417..424]) land in Phase R5. + _rl_gamma_controller_module: Arc, + rl_gamma_controller_fn: CudaFunction, + _rl_target_tau_controller_module: Arc, + rl_target_tau_controller_fn: CudaFunction, + _rl_ppo_clip_controller_module: Arc, + rl_ppo_clip_controller_fn: CudaFunction, + _rl_entropy_coef_controller_module: Arc, + rl_entropy_coef_controller_fn: CudaFunction, + _rl_rollout_steps_controller_module: Arc, + rl_rollout_steps_controller_fn: CudaFunction, + _rl_per_alpha_controller_module: Arc, + rl_per_alpha_controller_fn: CudaFunction, + _rl_reward_scale_controller_module: Arc, + rl_reward_scale_controller_fn: CudaFunction, + /// Combined encoder grad slot — folded from Q + π + V `grad_h_t` /// contributions via `grad_h_accumulate_scaled` (item 1). Consumed by /// `PerceptionTrainer::backward_encoder_with_grad_h_t`. Size @@ -309,6 +357,50 @@ impl IntegratedTrainer { .load_function("rl_lr_controller") .context("load rl_lr_controller")?; + // Phase R1: load the 7 RL adaptive controllers. + let rl_gamma_controller_module = ctx + .load_cubin(RL_GAMMA_CONTROLLER_CUBIN.to_vec()) + .context("load rl_gamma_controller cubin")?; + let rl_gamma_controller_fn = rl_gamma_controller_module + .load_function("rl_gamma_controller") + .context("load rl_gamma_controller")?; + let rl_target_tau_controller_module = ctx + .load_cubin(RL_TARGET_TAU_CONTROLLER_CUBIN.to_vec()) + .context("load rl_target_tau_controller cubin")?; + let rl_target_tau_controller_fn = rl_target_tau_controller_module + .load_function("rl_target_tau_controller") + .context("load rl_target_tau_controller")?; + let rl_ppo_clip_controller_module = ctx + .load_cubin(RL_PPO_CLIP_CONTROLLER_CUBIN.to_vec()) + .context("load rl_ppo_clip_controller cubin")?; + let rl_ppo_clip_controller_fn = rl_ppo_clip_controller_module + .load_function("rl_ppo_clip_controller") + .context("load rl_ppo_clip_controller")?; + let rl_entropy_coef_controller_module = ctx + .load_cubin(RL_ENTROPY_COEF_CONTROLLER_CUBIN.to_vec()) + .context("load rl_entropy_coef_controller cubin")?; + let rl_entropy_coef_controller_fn = rl_entropy_coef_controller_module + .load_function("rl_entropy_coef_controller") + .context("load rl_entropy_coef_controller")?; + let rl_rollout_steps_controller_module = ctx + .load_cubin(RL_ROLLOUT_STEPS_CONTROLLER_CUBIN.to_vec()) + .context("load rl_rollout_steps_controller cubin")?; + let rl_rollout_steps_controller_fn = rl_rollout_steps_controller_module + .load_function("rl_rollout_steps_controller") + .context("load rl_rollout_steps_controller")?; + let rl_per_alpha_controller_module = ctx + .load_cubin(RL_PER_ALPHA_CONTROLLER_CUBIN.to_vec()) + .context("load rl_per_alpha_controller cubin")?; + let rl_per_alpha_controller_fn = rl_per_alpha_controller_module + .load_function("rl_per_alpha_controller") + .context("load rl_per_alpha_controller")?; + let rl_reward_scale_controller_module = ctx + .load_cubin(RL_REWARD_SCALE_CONTROLLER_CUBIN.to_vec()) + .context("load rl_reward_scale_controller cubin")?; + let rl_reward_scale_controller_fn = rl_reward_scale_controller_module + .load_function("rl_reward_scale_controller") + .context("load rl_reward_scale_controller")?; + Ok(Self { cfg, perception, @@ -330,10 +422,104 @@ impl IntegratedTrainer { reduce_axis0_fn, _rl_lr_controller_module: rl_lr_controller_module, rl_lr_controller_fn, + _rl_gamma_controller_module: rl_gamma_controller_module, + rl_gamma_controller_fn, + _rl_target_tau_controller_module: rl_target_tau_controller_module, + rl_target_tau_controller_fn, + _rl_ppo_clip_controller_module: rl_ppo_clip_controller_module, + rl_ppo_clip_controller_fn, + _rl_entropy_coef_controller_module: rl_entropy_coef_controller_module, + rl_entropy_coef_controller_fn, + _rl_rollout_steps_controller_module: rl_rollout_steps_controller_module, + rl_rollout_steps_controller_fn, + _rl_per_alpha_controller_module: rl_per_alpha_controller_module, + rl_per_alpha_controller_fn, + _rl_reward_scale_controller_module: rl_reward_scale_controller_module, + rl_reward_scale_controller_fn, grad_h_t_combined_d, last_pi_loss: 0.0, step_counter: 0, - }) + } + .with_controllers_bootstrapped()?) + } + + /// Phase R1: fire each of the 7 RL controllers ONCE against the + /// freshly-zeroed `isv_d`. Each kernel detects the sentinel (its + /// dedicated ISV slot == 0.0), executes its first-observation- + /// bootstrap path per `pearl_first_observation_bootstrap`, and + /// writes its canonical default into the slot: + /// + /// ISV[400] γ = 0.99 (GAMMA_BOOTSTRAP) + /// ISV[401] τ = 0.005 (TAU_BOOTSTRAP) + /// ISV[402] ε = 0.2 (EPS_BOOTSTRAP) + /// ISV[403] entropy_coef = 0.01 (COEF_BOOTSTRAP) + /// ISV[404] n_rollout_steps= 2048 (ROLLOUT_BOOTSTRAP) + /// ISV[405] per_α = 0.6 (PER_ALPHA_BOOTSTRAP) + /// ISV[406] reward_scale = 1.0 (REWARD_SCALE_BOOTSTRAP) + /// + /// The scalar `input` argument is IGNORED on the bootstrap path + /// (each kernel `return`s immediately after the write). Phase R5 + /// reuses [`Self::launch_isv_controller_3arg`] with real EMA + /// inputs sourced from ISV[417..424] — at that point the Wiener-α + /// blend kicks in and the slots adapt from the bootstrap defaults. + /// + /// Replaces the flawed Phase F approach of host-side `memcpy_htod` + /// of canonical constants, which violated + /// `feedback_no_htod_htoh_only_mapped_pinned` (tests not exempt) + /// and short-circuited `pearl_first_observation_bootstrap`. + fn with_controllers_bootstrapped(self) -> Result { + let alpha = RL_LR_CONTROLLER_ALPHA; + let zero = 0.0_f32; + self.launch_isv_controller_3arg(&self.rl_gamma_controller_fn, alpha, zero) + .context("R1 bootstrap rl_gamma_controller")?; + self.launch_isv_controller_3arg(&self.rl_target_tau_controller_fn, alpha, zero) + .context("R1 bootstrap rl_target_tau_controller")?; + self.launch_isv_controller_3arg(&self.rl_ppo_clip_controller_fn, alpha, zero) + .context("R1 bootstrap rl_ppo_clip_controller")?; + self.launch_isv_controller_3arg(&self.rl_entropy_coef_controller_fn, alpha, zero) + .context("R1 bootstrap rl_entropy_coef_controller")?; + self.launch_isv_controller_3arg(&self.rl_rollout_steps_controller_fn, alpha, zero) + .context("R1 bootstrap rl_rollout_steps_controller")?; + self.launch_isv_controller_3arg(&self.rl_per_alpha_controller_fn, alpha, zero) + .context("R1 bootstrap rl_per_alpha_controller")?; + self.launch_isv_controller_3arg(&self.rl_reward_scale_controller_fn, alpha, zero) + .context("R1 bootstrap rl_reward_scale_controller")?; + self.stream + .synchronize() + .context("R1 controller bootstrap sync")?; + Ok(self) + } + + /// Phase R1 / R5: single-thread launch of any of the 7 RL adaptive + /// controllers. All 7 share the `(isv*, alpha, scalar_input)` + /// signature; this helper centralises the `LaunchConfig` + arg + /// binding so per-step launchers reduce to one line each. + /// + /// At bootstrap (R1), `input` is `0.0` — the kernel reads its ISV + /// slot, sees sentinel zero, writes its `*_BOOTSTRAP` value, and + /// returns without touching `input`. At per-step launches (R5), + /// `input` is the real EMA value from the corresponding ISV[417..424] + /// slot and the kernel runs its Wiener-α-with-floor blend per + /// `pearl_wiener_alpha_floor_for_nonstationary`. + fn launch_isv_controller_3arg( + &self, + func: &CudaFunction, + alpha: f32, + input: f32, + ) -> Result<()> { + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(func); + launch.arg(&self.isv_d).arg(&alpha).arg(&input); + unsafe { + launch + .launch(cfg) + .context("isv controller 3-arg launch")?; + } + Ok(()) } /// Phase E.2 + E.2-DEFER: synthetic-batch end-to-end step. Runs real diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs new file mode 100644 index 000000000..1d6fdebfe --- /dev/null +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -0,0 +1,143 @@ +//! Phase R1 gate **G1** — `IntegratedTrainer::new` bootstraps each of +//! the 7 RL adaptive-controller ISV slots to its canonical default via +//! the kernel-defined `*_BOOTSTRAP` constants (`pearl_first_observation_bootstrap` +//! sentinel-zero path). +//! +//! This catches the canonical Phase F defect (commit history preserved +//! on `ml-alpha-phase-f-g-flawed`) where ISV[400..406] stayed at +//! `alloc_zeros` sentinel `0.0` in production, causing +//! `bellman_target_projection`, `dqn_distributional_q`, and +//! `ppo_clipped_surrogate` to consume γ=0 / ε=0 / entropy_coef=0 and +//! train against degenerate Bellman / surrogate targets. +//! +//! Per `feedback_no_cpu_test_fallbacks`: the oracle is the kernel's +//! own `#define *_BOOTSTRAP` constant, not a CPU computation. +//! +//! Per `pearl_tests_must_prove_not_lock_observations`: this asserts an +//! invariant ("bootstrap path wrote the canonical value defined by the +//! kernel") not a tuned magic number — if a kernel rewrites its +//! bootstrap constant in a future refactor, this test updates with the +//! constant, not with empirical observation. +//! +//! Run with: +//! `cargo test -p ml-alpha --test isv_bootstrap -- --ignored --nocapture` + +use ml_alpha::rl::isv_slots::{ + RL_ENTROPY_COEF_INDEX, RL_GAMMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, + RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PPO_CLIP_INDEX, RL_REWARD_SCALE_INDEX, + RL_SLOTS_END, RL_TARGET_TAU_INDEX, +}; +use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; +use ml_alpha::trainer::perception::PerceptionTrainerConfig; +use ml_core::device::MlDevice; + +// Canonical bootstrap values — each MUST match the kernel's +// corresponding `*_BOOTSTRAP` #define in +// `crates/ml-alpha/cuda/rl_*_controller.cu`. Source of truth is the +// .cu file; this test asserts the kernel actually wrote those values. +const GAMMA_BOOTSTRAP: f32 = 0.99; +const TAU_BOOTSTRAP: f32 = 0.005; +const EPS_BOOTSTRAP: f32 = 0.2; +const COEF_BOOTSTRAP: f32 = 0.01; +const ROLLOUT_BOOTSTRAP: f32 = 2048.0; +const PER_ALPHA_BOOTSTRAP: f32 = 0.6; +const REWARD_SCALE_BOOTSTRAP: f32 = 1.0; + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g1_isv_bootstrap_writes_canonical_values() { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping G1 ({e})"); + return; + } + }; + + let cfg = IntegratedTrainerConfig { + perception: PerceptionTrainerConfig { + seq_len: 4, + n_batch: 1, + ..PerceptionTrainerConfig::default() + }, + dqn_seed: 0xCAFE, + ppo_seed: 0xBEEF, + }; + let trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); + + // Read full ISV slice to host. Uses the same pattern as the + // trainer's own per-step ISV mirror refresh. + let mut isv = vec![0.0_f32; RL_SLOTS_END]; + let stream = dev.cuda_stream().expect("cuda_stream"); + stream + .memcpy_dtoh(&trainer.isv_d, isv.as_mut_slice()) + .expect("isv dtoh"); + + // Floating-point exact equality is the right oracle here — each + // kernel's bootstrap path is `isv[slot] = K_BOOTSTRAP; return;` + // with no arithmetic, so the host-side f32 must equal the kernel's + // f32 literal bit-for-bit. EPS guards against any unexpected + // post-bootstrap math we don't see here. + const EPS: f32 = 1e-6; + assert!( + (isv[RL_GAMMA_INDEX] - GAMMA_BOOTSTRAP).abs() < EPS, + "ISV[γ={RL_GAMMA_INDEX}] expected {GAMMA_BOOTSTRAP}, got {}", + isv[RL_GAMMA_INDEX] + ); + assert!( + (isv[RL_TARGET_TAU_INDEX] - TAU_BOOTSTRAP).abs() < EPS, + "ISV[τ={RL_TARGET_TAU_INDEX}] expected {TAU_BOOTSTRAP}, got {}", + isv[RL_TARGET_TAU_INDEX] + ); + assert!( + (isv[RL_PPO_CLIP_INDEX] - EPS_BOOTSTRAP).abs() < EPS, + "ISV[ε={RL_PPO_CLIP_INDEX}] expected {EPS_BOOTSTRAP}, got {}", + isv[RL_PPO_CLIP_INDEX] + ); + assert!( + (isv[RL_ENTROPY_COEF_INDEX] - COEF_BOOTSTRAP).abs() < EPS, + "ISV[entropy_coef={RL_ENTROPY_COEF_INDEX}] expected {COEF_BOOTSTRAP}, got {}", + isv[RL_ENTROPY_COEF_INDEX] + ); + assert!( + (isv[RL_N_ROLLOUT_STEPS_INDEX] - ROLLOUT_BOOTSTRAP).abs() < EPS, + "ISV[n_rollout_steps={RL_N_ROLLOUT_STEPS_INDEX}] expected {ROLLOUT_BOOTSTRAP}, got {}", + isv[RL_N_ROLLOUT_STEPS_INDEX] + ); + assert!( + (isv[RL_PER_ALPHA_INDEX] - PER_ALPHA_BOOTSTRAP).abs() < EPS, + "ISV[per_α={RL_PER_ALPHA_INDEX}] expected {PER_ALPHA_BOOTSTRAP}, got {}", + isv[RL_PER_ALPHA_INDEX] + ); + assert!( + (isv[RL_REWARD_SCALE_INDEX] - REWARD_SCALE_BOOTSTRAP).abs() < EPS, + "ISV[reward_scale={RL_REWARD_SCALE_INDEX}] expected {REWARD_SCALE_BOOTSTRAP}, got {}", + isv[RL_REWARD_SCALE_INDEX] + ); + + // Invariant: the 7 EMA-input slots (417..424) are NOT yet + // bootstrapped. Phase R3 wires the EMA producer kernels that fill + // these; until then they MUST remain at `alloc_zeros` sentinel 0. + // Asserting this here protects against accidental controller-side + // writes to the wrong slot (a defect a one-character bug could + // introduce, given the 7 slots are sequential). + for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END { + assert_eq!( + isv[slot], 0.0, + "ISV[{slot}] expected sentinel 0.0 (R3 wires EMA producers), got {}", + isv[slot] + ); + } + + eprintln!( + "G1 OK — bootstraps: γ={:.4} τ={:.4} ε={:.4} entropy_coef={:.4} \ + n_rollout_steps={:.1} per_α={:.4} reward_scale={:.4}", + isv[RL_GAMMA_INDEX], + isv[RL_TARGET_TAU_INDEX], + isv[RL_PPO_CLIP_INDEX], + isv[RL_ENTROPY_COEF_INDEX], + isv[RL_N_ROLLOUT_STEPS_INDEX], + isv[RL_PER_ALPHA_INDEX], + isv[RL_REWARD_SCALE_INDEX], + ); +}