From d6a2e996703b7d1242767d8d53fb606d75494244 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 00:56:52 +0200 Subject: [PATCH] =?UTF-8?q?feat(isv):=20feature-level=20ISV=20gating=20?= =?UTF-8?q?=E2=80=94=20which=20features=20matter=20per=20regime?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit isv_feature_gate kernel: ISV embedding [8] → per-feature gate [SH2]. h_s2 *= sigmoid(w_gate @ isv_emb + b_gate). Near pass-through init (bias=2.0 → sigmoid≈0.88). Model learns to amplify/suppress features based on ISV state. 2 new tensors [78-79]. NUM_WEIGHT_TENSORS: 78→80. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 22 +++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 57 ++++++++++++++++++- 2 files changed, 76 insertions(+), 3 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 119675996..fcafb9787 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -5102,3 +5102,25 @@ extern "C" __global__ void recursive_confidence_backward( atomicAdd(&d_h_s2[(long long)i * SH2 + k], d_sigmoid * w_conf[k]); } } + +/* ================================================================== */ +/* Kernel: isv_feature_gate — ISV embedding modulates h_s2 features */ +/* ================================================================== */ +extern "C" __global__ void isv_feature_gate( + float* __restrict__ h_s2, /* [B, SH2] in-place */ + const float* __restrict__ isv_embedding, /* [8] shared (from isv_forward) */ + const float* __restrict__ w_feature_gate, /* [SH2, 8] */ + const float* __restrict__ b_feature_gate, /* [SH2] */ + int B, int SH2 +) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= B) return; + + for (int j = 0; j < SH2; j++) { + float gate_val = b_feature_gate[j]; + for (int k = 0; k < 8; k++) + gate_val += w_feature_gate[(long long)j * 8 + k] * isv_embedding[k]; + float gate = 1.0f / (1.0f + expf(-gate_val)); + h_s2[(long long)i * SH2 + j] *= gate; + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 2be5dddef..f801c7608 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -85,7 +85,7 @@ const ISV_K: usize = 4; // Temporal ISV history length const ISV_DIM: usize = 12; // ISV signal count (8 core + 4 regime awareness) const ISV_EMB_DIM: usize = 8; // ISV embedding output dimension (FC2 output, gate/gamma input) /// First ISV tensor index in the flat param buffer. -/// ISV weights (68-77) are online-only — NOT synced to the target network. +/// ISV weights (68-79) are online-only — NOT synced to the target network. const FIRST_ISV_TENSOR: usize = 68; const NEW_COMPONENT_WARMUP_STEPS: u32 = 500; @@ -373,7 +373,7 @@ 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. -pub(crate) const NUM_WEIGHT_TENSORS: usize = 78; +pub(crate) const NUM_WEIGHT_TENSORS: usize = 80; /// Compute the size (element count) of each weight tensor. /// @@ -388,6 +388,7 @@ pub(crate) const NUM_WEIGHT_TENSORS: usize = 78; /// Tensors 56-63: Multi-horizon value heads (5-bar and 20-bar). /// Tensors 64-67: Learned risk management (5th branch). /// Tensors 68-77: Introspective State Vector (ISV) encoder + conditioning. +/// Tensors 78-79: ISV feature gate (per-feature h_s2 modulation). /// /// When bottleneck is active, w_s1 input dimension changes from state_dim to /// (bottleneck_dim + portfolio_dim) where portfolio_dim = state_dim - market_dim. @@ -490,6 +491,9 @@ pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT 1, // [75] b_isv_gamma [1] cfg.shared_h2, // [76] w_conf_fc [SH2] 1, // [77] b_conf_fc [1] + // ── ISV Feature Gating (Phase 3) ── + cfg.shared_h2 * ISV_EMB_DIM, // [78] w_feature_gate [SH2, 8] + cfg.shared_h2, // [79] b_feature_gate [SH2] ] } @@ -1350,6 +1354,7 @@ pub struct GpuDqnTrainer { // ── ISV forward outputs ── isv_forward_kernel: CudaFunction, + isv_feature_gate_kernel: CudaFunction, fill_gamma_buf_kernel: CudaFunction, isv_embedding_buf: CudaSlice, // [ISV_EMB_DIM] = [8] branch_gate_buf: CudaSlice, // [4] @@ -2088,6 +2093,35 @@ impl GpuDqnTrainer { Ok(()) } + /// ISV feature gate: modulate h_s2 features based on ISV embedding. + /// h_s2[i,j] *= sigmoid(w_feature_gate[j,:] @ isv_embedding + b_feature_gate[j]) + /// Operates in-place on save_h_s2. Must be called AFTER launch_isv_forward. + pub(crate) fn launch_isv_feature_gate(&self, batch_size: usize) -> Result<(), MLError> { + let param_sizes = compute_param_sizes(&self.config); + let w_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 78); + let b_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 79); + 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; + + unsafe { + self.stream.launch_builder(&self.isv_feature_gate_kernel) + .arg(&self.save_h_s2) // h_s2 in-place + .arg(&self.isv_embedding_buf) + .arg(&w_gate) + .arg(&b_gate) + .arg(&b_i32) + .arg(&sh2) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("isv_feature_gate: {e}")))?; + } + Ok(()) + } + /// Recursive confidence forward: predict own TD-error from h_s2. /// h_s2 → sigmoid(w_conf @ h + b_conf) → predicted_error [B]. pub(crate) fn launch_recursive_confidence_forward(&self, batch_size: usize) -> Result<(), MLError> { @@ -4258,6 +4292,8 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("isv_signal_update load: {e}")))?; let isv_forward_kernel = exp_module_for_mag.load_function("isv_forward") .map_err(|e| MLError::ModelError(format!("isv_forward load: {e}")))?; + let isv_feature_gate_kernel = exp_module_for_mag.load_function("isv_feature_gate") + .map_err(|e| MLError::ModelError(format!("isv_feature_gate load: {e}")))?; let fill_gamma_buf_kernel = exp_module_for_mag.load_function("fill_gamma_buf") .map_err(|e| MLError::ModelError(format!("fill_gamma_buf load: {e}")))?; let recursive_conf_fwd_kernel = exp_module_for_mag.load_function("recursive_confidence_forward") @@ -5759,6 +5795,7 @@ impl GpuDqnTrainer { lagged_td_error_dev_ptr, isv_signal_update_kernel, isv_forward_kernel, + isv_feature_gate_kernel, fill_gamma_buf_kernel, isv_embedding_buf, branch_gate_buf, @@ -7011,6 +7048,9 @@ impl GpuDqnTrainer { // ISV forward: encoder MLP → branch gate + gamma mod self.launch_isv_forward()?; + // ISV feature gate: modulate h_s2 features based on ISV regime embedding + self.launch_isv_feature_gate(batch_size)?; + // Recursive confidence: predict own TD-error from h_s2 self.launch_recursive_confidence_forward(batch_size)?; @@ -9132,7 +9172,7 @@ impl GpuDqnTrainer { /// matches the GOFF_* defines exactly, with `align4()` padding per tensor. /// /// Both params_buf and target_params_buf receive identical weights — the EMA - /// kernel will diverge them during training. ISV weights (indices 68-77) in + /// kernel will diverge them during training. ISV weights (indices 68-79) in /// target_params_buf are inert: EMA skips them, and the target forward never /// reads them. pub(crate) fn xavier_init_params_buf(&mut self) -> Result<(), MLError> { @@ -9230,6 +9270,9 @@ impl GpuDqnTrainer { (0, 0), // [75] b_isv_gamma (zero) (1, cfg.shared_h2), // [76] w_conf_fc (Xavier) (0, 0), // [77] b_conf_fc (zero) + // ── ISV Feature Gating (Phase 3) ── + (cfg.shared_h2, ISV_EMB_DIM), // [78] w_feature_gate (Xavier) + (0, 0), // [79] b_feature_gate (special init: 2.0) ]; // Build flat host buffer: Xavier init for weights, zeros for biases + padding. @@ -9290,6 +9333,14 @@ impl GpuDqnTrainer { } } + // Feature gate bias: init to 2.0 for near-pass-through (sigmoid(2)≈0.88) + { + let b_gate_offset: usize = sizes[..79].iter().map(|&s| align4(s)).sum(); + for j in 0..cfg.shared_h2 { + host_buf[b_gate_offset + j] = 2.0; + } + } + debug_assert_eq!(offset, total, "xavier_init: offset mismatch with total_params"); // HtoD to both params_buf and target_params_buf (same initial weights).