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

feat(plan): trade_plan_forward kernel + 4 weight tensors (82->86)

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-17 01:17:57 +02:00
parent 242c18f4ab
commit 6cf09e0806
3 changed files with 107 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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]));
}

65
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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<f32>, // [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(&param_sizes, 82);
let b_fc = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 83);
let w_out = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 84);
let b_out = self.ptrs.params_ptr + padded_byte_offset(&param_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::<f32>(b)
.map_err(|e| MLError::ModelError(format!("predicted_error_buf alloc: {e}")))?;
let plan_params_buf = stream.alloc_zeros::<f32>(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<f32> { &self.states_buf }

4
crates/ml/src/trainers/dqn/fused_training.rs
View File

@@ -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 {