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

feat(D7/N7 Part B): wire contrarian Q-negation into Boltzmann action selection

Browse Source 
Add `contrarian_active` int parameter to `experience_action_select` kernel.
When non-zero, a `q_sign = -1.0f` multiplier is applied to Q values across
all 4 branches (direction/magnitude/order/urgency) before Boltzmann softmax,
converting argmax-favoring sampling to argmin-favoring without touching
temperature, epsilon, conviction filter, masking, or sampling logic.
When zero, q_sign = +1.0f — behavior is bit-identical to before.

Wire: GpuExperienceCollector gains `contrarian_active_cache: u8` field plus
`set_contrarian_active()` / `contrarian_active()` accessors. The flag is
appended as the last arg at the action_select launch site. training_loop.rs
propagates `self.contrarian_active` to the collector immediately after the
D7 Part A state machine updates.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-20 20:51:11 +02:00
parent ac58b9d402
commit 598c1d57f7
3 changed files with 52 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -761,7 +761,8 @@ extern "C" __global__ void experience_action_select(
float eps_urg_mult, /* per-branch epsilon multiplier: urgency */
int timestep, /* current timestep for stateless RNG */
const float* __restrict__ per_sample_epsilon, /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */
const float* __restrict__ isv_signals_ptr /* [8] pinned ISV signals for adaptive hold. NULL = static. */
const float* __restrict__ isv_signals_ptr, /* [8] pinned ISV signals for adaptive hold. NULL = static. */
int contrarian_active /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -804,6 +805,14 @@ extern "C" __global__ void experience_action_select(
const float* q_b2 = q_b1 + b1_size;
const float* q_b3 = q_b2 + b2_size;
/* D7/N7: contrarian sign flip.
* When contrarian_active != 0, negate Q values before Boltzmann softmax.
* softmax(-q/tau) is argmin-favoring — the model picks its least-preferred
* action, breaking systematic anti-correlation attractors.
* When contrarian_active == 0, q_sign = +1.0f → bit-identical to previous
* behavior (identity multiplication). */
float q_sign = (contrarian_active != 0) ? -1.0f : 1.0f;
int dir_idx, mag_idx, a2, a3;
/* Hold enforcement removed — replaced by cost-driven hold timing.
@@ -833,10 +842,10 @@ extern "C" __global__ void experience_action_select(
* Q-values differentiate (converges to argmax deterministically).
* When Q-values are flat, spreads evenly instead of flipping on noise. */
/* Training mode: Boltzmann softmax over direction Q-values */
float q_max_d = (q_b0[0]);
float q_max_d = q_sign * (q_b0[0]);
float q_min_d = q_max_d;
for (int a = 1; a < b0_size; a++) {
float qv = (q_b0[a]);
float qv = q_sign * (q_b0[a]);
q_max_d = fmaxf(q_max_d, qv);
q_min_d = fminf(q_min_d, qv);
}
@@ -845,7 +854,7 @@ extern "C" __global__ void experience_action_select(
float sum_e = 0.0f;
float exps_d[4]; /* b0_size=4: Short/Hold/Long/Flat */
for (int a = 0; a < b0_size; a++) {
float qv = (q_b0[a]);
float qv = q_sign * (q_b0[a]);
exps_d[a] = expf((qv - q_max_d) / tau_d);
sum_e += exps_d[a];
}
@@ -889,10 +898,10 @@ extern "C" __global__ void experience_action_select(
} else {
/* Adaptive temperature: scale by Q-range so Boltzmann is meaningful
* regardless of absolute Q magnitude. Floor at 0.01 to avoid div/0. */
float q_max_m = (q_b1[0]);
float q_max_m = q_sign * (q_b1[0]);
float q_min_m = q_max_m;
for (int a = 1; a < b1_size; a++) {
float qv = (q_b1[a]);
float qv = q_sign * (q_b1[a]);
q_max_m = fmaxf(q_max_m, qv);
q_min_m = fminf(q_min_m, qv);
}
@@ -906,7 +915,7 @@ extern "C" __global__ void experience_action_select(
float sum_e = 0.0f;
float exps[MAX_MAGNITUDE_ACTIONS];
for (int a = 0; a < b1_size && a < MAX_MAGNITUDE_ACTIONS; a++) {
float qv = (q_b1[a]);
float qv = q_sign * (q_b1[a]);
exps[a] = expf((qv - q_max_m) / tau);
sum_e += exps[a];
}
@@ -925,10 +934,11 @@ extern "C" __global__ void experience_action_select(
a2 = (r >= b2_size) ? b2_size - 1 : r;
} else {
/* Boltzmann softmax over order Q-values */
float q_max_ord = q_b2[0], q_min_ord = q_b2[0];
float q_max_ord = q_sign * q_b2[0], q_min_ord = q_max_ord;
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 qv_ord = q_sign * q_b2[a];
q_max_ord = fmaxf(q_max_ord, qv_ord);
q_min_ord = fminf(q_min_ord, qv_ord);
}
/* Higher temperature floor (0.5 vs 0.01) for order: Q-values for different
* order types start nearly identical (same reward signal). The low 0.01 floor
@@ -939,7 +949,7 @@ extern "C" __global__ void experience_action_select(
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);
exps_ord[a] = expf((q_sign * q_b2[a] - q_max_ord) / tau_ord);
sum_e += exps_ord[a];
}
float ro = philox_uniform(i, timestep, rng_ctr++) * sum_e;
@@ -957,16 +967,17 @@ extern "C" __global__ void experience_action_select(
a3 = (r >= b3_size) ? b3_size - 1 : r;
} else {
/* Boltzmann softmax over urgency Q-values */
float q_max_urg = q_b3[0], q_min_urg = q_b3[0];
float q_max_urg = q_sign * q_b3[0], q_min_urg = q_max_urg;
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 qv_urg = q_sign * q_b3[a];
q_max_urg = fmaxf(q_max_urg, qv_urg);
q_min_urg = fminf(q_min_urg, qv_urg);
}
float tau_urg = fmaxf(q_max_urg - q_min_urg, 0.5f);
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);
exps_urg[a] = expf((q_sign * q_b3[a] - q_max_urg) / tau_urg);
sum_e += exps_urg[a];
}
float ru = philox_uniform(i, timestep, rng_ctr++) * sum_e;

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

@@ -737,6 +737,12 @@ pub struct GpuExperienceCollector {
learning_health_cache: f32,
/// D4/N4: last effective cf_ratio computed from learning_health_cache.
last_cf_ratio_eff: f32,
/// D7/N7: Contrarian override flag cached from trainer. When non-zero, the
/// experience kernel negates Q values before Boltzmann softmax, converting
/// argmax-favoring sampling to argmin-favoring. Used briefly to escape Q-uniform
/// attractors when the policy is systematically anti-correlated with market.
contrarian_active_cache: u8,
}
impl Drop for GpuExperienceCollector {
@@ -1305,6 +1311,7 @@ impl GpuExperienceCollector {
td_lambda_kernel,
learning_health_cache: 1.0, // D4/N4: assume healthy at construction
last_cf_ratio_eff: 0.5, // D4/N4: standard cf_ratio at healthy state
contrarian_active_cache: 0, // D7/N7: off by default
reward_rank_kernel,
reward_compute_abs_sharpe_kernel,
bitonic_sort_step_kernel,
@@ -1358,6 +1365,18 @@ impl GpuExperienceCollector {
self.last_cf_ratio_eff
}
/// D7/N7: propagate contrarian flag from trainer (called once per epoch when
/// the state machine updates). When active, the action-selection kernel negates
/// Q values before Boltzmann softmax → argmin-favoring sampling.
pub fn set_contrarian_active(&mut self, active: bool) {
self.contrarian_active_cache = if active { 1 } else { 0 };
}
/// D7/N7: return the cached contrarian flag.
pub fn contrarian_active(&self) -> u8 {
self.contrarian_active_cache
}
/// D4/N4: compute cf_ratio from cached health.
fn current_cf_ratio(&self) -> f32 {
(0.5_f32 + 0.3_f32 * (1.0 - self.learning_health_cache)).clamp(0.0, 1.0)
@@ -2371,6 +2390,7 @@ impl GpuExperienceCollector {
.arg(&(t as i32)) // timestep for stateless Philox RNG
.arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule)
.arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static)
.arg(&(self.contrarian_active_cache as i32)) // D7/N7: Q-negation flag
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_action_select t={t}: {e}"

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

@@ -2010,6 +2010,12 @@ impl DQNTrainer {
}
self.last_contrarian_active = Some(self.contrarian_active);
// D7/N7 Part B: forward the flag to the experience collector so the
// kernel negates Q values during Boltzmann sampling.
if let Some(ref mut collector) = self.gpu_experience_collector {
collector.set_contrarian_active(self.contrarian_active);
}
}
// HEALTH_DIAG: components are [0, 1] normalized. effective = hyperparams after health-adaptation. novels = mechanism states.