diff --git a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu index 32cc8e812..55ad7b004 100644 --- a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu +++ b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu @@ -39,8 +39,10 @@ #define RL_Q_DISTILL_LAMBDA_INDEX 486 #define RL_Q_DISTILL_TEMPERATURE_INDEX 487 #define RL_Q_DISTILL_KL_EMA_INDEX 488 - -#define KL_EMA_ALPHA 0.05f +// audit-isv 2026-05-24: KL_EMA_ALPHA was hardcoded 0.05f — caught +// by scripts/audit-isv.sh on the first manifest-driven dogfood +// run. ISV-resident now per SP20 §0.1. +#define RL_Q_DISTILL_KL_EMA_ALPHA_INDEX 493 extern "C" __global__ void rl_q_pi_distill_grad( const float* __restrict__ q_logits, // [B × N_ACTIONS × Q_N_ATOMS] @@ -167,8 +169,9 @@ extern "C" __global__ void rl_q_pi_distill_grad( if (b == 0 && a == 0) { const float kl_prev = isv[RL_Q_DISTILL_KL_EMA_INDEX]; + const float alpha = isv[RL_Q_DISTILL_KL_EMA_ALPHA_INDEX]; const float kl_new = (kl_prev == 0.0f) ? s_kl - : (1.0f - KL_EMA_ALPHA) * kl_prev + KL_EMA_ALPHA * s_kl; + : (1.0f - alpha) * kl_prev + alpha * s_kl; isv[RL_Q_DISTILL_KL_EMA_INDEX] = kl_new; } } diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 56e7520bf..00a6942d9 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -45,7 +45,9 @@ //! | 489-490 | 2 negative-tail tracking slots (raw + EMA) for adaptive RATIO | apply_reward_scale + rl_reward_clamp_controller | //! | 491-492 | 2 controller-anchor slots (KL distill target + reward_scale MIN) | rl_isv_write (seed) + adaptive controllers | //! -//! Total: 93 slots; `RL_SLOTS_END = 493`. +//! | 493 | Q→π distill KL_EMA blend alpha | rl_isv_write (seed) + rl_q_pi_distill_grad | +//! +//! Total: 94 slots; `RL_SLOTS_END = 494`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -724,6 +726,13 @@ pub const RL_Q_DISTILL_KL_TARGET_INDEX: usize = 491; /// Seeded by `rl_isv_write`. pub const RL_REWARD_SCALE_MIN_INDEX: usize = 492; +/// Q→π distill KL_EMA blend α — was hardcoded `0.05f` in +/// `rl_q_pi_distill_grad.cu`. Caught by `scripts/audit-isv.sh` on +/// the first manifest-driven dogfood run pre-SP20. Per SP20 §0.1 +/// every numerical constant in new (and existing, when discovered) +/// kernels is ISV-resident. Seeded 0.05 (preserves prior behavior). +pub const RL_Q_DISTILL_KL_EMA_ALPHA_INDEX: usize = 493; + /// 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 = 493; +pub const RL_SLOTS_END: usize = 494; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index d7a5e9935..28be1c903 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1251,7 +1251,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 35] = [ + let isv_constants: [(usize, f32); 36] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1283,6 +1283,9 @@ impl IntegratedTrainer { // one more order of magnitude before pegging when // mean_abs_pnl spikes. (crate::rl::isv_slots::RL_REWARD_SCALE_MIN_INDEX, 1e-4), + // Q→π distill KL_EMA blend α (was hardcoded 0.05f in + // rl_q_pi_distill_grad.cu, caught by audit-isv.sh dogfood). + (crate::rl::isv_slots::RL_Q_DISTILL_KL_EMA_ALPHA_INDEX, 0.05), (crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01), (crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99), (crate::rl::isv_slots::RL_PLATEAU_PATIENCE_INDEX, 1000.0),