gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat: per-branch epsilon — magnitude 2× exploration, Boltzmann for order/urgency

Browse Source 
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-11 12:39:02 +02:00
parent d6ab83d37c
commit 59695d67b7
8 changed files with 152 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files

72
crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu
View File

@@ -194,7 +194,9 @@ extern "C" __global__ void branching_action_select(
const float* __restrict__ q_urgency, /* [batch_size, 3] */
unsigned int* rng_states, /* [batch_size] */
unsigned int* actions_out, /* [batch_size] -- factored index 0-44 */
const float epsilon,
const float epsilon_exp, /* per-branch epsilon: exposure (dir+mag) */
const float epsilon_ord, /* per-branch epsilon: order type */
const float epsilon_urg, /* per-branch epsilon: urgency */
const int batch_size,
const float q_gap_threshold, /* min Q-gap for trade entry (0.0 = disabled) */
const float* __restrict__ bonus_exposure, /* [5] UCB count bonus per exposure action (NULL = disabled) */
@@ -205,12 +207,14 @@ extern "C" __global__ void branching_action_select(
if (idx >= batch_size) return;
unsigned int rng = rng_states[idx];
float eps_bf = bf16(epsilon);
float eps_exp_bf = bf16(epsilon_exp);
float eps_ord_bf = bf16(epsilon_ord);
float eps_urg_bf = bf16(epsilon_urg);
/* Head 1: Exposure (5 actions) -- with UCB bonus + Q-gap conviction filter */
int exposure;
float r1 = bf16(gpu_random(&rng));
if (r1 < eps_bf) {
if (r1 < eps_exp_bf) {
exposure = (int)(gpu_random(&rng) * 5.0f);
if (exposure >= 5) exposure = 4;
} else {
@@ -237,35 +241,73 @@ extern "C" __global__ void branching_action_select(
}
}
/* Head 2: Order type (3 actions) -- with UCB bonus */
/* Head 2: Order type (3 actions) -- Boltzmann softmax with UCB bonus */
int order;
float r2 = bf16(gpu_random(&rng));
if (r2 < eps_bf) {
if (r2 < eps_ord_bf) {
/* Random exploration */
order = (int)(gpu_random(&rng) * 3.0f);
if (order >= 3) order = 2;
} else {
/* Boltzmann softmax over order Q-values */
const float* qo = q_order + idx * 3;
float best = qo[0] + ((bonus_order != NULL) ? bonus_order[0] : bf16_zero());
order = 0;
float q_max_o = qo[0] + ((bonus_order != NULL) ? bonus_order[0] : bf16_zero());
float q_min_o = q_max_o;
float qo_adj[3];
qo_adj[0] = q_max_o;
for (int a = 1; a < 3; a++) {
float q_plus_bonus = qo[a] + ((bonus_order != NULL) ? bonus_order[a] : bf16_zero());
if (q_plus_bonus > best) { best = q_plus_bonus; order = a; }
qo_adj[a] = qo[a] + ((bonus_order != NULL) ? bonus_order[a] : bf16_zero());
q_max_o = fmaxf(q_max_o, qo_adj[a]);
q_min_o = fminf(q_min_o, qo_adj[a]);
}
float tau_ord = fmaxf(q_max_o - q_min_o, 0.01f);
float sum_e = 0.0f;
float exps_ord[3];
for (int a = 0; a < 3; a++) {
exps_ord[a] = expf((qo_adj[a] - q_max_o) / tau_ord);
sum_e += exps_ord[a];
}
float ro = gpu_random(&rng) * sum_e;
float cum = 0.0f;
order = 2;
for (int a = 0; a < 3; a++) {
cum += exps_ord[a];
if (ro < cum) { order = a; break; }
}
}
/* Head 3: Urgency (3 actions) -- with UCB bonus */
/* Head 3: Urgency (3 actions) -- Boltzmann softmax with UCB bonus */
int urgency;
float r3 = bf16(gpu_random(&rng));
if (r3 < eps_bf) {
if (r3 < eps_urg_bf) {
/* Random exploration */
urgency = (int)(gpu_random(&rng) * 3.0f);
if (urgency >= 3) urgency = 2;
} else {
/* Boltzmann softmax over urgency Q-values */
const float* qu = q_urgency + idx * 3;
float best = qu[0] + ((bonus_urgency != NULL) ? bonus_urgency[0] : bf16_zero());
urgency = 0;
float q_max_u = qu[0] + ((bonus_urgency != NULL) ? bonus_urgency[0] : bf16_zero());
float q_min_u = q_max_u;
float qu_adj[3];
qu_adj[0] = q_max_u;
for (int a = 1; a < 3; a++) {
float q_plus_bonus = qu[a] + ((bonus_urgency != NULL) ? bonus_urgency[a] : bf16_zero());
if (q_plus_bonus > best) { best = q_plus_bonus; urgency = a; }
qu_adj[a] = qu[a] + ((bonus_urgency != NULL) ? bonus_urgency[a] : bf16_zero());
q_max_u = fmaxf(q_max_u, qu_adj[a]);
q_min_u = fminf(q_min_u, qu_adj[a]);
}
float tau_urg = fmaxf(q_max_u - q_min_u, 0.01f);
float sum_e = 0.0f;
float exps_urg[3];
for (int a = 0; a < 3; a++) {
exps_urg[a] = expf((qu_adj[a] - q_max_u) / tau_urg);
sum_e += exps_urg[a];
}
float ru = gpu_random(&rng) * sum_e;
float cum = 0.0f;
urgency = 2;
for (int a = 0; a < 3; a++) {
cum += exps_urg[a];
if (ru < cum) { urgency = a; break; }
}
}

70
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -731,7 +731,10 @@ extern "C" __global__ void experience_action_select(
float q_gap_threshold
,const float* __restrict__ portfolio_states, /* [N, 23] read-only */
int min_hold_bars,
float max_position
float max_position,
float eps_exp_mult, /* per-branch epsilon multiplier: exposure (dir+mag) */
float eps_ord_mult, /* per-branch epsilon multiplier: order type */
float eps_urg_mult /* per-branch epsilon multiplier: urgency */
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -743,12 +746,13 @@ extern "C" __global__ void experience_action_select(
float epsilon = eps_end + 0.5f * (eps_start - eps_end)
* (1.0f + cosf(3.14159265f * fminf(cosine_progress, 1.0f)));
/* Gem 2: Per-branch epsilon — direction AND magnitude explore fully.
* Both are hard problems: direction chooses side, magnitude chooses size.
* Order/urgency get reduced epsilon: 10% of epsilon range. */
float eps_dir = epsilon;
float eps_mag = epsilon; /* magnitude needs full exploration — same as direction */
float eps_other = eps_end + 0.1f * (epsilon - eps_end);
/* Per-branch epsilon via config multipliers (ABI-2).
* Direction and magnitude share eps_exp_mult (magnitude only active when directional).
* Order and urgency have independent multipliers for Boltzmann exploration. */
float eps_dir = fminf(epsilon * eps_exp_mult, 1.0f);
float eps_mag = fminf(epsilon * eps_exp_mult, 1.0f);
float eps_ord = fminf(epsilon * eps_ord_mult, 1.0f);
float eps_urg = fminf(epsilon * eps_urg_mult, 1.0f);
int q_stride = b0_size + b1_size + b2_size + b3_size;
const float* q_b0 = q_values + (long long)i * q_stride;
@@ -901,20 +905,60 @@ extern "C" __global__ void experience_action_select(
}
} /* end else (not in_hold) */
/* Branch 2: order type — Gem 2: uses eps_other */
if (lcg_random(&rng) < eps_other) {
/* Branch 2: order type — Boltzmann softmax (ABI-2) */
if (lcg_random(&rng) < eps_ord) {
/* Random exploration */
int r = (int)(lcg_random(&rng) * (float)b2_size);
a2 = (r >= b2_size) ? b2_size - 1 : r;
} else {
a2 = argmax_n(q_b2, b2_size);
/* Boltzmann softmax over order Q-values */
float q_max_ord = q_b2[0], q_min_ord = q_b2[0];
for (int a = 1; a < b2_size; a++) {
q_max_ord = fmaxf(q_max_ord, q_b2[a]);
q_min_ord = fminf(q_min_ord, q_b2[a]);
}
float tau_ord = fmaxf(q_max_ord - q_min_ord, 0.01f);
float sum_e = 0.0f;
float exps_ord[3];
for (int a = 0; a < b2_size; a++) {
exps_ord[a] = expf((q_b2[a] - q_max_ord) / tau_ord);
sum_e += exps_ord[a];
}
float ro = lcg_random(&rng) * sum_e;
float cum = 0.0f;
a2 = b2_size - 1;
for (int a = 0; a < b2_size; a++) {
cum += exps_ord[a];
if (ro < cum) { a2 = a; break; }
}
}
/* Branch 3: urgency — Gem 2: uses eps_other */
if (lcg_random(&rng) < eps_other) {
/* Branch 3: urgency — Boltzmann softmax (ABI-2) */
if (lcg_random(&rng) < eps_urg) {
/* Random exploration */
int r = (int)(lcg_random(&rng) * (float)b3_size);
a3 = (r >= b3_size) ? b3_size - 1 : r;
} else {
a3 = argmax_n(q_b3, b3_size);
/* Boltzmann softmax over urgency Q-values */
float q_max_urg = q_b3[0], q_min_urg = q_b3[0];
for (int a = 1; a < b3_size; a++) {
q_max_urg = fmaxf(q_max_urg, q_b3[a]);
q_min_urg = fminf(q_min_urg, q_b3[a]);
}
float tau_urg = fmaxf(q_max_urg - q_min_urg, 0.01f);
float sum_e = 0.0f;
float exps_urg[3];
for (int a = 0; a < b3_size; a++) {
exps_urg[a] = expf((q_b3[a] - q_max_urg) / tau_urg);
sum_e += exps_urg[a];
}
float ru = lcg_random(&rng) * sum_e;
float cum = 0.0f;
a3 = b3_size - 1;
for (int a = 0; a < b3_size; a++) {
cum += exps_urg[a];
if (ru < cum) { a3 = a; break; }
}
}
/* Compose factored action: 4-branch encoding

6
crates/ml/src/cuda_pipeline/gpu_action_selector.rs
View File

@@ -140,7 +140,7 @@ impl GpuActionSelector {
}
pub fn select_actions_branching(&mut self, q_exposure: &CudaSlice<f32>, q_order: &CudaSlice<f32>,
q_urgency: &CudaSlice<f32>, epsilon: f32, batch_size: usize,
q_urgency: &CudaSlice<f32>, epsilon_exp: f32, epsilon_ord: f32, epsilon_urg: f32, batch_size: usize,
) -> Result<CudaSlice<u32>, MLError> {
if batch_size == 0 { return self.stream.alloc_zeros::<u32>(0).map_err(|e| MLError::ModelError(format!("alloc empty: {e}"))); }
if batch_size > self.max_batch_size { return Err(MLError::ModelError(format!("batch_size {} exceeds max {}", batch_size, self.max_batch_size))); }
@@ -152,7 +152,9 @@ impl GpuActionSelector {
unsafe {
self.stream.launch_builder(&self.branching_kernel_func)
.arg(q_exposure).arg(q_order).arg(q_urgency)
.arg(&mut self.rng_states).arg(&mut self.actions_buf).arg(&epsilon).arg(&bs_i32).arg(&self.q_gap_threshold)
.arg(&mut self.rng_states).arg(&mut self.actions_buf)
.arg(&epsilon_exp).arg(&epsilon_ord).arg(&epsilon_urg)
.arg(&bs_i32).arg(&self.q_gap_threshold)
.arg(&be_ptr).arg(&bo_ptr).arg(&bu_ptr) // UCB count bonuses (0 = NULL = disabled)
.launch(config).map_err(|e| MLError::ModelError(format!("branching kernel launch: {e}")))?;
}

3
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -1167,6 +1167,9 @@ impl GpuBacktestEvaluator {
.arg(&null_portfolio)
.arg(&eval_min_hold)
.arg(&eval_max_pos)
.arg(&1.0_f32) // eps_exp_mult (no-op at eps=0)
.arg(&1.0_f32) // eps_ord_mult (no-op at eps=0)
.arg(&1.0_f32) // eps_urg_mult (no-op at eps=0)
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"chunked action_select chunk_start={chunk_start}: {e}"

12
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -297,6 +297,12 @@ pub struct ExperienceCollectorConfig {
pub eps_start: f32,
/// v8: Epsilon schedule end value.
pub eps_end: f32,
/// Per-branch epsilon multiplier: exposure (dir+mag). Default 1.0.
pub eps_exp_mult: f32,
/// Per-branch epsilon multiplier: order type. Default 1.5.
pub eps_ord_mult: f32,
/// Per-branch epsilon multiplier: urgency. Default 1.5.
pub eps_urg_mult: f32,
/// v8: TD(λ) trace parameter. 0.0 = 1-step TD (n-step only), >0.01 = TD(λ) overwrites n-step.
pub td_lambda: f32,
/// v8: Maximum trace length for truncated TD(λ).
@@ -375,6 +381,9 @@ impl Default for ExperienceCollectorConfig {
total_epochs: 20,
eps_start: 0.3,
eps_end: 0.02,
eps_exp_mult: 1.0,
eps_ord_mult: 1.5,
eps_urg_mult: 1.5,
td_lambda: 0.0,
max_trace_length: 7,
hindsight_fraction: 0.0,
@@ -1879,6 +1888,9 @@ impl GpuExperienceCollector {
.arg(&self.portfolio_states) // portfolio state for hold masking
.arg(&min_hold_bars_i32) // min_hold_bars
.arg(&max_pos) // max_position for direction index
.arg(&config.eps_exp_mult) // per-branch epsilon multiplier: exposure
.arg(&config.eps_ord_mult) // per-branch epsilon multiplier: order
.arg(&config.eps_urg_mult) // per-branch epsilon multiplier: urgency
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_action_select t={t}: {e}"

12
crates/ml/src/trainers/dqn/config.rs
View File

@@ -425,6 +425,14 @@ pub struct DQNHyperparameters {
pub epsilon_end: f64,
/// Exploration decay rate
pub epsilon_decay: f64,
/// Per-branch epsilon multiplier: direction (relative to base epsilon). Default 1.0.
pub epsilon_direction_mult: f64,
/// Per-branch epsilon multiplier: magnitude (relative to base epsilon). Default 2.0 — needs more exploration.
pub epsilon_magnitude_mult: f64,
/// Per-branch epsilon multiplier: order type (relative to base epsilon). Default 1.5.
pub epsilon_order_mult: f64,
/// Per-branch epsilon multiplier: urgency (relative to base epsilon). Default 1.5.
pub epsilon_urgency_mult: f64,
/// Replay buffer capacity
pub buffer_size: usize,
/// Fraction of free VRAM to allocate for replay buffer (0.0 = use static buffer_size).
@@ -1265,6 +1273,10 @@ impl DQNHyperparameters {
epsilon_start: 0.3,
epsilon_end: 0.02,
epsilon_decay: 0.995,
epsilon_direction_mult: 1.0,
epsilon_magnitude_mult: 2.0,
epsilon_order_mult: 1.5,
epsilon_urgency_mult: 1.5,
buffer_size: 500000, // WAVE 24: Increased from 100K for better sample diversity (anti-overfitting)
replay_buffer_vram_fraction: 0.70,
min_replay_size: 1000,

5
crates/ml/src/trainers/dqn/trainer/action.rs
View File

@@ -99,7 +99,10 @@ impl DQNTrainer {
}
let factored_slice = if let Some((ref q_exp, ref q_ord, ref q_urg)) = branching_q_slices {
selector.select_actions_branching(q_exp, q_ord, q_urg, epsilon, batch_size).map_err(|e| anyhow::anyhow!("GPU branching action selection failed: {e}"))?
let eps_exp = epsilon * self.hyperparams.epsilon_magnitude_mult as f32;
let eps_ord = epsilon * self.hyperparams.epsilon_order_mult as f32;
let eps_urg = epsilon * self.hyperparams.epsilon_urgency_mult as f32;
selector.select_actions_branching(q_exp, q_ord, q_urg, eps_exp, eps_ord, eps_urg, batch_size).map_err(|e| anyhow::anyhow!("GPU branching action selection failed: {e}"))?
} else {
let exposure_slice = selector.select_actions(&batch_q_values, epsilon, batch_size, 5).map_err(|e| anyhow::anyhow!("GPU fused action selection failed: {e}"))?;
selector.route_exposure_to_factored(&exposure_slice, batch_size, self.hyperparams.avg_spread as f32, self.hyperparams.avg_spread as f32, self.vol_ema as f32, self.median_vol as f32)

3
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -1098,6 +1098,9 @@ impl DQNTrainer {
total_epochs: self.hyperparams.epochs as i32,
eps_start: self.hyperparams.epsilon_start as f32,
eps_end: self.hyperparams.epsilon_end as f32,
eps_exp_mult: self.hyperparams.epsilon_magnitude_mult as f32,
eps_ord_mult: self.hyperparams.epsilon_order_mult as f32,
eps_urg_mult: self.hyperparams.epsilon_urgency_mult as f32,
td_lambda: self.hyperparams.td_lambda as f32,
max_trace_length: self.hyperparams.max_trace_length as i32,
hindsight_fraction: self.hyperparams.hindsight_fraction as f32,