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

feat(isv): feature-level ISV gating — which features matter per regime

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-17 00:56:52 +02:00
parent 750b228ed9
commit d6a2e99670
2 changed files with 76 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

@@ -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<f32>, // [ISV_EMB_DIM] = [8]
branch_gate_buf: CudaSlice<f32>, // [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(&param_sizes, 78);
let b_gate = self.ptrs.params_ptr + padded_byte_offset(&param_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).