From 6cf09e08066fa20c71fbb1600be1cfc485657cce Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 01:17:57 +0200 Subject: [PATCH] feat(plan): trade_plan_forward kernel + 4 weight tensors (82->86) Plan MLP: h_s2 -> hidden[AH] -> plan_params[B, 6]. Outputs: target_bars[3-25], profit_target[0.05-1.5%], stop_loss[0.02-0.5%], scale_aggression[0-1], conviction[0-1], asymmetry[0.5-3.0]. NUM_WEIGHT_TENSORS: 82->86. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 41 ++++++++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 65 ++++++++++++++++++- crates/ml/src/trainers/dqn/fused_training.rs | 4 ++ 3 files changed, 107 insertions(+), 3 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 4cfcfb51e..f75bc1a9f 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -5138,3 +5138,44 @@ extern "C" __global__ void isv_feature_gate( h_s2[(long long)i * SH2 + j] *= gate; } } + +/* ================================================================== */ +/* Kernel: trade_plan_forward — plan head MLP → 6 plan parameters */ +/* ================================================================== */ +extern "C" __global__ void trade_plan_forward( + const float* __restrict__ h_s2, /* [B, SH2] */ + const float* __restrict__ w_plan_fc, /* [AH, SH2] */ + const float* __restrict__ b_plan_fc, /* [AH] */ + const float* __restrict__ w_plan_out, /* [6, AH] */ + const float* __restrict__ b_plan_out, /* [6] */ + float* __restrict__ plan_params, /* [B, 6] output */ + int B, int SH2, int AH +) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= B) return; + + const float* h = h_s2 + (long long)i * SH2; + float hidden[128]; + + for (int j = 0; j < AH; j++) { + float val = b_plan_fc[j]; + for (int k = 0; k < SH2; k++) + val += w_plan_fc[(long long)j * SH2 + k] * h[k]; + hidden[j] = fmaxf(val, 0.0f); + } + + float* out = plan_params + i * 6; + for (int p = 0; p < 6; p++) { + float val = b_plan_out[p]; + for (int j = 0; j < AH; j++) + val += w_plan_out[(long long)p * AH + j] * hidden[j]; + out[p] = val; + } + + out[0] = 3.0f + 22.0f / (1.0f + expf(-out[0])); + out[1] = 0.0005f + 0.0145f / (1.0f + expf(-out[1])); + out[2] = 0.0002f + 0.0048f / (1.0f + expf(-out[2])); + out[3] = 1.0f / (1.0f + expf(-out[3])); + out[4] = 1.0f / (1.0f + expf(-out[4])); + out[5] = 0.5f + 2.5f / (1.0f + expf(-out[5])); +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 0900cc358..4afad7f51 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -373,8 +373,9 @@ impl Default for CausalInterventionConfig { /// + 2 regime branch gate (w_regime + b_regime) + 4 adaptive atom spacing /// + 8 multi-horizon value heads (5-bar and 20-bar) /// + 4 learned risk management (5th branch) = 68 -/// + 10 ISV encoder/conditioning + 2 feature gate + 2 temporal routing = 82. -pub(crate) const NUM_WEIGHT_TENSORS: usize = 82; +/// + 10 ISV encoder/conditioning + 2 feature gate + 2 temporal routing = 82 +/// + 4 trade plan head = 86. +pub(crate) const NUM_WEIGHT_TENSORS: usize = 86; /// Compute the size (element count) of each weight tensor. /// @@ -391,6 +392,7 @@ pub(crate) const NUM_WEIGHT_TENSORS: usize = 82; /// Tensors 68-77: Introspective State Vector (ISV) encoder + conditioning. /// Tensors 78-79: ISV feature gate (per-feature h_s2 modulation). /// Tensors 80-81: ISV temporal routing (per-feature temporal depth). +/// Tensors 82-85: Trade plan head (h_s2 → 6 plan parameters). /// /// When bottleneck is active, w_s1 input dimension changes from state_dim to /// (bottleneck_dim + portfolio_dim) where portfolio_dim = state_dim - market_dim. @@ -499,6 +501,11 @@ pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT // ── ISV Temporal Routing (Phase 3) ── cfg.shared_h2 * ISV_EMB_DIM, // [80] w_temporal_route [SH2, 8] cfg.shared_h2, // [81] b_temporal_route [SH2] + // ── Trade Plan Head ── + cfg.adv_h * cfg.shared_h2, // [82] w_plan_fc [AH, SH2] + cfg.adv_h, // [83] b_plan_fc [AH] + 6 * cfg.adv_h, // [84] w_plan_out [6, AH] + 6, // [85] b_plan_out [6] ] } @@ -1372,6 +1379,10 @@ pub struct GpuDqnTrainer { // ── Recursive confidence kernels ── recursive_conf_fwd_kernel: CudaFunction, recursive_conf_bwd_kernel: CudaFunction, + + // ── Trade plan head ── + plan_params_buf: CudaSlice, // [B, 6] + trade_plan_fwd_kernel: CudaFunction, } impl GpuDqnTrainer { @@ -2183,6 +2194,37 @@ impl GpuDqnTrainer { Ok(()) } + /// Trade plan forward: h_s2 → hidden[AH] (ReLU) → plan_params[B, 6]. + /// Output activations: target_bars[3-25], profit_target[0.05-1.5%], + /// stop_loss[0.02-0.5%], scale_aggression[0-1], conviction[0-1], asymmetry[0.5-3.0]. + pub(crate) fn launch_trade_plan_forward(&self, batch_size: usize) -> Result<(), MLError> { + let param_sizes = compute_param_sizes(&self.config); + let w_fc = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 82); + let b_fc = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 83); + let w_out = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 84); + let b_out = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 85); + let blocks = ((batch_size as u32 + 255) / 256).max(1); + let b_i32 = batch_size as i32; + let sh2 = self.config.shared_h2 as i32; + let ah = self.config.adv_h as i32; + + unsafe { + self.stream.launch_builder(&self.trade_plan_fwd_kernel) + .arg(&self.save_h_s2) + .arg(&w_fc).arg(&b_fc) + .arg(&w_out).arg(&b_out) + .arg(&self.plan_params_buf) + .arg(&b_i32).arg(&sh2).arg(&ah) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("trade_plan_forward: {e}")))?; + } + Ok(()) + } + /// Recursive confidence backward: MSE loss gradient into trunk + conf weight gradients. /// Accumulates into grad_buf (same buffer Adam reads) and bw_d_h_s2 trunk gradient. pub(crate) fn launch_recursive_confidence_backward(&self, batch_size: usize) -> Result<(), MLError> { @@ -4334,7 +4376,9 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("recursive_confidence_forward load: {e}")))?; let recursive_conf_bwd_kernel = exp_module_for_mag.load_function("recursive_confidence_backward") .map_err(|e| MLError::ModelError(format!("recursive_confidence_backward load: {e}")))?; - info!("GpuDqnTrainer: mag_concat + strided_accumulate/scatter + concat_ofi + regime_gate + adaptive_atom + atom_grad + q_anchor + regime_dropout + G5/G6/G10/G12 + risk_budget + isv_signal_update + isv_forward + fill_gamma_buf kernels loaded"); + let trade_plan_fwd_kernel = exp_module_for_mag.load_function("trade_plan_forward") + .map_err(|e| MLError::ModelError(format!("trade_plan_forward load: {e}")))?; + info!("GpuDqnTrainer: mag_concat + strided_accumulate/scatter + concat_ofi + regime_gate + adaptive_atom + atom_grad + q_anchor + regime_dropout + G5/G6/G10/G12 + risk_budget + isv_signal_update + isv_forward + fill_gamma_buf + trade_plan_forward kernels loaded"); // ── G5: Epistemic-gated magnitude — pinned var_ema threshold ─ let (var_ema_pinned, var_ema_dev_ptr) = { @@ -5442,6 +5486,8 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("gamma_buf alloc: {e}")))?; let predicted_error_buf = stream.alloc_zeros::(b) .map_err(|e| MLError::ModelError(format!("predicted_error_buf alloc: {e}")))?; + let plan_params_buf = stream.alloc_zeros::(b * 6) + .map_err(|e| MLError::ModelError(format!("plan_params_buf alloc: {e}")))?; Ok(Self { config, @@ -5844,6 +5890,8 @@ impl GpuDqnTrainer { predicted_error_buf, recursive_conf_fwd_kernel, recursive_conf_bwd_kernel, + plan_params_buf, + trade_plan_fwd_kernel, }) } @@ -7098,6 +7146,9 @@ impl GpuDqnTrainer { // Recursive confidence: predict own TD-error from h_s2 self.launch_recursive_confidence_forward(batch_size)?; + // Trade plan forward: h_s2 → plan_params [B, 6] + self.launch_trade_plan_forward(batch_size)?; + // Risk budget forward: h_s2 → risk_budget R ∈ (0,1) before Q-value computation self.risk_budget_forward(batch_size)?; @@ -9320,6 +9371,11 @@ impl GpuDqnTrainer { // ── ISV Temporal Routing (Phase 3) ── (cfg.shared_h2, ISV_EMB_DIM), // [80] w_temporal_route (Xavier) (0, 0), // [81] b_temporal_route (special init: 2.0) + // ── Trade Plan Head ── + (cfg.adv_h, cfg.shared_h2), // [82] w_plan_fc (Xavier) + (0, 0), // [83] b_plan_fc (zero) + (6, cfg.adv_h), // [84] w_plan_out (Xavier) + (0, 0), // [85] b_plan_out (zero) ]; // Build flat host buffer: Xavier init for weights, zeros for biases + padding. @@ -9661,6 +9717,9 @@ impl GpuDqnTrainer { /// Raw pointer to states_buf for pad_states kernel. pub(crate) fn states_buf_ptr(&self) -> u64 { self.states_buf.raw_ptr() } + /// Raw pointer to plan_params_buf [B, 6] for trade plan integration. + pub fn plan_params_buf_ptr(&self) -> u64 { self.plan_params_buf.raw_ptr() } + /// F32 states buffer ref for cold-path Q-value computation. pub(crate) fn states_buf_f32(&self) -> &CudaSlice { &self.states_buf } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index fb2dd0087..9875c4fa4 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -2241,6 +2241,10 @@ impl FusedTrainingCtx { pub(crate) fn num_atoms(&self) -> usize { self.trainer.config().num_atoms } pub(crate) fn update_adaptive_clip(&mut self, grad_norm: f32) { self.trainer.update_adaptive_clip(grad_norm); } pub(crate) fn adaptive_clip_value(&self) -> f32 { self.trainer.adaptive_clip_value() } + /// Raw pointer to plan_params_buf [B, 6] for trade plan integration. + pub(crate) fn plan_params_buf_ptr(&self) -> u64 { + self.trainer.plan_params_buf_ptr() + } /// Raw device pointer to the flat f32 target params buffer (for DtoD sync). pub(crate) fn target_params_flat_ptr(&self) -> u64 {