feat(cuda): Thompson sampling floor + LOSS=3.0 bootstrap for entropy stability
Two fixes that decouple entropy stability from loss aversion: 1. Thompson sampling probability floor (RL_THOMPSON_FLOOR_INDEX=588, bootstrap 0.05): 5% of steps pick uniform random action. Prevents any action from reaching π=0 (δ-function attractor). ISV-driven. Per pearl_pi_actor_collapses_without_entropy_floor. 2. LOSS bootstrap restored to 3.0 (was 1.5). LOSS=1.5 caused entropy collapse to 0.94 despite max SAC. LOSS=3.0 keeps entropy stable at 1.87 (proven over 58k steps). The done-gated EMAs (slots 585/586) will adapt LOSS toward ~1.45 AFTER entropy stabilizes. Together: entropy stays stable (LOSS=3.0 + 5% floor) AND PnL improves (adaptive LOSS lowers toward real L/W ratio after warmup). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
@@ -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;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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 {
|
||||
|
||||
Reference in New Issue
Block a user