diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index f75bc1a9f..222cc502a 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -987,6 +987,17 @@ extern "C" __global__ void experience_action_select( } } + /* Plan direction lock: during active plan, force current direction */ + if (portfolio_states != NULL) { + int ps_base_plan = i * PORTFOLIO_STRIDE; + int has_plan_active = (portfolio_states[ps_base_plan + 23] > 0.5f); + if (has_plan_active) { + float cur_pos_plan = portfolio_states[ps_base_plan + 0]; + if (cur_pos_plan > 0.001f) dir_idx = 2; /* Long locked */ + else if (cur_pos_plan < -0.001f) dir_idx = 0; /* Short locked */ + } + } + /* Compose factored action: 4-branch encoding * action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency */ int action_idx = dir_idx * b1_size * b2_size * b3_size @@ -1099,7 +1110,8 @@ extern "C" __global__ void experience_env_step( int total_actions_for_var, /* b0+b1+b2+b3 to index q_variance */ const float* __restrict__ cost_anneal_ptr, /* [1] pinned device-mapped: 0.0=no costs, 1.0=full costs */ const float* __restrict__ commit_lambda_buf, /* [N] per-sample commitment lambda. NULL = use 0.01. */ - const float* __restrict__ isv_signals_ptr /* [8] pinned device-mapped ISV signals. NULL = static hold. */ + const float* __restrict__ isv_signals_ptr, /* [8] pinned device-mapped ISV signals. NULL = static hold. */ + const float* __restrict__ plan_params_ptr /* [N, 6] trade plan params. NULL = no plan. */ ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1427,6 +1439,51 @@ extern "C" __global__ void experience_env_step( tx_cost_multiplier, spread_cost, max_position, order_type_idx, spread_scale); + /* ── Trade Plan: activation on Flat → Positioned ── */ + int was_flat_plan = (fabsf(ps0_f) < 0.001f); + int now_positioned_plan = (fabsf(position) > 0.001f); + + if (was_flat_plan && now_positioned_plan && plan_params_ptr != NULL) { + const float* pp = plan_params_ptr + i * 6; + ps[23] = pp[0]; /* target_bars */ + ps[24] = pp[1]; /* profit_target */ + ps[25] = pp[2]; /* stop_loss */ + ps[26] = pp[3]; /* scale_aggression */ + ps[27] = pp[4]; /* conviction */ + ps[28] = pp[5]; /* asymmetry */ + /* Apply conviction to position size */ + position *= fmaxf(ps[27], 0.1f); + } + + /* ── Trade Plan: enforcement (auto-exit conditions) ── */ + int has_plan = (ps[23] > 0.5f); + if (has_plan && fabsf(position) > 0.001f) { + float pnl_pct = (raw_close - entry_price) / fmaxf(fabsf(entry_price), 1.0f) + * ((position > 0.0f) ? 1.0f : -1.0f); + + /* ISV-modulated thresholds */ + float stability = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f; + float eff_profit = ps[24] * ps[28] * (0.5f + 0.5f * stability); + float eff_stop = ps[25] * fmaxf(stability, 0.5f); + + int plan_exit = 0; + if (pnl_pct >= eff_profit) plan_exit = 1; + if (pnl_pct <= -eff_stop) plan_exit = 1; + if (hold_time >= ps[23]) plan_exit = 1; + if (stability < 0.3f && hold_time > 3.0f) plan_exit = 1; + + /* Scale schedule: ramp over first 3 bars */ + if (hold_time < 3.0f && !plan_exit) { + float scale = ps[26] + (1.0f - ps[26]) * hold_time / 3.0f; + position *= fminf(scale, 1.0f); + } + + if (plan_exit) { + position = 0.0f; + for (int s = 23; s <= 29; s++) ps[s] = 0.0f; + } + } + /* ---- Mark-to-market PnL for this timestep ---- */ /* Per-bar mark-to-market P&L removed: used raw_next (future price), * creating 1-bar action-reward misalignment. Rewards are now diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index be7df108c..99e85ad91 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -678,6 +678,10 @@ pub struct GpuExperienceCollector { /// Set via set_isv_signals_ptr(). 0 = NULL (static hold). isv_signals_dev_ptr: u64, + /// Trade plan params dev_ptr from fused training context. + /// Set via set_plan_params_ptr(). 0 = NULL (no plan). + plan_params_dev_ptr: u64, + /// v8: Per-bar curriculum difficulty scoring kernel. difficulty_scores_kernel: CudaFunction, /// v8: Hindsight experience relabeling kernel (optimal exit). @@ -1203,6 +1207,7 @@ impl GpuExperienceCollector { cost_anneal_pinned, cost_anneal_dev_ptr, isv_signals_dev_ptr: 0, // NULL until trainer sets it + plan_params_dev_ptr: 0, // NULL until trainer sets it difficulty_scores_kernel, hindsight_relabel_kernel, td_lambda_kernel, @@ -1227,6 +1232,13 @@ impl GpuExperienceCollector { self.isv_signals_dev_ptr = dev_ptr; } + /// Set trade plan params device pointer for plan activation/enforcement. + /// The env_step kernel reads plan params to copy into portfolio state on + /// Flat→Positioned transitions. Pass 0 to disable (NULL = no plan). + pub fn set_plan_params_ptr(&mut self, dev_ptr: u64) { + self.plan_params_dev_ptr = dev_ptr; + } + /// Set CVaR position scaling from IQN dual-head. /// The device pointer will be passed to the env_step kernel. /// Call with 0 to disable (NULL pointer = no scaling). @@ -2323,6 +2335,7 @@ impl GpuExperienceCollector { .arg(&self.cost_anneal_dev_ptr) // G2: transaction cost curriculum .arg(&0u64) // commit_lambda_buf: NULL = use 0.01 fallback (risk branch not in experience path yet) .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0 = NULL = static hold) + .arg(&self.plan_params_dev_ptr) // trade plan params (0 = NULL = no plan) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step 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 2912c67d4..c9ae3ac07 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -526,6 +526,14 @@ impl DQNTrainer { } } + // Trade plan params for plan activation/enforcement + if let Some(ref fused) = self.fused_ctx { + let plan_ptr = fused.plan_params_buf_ptr(); + if let Some(ref mut collector) = self.gpu_experience_collector { + collector.set_plan_params_ptr(plan_ptr); + } + } + // IQR exploration bonus from IQN quantile spread if let Some(ref fused) = self.fused_ctx { let iqr_ptr = fused.iqn_iqr_ptr().unwrap_or(0);