diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 07107551e..23df9e388 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -392,7 +392,9 @@ extern "C" __global__ void experience_state_gather( /* #11 Time-reversal: reverse bar indexing for selected episodes. * If > 0, episodes with index % time_reversal_mod == 0 are played backwards. * 5 = 20% of episodes reversed, 3 = 33%, 0 = disabled. */ - int time_reversal_mod + int time_reversal_mod, + const float* __restrict__ plan_params_ptr, /* [N, 6] trade plan params. NULL = no plan. */ + const float* __restrict__ isv_signals_ptr /* [12] pinned device-mapped ISV signals. NULL = static. */ ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -572,7 +574,7 @@ extern "C" __global__ void experience_state_gather( : 0.0f; int portfolio_base = market_dim; - if (portfolio_base + 11 < state_dim) { + if (portfolio_base + 13 < state_dim) { out[portfolio_base + 0] = f_position; /* raw position */ out[portfolio_base + 1] = f_unrealized_pnl / f_equity; /* trade P&L signal */ out[portfolio_base + 2] = f_drawdown; /* risk: how deep are we? */ @@ -597,6 +599,20 @@ extern "C" __global__ void experience_state_gather( ? fminf(-f_unrealized_pnl / (plan_stop * f_equity + 1e-6f), 2.0f) /* pnl vs stop */ : 0.0f; out[portfolio_base + 11] = plan_conv; /* conviction at entry [0, 1] */ + + /* P11: Conviction drift — current conviction minus entry conviction. + * Negative drift = model's thesis is weakening → exit signal. + * Only meaningful when a plan is active (ps[23] > 0.5). */ + out[portfolio_base + 12] = (plan_params_ptr != NULL && ps[23] > 0.5f) + ? (plan_params_ptr[i * 6 + 4] - ps[27]) /* conviction drift */ + : 0.0f; + + /* P12: Regime shift since entry — |regime_stability_now - regime_stability_at_entry|. + * Large shift = regime changed mid-trade, plan may be invalid. + * Entry regime stored in ps[29] during plan activation. */ + out[portfolio_base + 13] = (isv_signals_ptr != NULL && ps[23] > 0.5f) + ? fabsf(isv_signals_ptr[11] - ps[29]) /* regime shift */ + : 0.0f; } /* -- Multi-timeframe features: [market_dim+8 .. market_dim+8+16) -- @@ -1481,6 +1497,8 @@ extern "C" __global__ void experience_env_step( ps[26] = pp[3]; /* scale_aggression */ ps[27] = pp[4]; /* conviction */ ps[28] = pp[5]; /* asymmetry */ + /* Store entry regime_stability for drift detection (P12). */ + ps[29] = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f; /* Apply conviction to position size — scale by readiness for smooth ramp */ position *= fmaxf(ps[27] * readiness, 0.1f); } diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index fd24cfd36..b9e8a9e0a 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -1919,6 +1919,10 @@ impl GpuExperienceCollector { .map(|b| b.device_ptr(&self.stream).0) .unwrap_or(0u64); + // Plan params and ISV signals pointers for conviction drift + regime shift features + let plan_ptr = self.plan_params_dev_ptr; + let isv_ptr = self.isv_signals_dev_ptr; + self.stream .launch_builder(&self.state_gather_kernel) .arg(market_features_buf) @@ -1935,6 +1939,8 @@ impl GpuExperienceCollector { .arg(&vol_norm) // #13 vol normalization .arg(&mirror_i32) // #10 mirror universe .arg(&config.time_reversal_mod) // #11 time-reversal + .arg(&plan_ptr) // plan_params_ptr (NULL=no plan) + .arg(&isv_ptr) // isv_signals_ptr (NULL=static) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_state_gather t={t}: {e}"