diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index e85ce041c..b74d69f95 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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; diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 5e8c37d64..c3bc6126f 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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}" diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 9a69e7334..70f01327f 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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.