diff --git a/crates/ml-alpha/cuda/rl_action_kernel.cu b/crates/ml-alpha/cuda/rl_action_kernel.cu index b4da974a5..ee8517af6 100644 --- a/crates/ml-alpha/cuda/rl_action_kernel.cu +++ b/crates/ml-alpha/cuda/rl_action_kernel.cu @@ -35,6 +35,7 @@ #define N_ACTIONS 11 #define Q_N_ATOMS 21 +#define RL_THOMPSON_FLOOR_INDEX 588 __device__ static uint32_t xorshift32(uint32_t* state) { uint32_t x = *state; @@ -54,6 +55,7 @@ __device__ static uint32_t xorshift32(uint32_t* state) { // q_logits [b_size, N_ACTIONS, Q_N_ATOMS] row-major // atom_supports [Q_N_ATOMS] e.g. linspace(Q_V_MIN, Q_V_MAX, Q_N_ATOMS) // prng_state [b_size] per-batch xorshift32 state (mutated in place) +// isv [>=RL_THOMPSON_FLOOR_INDEX+1] ISV bus (read-only) // Outputs: // actions [b_size] Thompson-sampled action index per batch extern "C" __global__ void rl_action_kernel( @@ -61,6 +63,7 @@ extern "C" __global__ void rl_action_kernel( const float* __restrict__ atom_supports, uint32_t* __restrict__ prng_state, int* __restrict__ actions, + const float* __restrict__ isv, int b_size ) { const int b = blockIdx.x; @@ -112,7 +115,8 @@ extern "C" __global__ void rl_action_kernel( __syncthreads(); // Thread 0: argmax over per-action sampled returns; write action; - // advance per-batch PRNG state for next call. + // apply epsilon-greedy floor; advance per-batch PRNG state for + // next call. if (a == 0) { int best_a = 0; float best_v = sampled_returns[0]; @@ -125,8 +129,18 @@ extern "C" __global__ void rl_action_kernel( } actions[b] = best_a; + // Epsilon-greedy floor: with probability `floor` (ISV-driven), + // override with a uniform random action. Prevents entropy + // collapse to a delta-function attractor when Q sharpens — + // per `pearl_pi_actor_collapses_without_entropy_floor`. + const float floor = __ldg(isv + RL_THOMPSON_FLOOR_INDEX); uint32_t adv = base; xorshift32(&adv); + const float uniform = (float)(adv & 0xFFFFFFu) / (float)0xFFFFFFu; + if (uniform < floor) { + xorshift32(&adv); + actions[b] = (int)(adv % N_ACTIONS); + } prng_state[b] = adv; } } diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 052cd1609..735fa4f78 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -53,7 +53,9 @@ //! | 547 | Noisy linear σ_init (NoisyNet exploration) | NoisyLinear layers | //! | 548 | Device-resident step counter (graph-safe) | rl_increment_step | //! -//! Total: 149 slots; `RL_SLOTS_END = 549`. +//! | 588 | Thompson sampler epsilon-greedy floor | rl_action_kernel (read-only) | +//! +//! Total: 150 slots; `RL_SLOTS_END = 589`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -1134,5 +1136,15 @@ pub const RL_DONE_LOSS_MAGNITUDE_EMA_INDEX: usize = 586; /// are overridden to Hold. pub const RL_MARGIN_FLOOR_USD_INDEX: usize = 587; +/// Thompson sampler epsilon-greedy floor. With probability `floor`, +/// the action kernel picks a UNIFORM random action instead of the +/// Thompson-sampled argmax. Prevents entropy collapse to a +/// delta-function attractor when the Q distribution sharpens — +/// per `pearl_pi_actor_collapses_without_entropy_floor`, SAC entropy +/// gradient vanishes at π→0. At floor=0.05 each action gets ≥0.45% +/// minimum sampling rate (floor/N_ACTIONS). Read by `rl_action_kernel` +/// from `isv[RL_THOMPSON_FLOOR_INDEX]`. Bootstrap 0.05. +pub const RL_THOMPSON_FLOOR_INDEX: usize = 588; + /// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 588; +pub const RL_SLOTS_END: usize = 589; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 2dede7d93..daca128f3 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -2673,7 +2673,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 112] = [ + let isv_constants: [(usize, f32); 113] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -2687,7 +2687,7 @@ impl IntegratedTrainer { // [-1, +1] span as the canonical floor — the ratchet // in `rl_reward_clamp_controller` only grows magnitude. (crate::rl::isv_slots::RL_REWARD_CLAMP_WIN_INDEX, 1.0), - (crate::rl::isv_slots::RL_REWARD_CLAMP_LOSS_INDEX, 1.5), + (crate::rl::isv_slots::RL_REWARD_CLAMP_LOSS_INDEX, 3.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_MARGIN_INDEX, 1.5), (crate::rl::isv_slots::RL_REWARD_CLAMP_RATIO_INDEX, 3.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX, 0.05), @@ -2843,6 +2843,10 @@ impl IntegratedTrainer { // Hard margin floor: $35k capital, ~$14.4k ES maintenance. // Max drawdown ~$20k before broker liquidates. (crate::rl::isv_slots::RL_MARGIN_FLOOR_USD_INDEX, 20000.0), + // Thompson sampler epsilon-greedy floor — 5% of steps + // pick a uniform random action, preventing entropy collapse + // per `pearl_pi_actor_collapses_without_entropy_floor`. + (crate::rl::isv_slots::RL_THOMPSON_FLOOR_INDEX, 0.05), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -3155,6 +3159,7 @@ impl IntegratedTrainer { args.push_ptr(self.atom_supports_d.raw_ptr()); args.push_ptr(self.prng_state_d.raw_ptr()); args.push_ptr(actions_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); args.push_i32(b_size_i); let mut ptrs = args.build_arg_ptrs(); unsafe {