feat(rl): R1 — ISV slot extension + 7-controller bootstrap (G1 gate)
Closes defect #1 from the flawed Phase F+G arc: ISV[400..406] were left at alloc_zeros sentinel 0 in production, causing bellman_target_projection (γ=0), ppo_clipped_surrogate (ε=0, entropy=0), and the C51 backward to train against degenerate targets that the MockLobEnv toy fixture (done=true every step, horizon=1) intrinsically could not detect. Three changes: 1. Port crates/ml-alpha/cuda/rl_reward_scale_controller.cu from the ml-alpha-phase-f-g-flawed reference branch (93 lines, unchanged). Add to build.rs KERNELS list; bump cache-bust to v25. 2. Extend src/rl/isv_slots.rs: add 7 new EMA-input slot constants (RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_MEAN_ABS_PNL_EMA_INDEX), RL_SLOTS_END goes 417 -> 424. These are reserved for the EMA producer kernels Phase R3 lands; in R1 they stay at sentinel 0 (asserted by the G1 test). 3. Wire all 7 RL adaptive controllers (γ / τ / ε / entropy_coef / n_rollout_steps / per_α / reward_scale) into IntegratedTrainer: - 7 cubin includes + 7 module/function fields - All 7 loaded in new() via the existing load_cubin pattern - New fn launch_isv_controller_3arg() centralises the shared (isv*, alpha, scalar_input) launch signature - New fn with_controllers_bootstrapped() consumes self and fires each controller once against the freshly-zeroed isv_d; each kernel's first-observation-bootstrap path (per pearl_first_observation_bootstrap) sees sentinel zero in its slot and writes its canonical *_BOOTSTRAP value: ISV[400] γ = 0.99 ISV[401] τ = 0.005 ISV[402] ε = 0.2 ISV[403] entropy_coef = 0.01 ISV[404] n_rollout_steps= 2048 ISV[405] per_α = 0.6 ISV[406] reward_scale = 1.0 - new() ends with `.with_controllers_bootstrapped()?` so every trainer construction site picks this up automatically. This replaces the flawed Phase F approach of host memcpy_htod-ing canonical constants into ISV, which violated feedback_no_htod_htoh_only_mapped_pinned (tests not exempt) AND short-circuited the canonical pearl_first_observation_bootstrap pattern every other adaptive controller in the codebase uses. The launch_isv_controller_3arg helper is reused by Phase R5's per-step controller launches with real EMA inputs sourced from ISV[417..424] — at that point the Wiener-α blend kicks in and the slots adapt away from the R1 bootstrap defaults. Gate G1 (crates/ml-alpha/tests/isv_bootstrap.rs): - Construct IntegratedTrainer - memcpy_dtoh full ISV slice to host - Assert ISV[400..406] equal each kernel's #define *_BOOTSTRAP - Assert ISV[417..424] still at sentinel 0 (R3 wires producers) Per feedback_no_cpu_test_fallbacks: the oracle is the kernel's own *_BOOTSTRAP constant, not a CPU computation. Per pearl_tests_must_prove_not_lock_observations: the test asserts an invariant (bootstrap path wrote the canonical value defined by the kernel), not a tuned magic number. Build clean: cargo check + cargo build --test isv_bootstrap on ml-alpha both green. CUDA-required, #[ignore]'d for non-GPU CI. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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");
|
||||
|
||||
93
crates/ml-alpha/cuda/rl_reward_scale_controller.cu
Normal file
93
crates/ml-alpha/cuda/rl_reward_scale_controller.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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<f32>,
|
||||
|
||||
/// 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<CudaModule>,
|
||||
rl_gamma_controller_fn: CudaFunction,
|
||||
_rl_target_tau_controller_module: Arc<CudaModule>,
|
||||
rl_target_tau_controller_fn: CudaFunction,
|
||||
_rl_ppo_clip_controller_module: Arc<CudaModule>,
|
||||
rl_ppo_clip_controller_fn: CudaFunction,
|
||||
_rl_entropy_coef_controller_module: Arc<CudaModule>,
|
||||
rl_entropy_coef_controller_fn: CudaFunction,
|
||||
_rl_rollout_steps_controller_module: Arc<CudaModule>,
|
||||
rl_rollout_steps_controller_fn: CudaFunction,
|
||||
_rl_per_alpha_controller_module: Arc<CudaModule>,
|
||||
rl_per_alpha_controller_fn: CudaFunction,
|
||||
_rl_reward_scale_controller_module: Arc<CudaModule>,
|
||||
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<Self> {
|
||||
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
|
||||
|
||||
143
crates/ml-alpha/tests/isv_bootstrap.rs
Normal file
143
crates/ml-alpha/tests/isv_bootstrap.rs
Normal file
@@ -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],
|
||||
);
|
||||
}
|
||||
Reference in New Issue
Block a user