From 8a746fceb28ea77a301714e06bf9d95a664c5bbd Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 00:34:49 +0200 Subject: [PATCH] =?UTF-8?q?feat(isv):=20regime=20awareness=20=E2=80=94=20I?= =?UTF-8?q?SV=5FDIM=208=E2=86=9212,=204=20regime=20signals?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit regime_velocity (EMA of ADX+CUSUM deltas), regime_disagreement (|norm_ADX - norm_CUSUM|), regime_transition_ema, regime_stability (1 - sigmoid(5*velocity)). Read from states buffer sample 0. isv_signal_update + isv_forward kernels updated. w_isv_fc1 grows 16*8→16*12. Pinned buffers resized automatically via ISV_DIM. Introduced ISV_EMB_DIM=8 to decouple embedding output from signal input dimension — FC2/gate/gamma stay at 8. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 68 ++++++++++++++----- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 37 +++++----- 2 files changed, 72 insertions(+), 33 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 069c10c47..44bec964a 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -4816,8 +4816,8 @@ extern "C" __global__ void risk_budget_backward( /* ================================================================== */ extern "C" __global__ void isv_signal_update( - float* __restrict__ isv_signals, /* [8] pinned device-mapped */ - float* __restrict__ isv_history, /* [K*8] pinned device-mapped */ + float* __restrict__ isv_signals, /* [12] pinned device-mapped */ + float* __restrict__ isv_history, /* [K*12] pinned device-mapped */ float* __restrict__ lagged_td_error, /* [1] pinned — for recursive confidence target */ const float* __restrict__ q_mean_ptr, /* [1] pinned — batch Q-mean */ const float* __restrict__ q_mean_ema_ptr, /* [1] pinned — Q-mean EMA */ @@ -4827,7 +4827,9 @@ extern "C" __global__ void isv_signal_update( const float* __restrict__ reward_ptr, /* [1] pinned — reward proxy */ const float* __restrict__ atom_util_ptr, /* [1] pinned — atom utilization */ const float* __restrict__ loss_ptr, /* [1] pinned — total loss */ - int K /* temporal history length (4) */ + int K, /* temporal history length (4) */ + const float* __restrict__ states_ptr, /* [B, state_dim] — for ADX/CUSUM regime signals. NULL = skip regime. */ + int state_dim /* to index ADX at [40], CUSUM at [41] */ ) { if (threadIdx.x != 0) return; @@ -4838,9 +4840,9 @@ extern "C" __global__ void isv_signal_update( /* Shift temporal history: newest at [0], oldest at [K-1] */ for (int t = K - 1; t > 0; t--) - for (int k = 0; k < 8; k++) - isv_history[t * 8 + k] = isv_history[(t - 1) * 8 + k]; - for (int k = 0; k < 8; k++) + for (int k = 0; k < 12; k++) + isv_history[t * 12 + k] = isv_history[(t - 1) * 12 + k]; + for (int k = 0; k < 12; k++) isv_history[k] = isv_signals[k]; /* [0] Q-drift: normalized deviation from EMA */ @@ -4870,6 +4872,38 @@ extern "C" __global__ void isv_signal_update( /* [7] Loss EMA */ isv_signals[7] = (1.0f - alpha) * isv_signals[7] + alpha * loss_ptr[0]; + + /* ── Regime awareness signals [8]-[11] from ADX/CUSUM ──────────── */ + if (states_ptr != 0 && state_dim > 41) { + float adx = states_ptr[0 * state_dim + 40]; + float cusum = states_ptr[0 * state_dim + 41]; + + /* [8] Regime velocity: EMA of ADX (used for delta computation next step) */ + float prev_adx = isv_signals[8]; + float adx_delta = fabsf(adx - prev_adx); + float prev_cusum_ema = isv_signals[10]; + float cusum_delta = fabsf(cusum - prev_cusum_ema); + float regime_vel = adx_delta + cusum_delta; + + isv_signals[8] = (1.0f - alpha) * isv_signals[8] + alpha * adx; + float velocity_ema = (1.0f - alpha) * prev_cusum_ema + alpha * regime_vel; + + /* [9] Regime disagreement: |norm(ADX) - norm(CUSUM)| */ + float sum_abs = fabsf(adx) + fabsf(cusum) + 1e-6f; + isv_signals[9] = fabsf(adx / sum_abs - cusum / sum_abs); + + /* [10] Regime transition EMA */ + isv_signals[10] = velocity_ema; + + /* [11] Regime stability: 1 - sigmoid(5 * velocity) */ + isv_signals[11] = 1.0f - 1.0f / (1.0f + expf(-5.0f * regime_vel)); + } else { + /* No states available — neutral defaults */ + isv_signals[8] = 0.0f; + isv_signals[9] = 0.0f; + isv_signals[10] = 0.0f; + isv_signals[11] = 0.5f; /* moderate stability */ + } } /* ================================================================== */ @@ -4877,10 +4911,10 @@ extern "C" __global__ void isv_signal_update( /* ================================================================== */ extern "C" __global__ void isv_forward( - const float* __restrict__ isv_signals, /* [8] pinned */ - const float* __restrict__ isv_history, /* [K*8] pinned */ - const float* __restrict__ isv_decay, /* [8] pinned learned decay */ - const float* __restrict__ w_fc1, /* [16, 8] */ + const float* __restrict__ isv_signals, /* [12] pinned */ + const float* __restrict__ isv_history, /* [K*12] pinned */ + const float* __restrict__ isv_decay, /* [12] pinned learned decay */ + const float* __restrict__ w_fc1, /* [16, 12] */ const float* __restrict__ b_fc1, /* [16] */ const float* __restrict__ w_fc2, /* [8, 16] */ const float* __restrict__ b_fc2, /* [8] */ @@ -4895,26 +4929,26 @@ extern "C" __global__ void isv_forward( ) { if (threadIdx.x != 0) return; - /* Step 1: Temporal ISV — decay-weighted average */ - float isv_temporal[8]; - for (int k = 0; k < 8; k++) { + /* Step 1: Temporal ISV — decay-weighted average over 12 signals */ + float isv_temporal[12]; + for (int k = 0; k < 12; k++) { float decay = 1.0f / (1.0f + expf(-isv_decay[k])); float weighted_sum = isv_signals[k]; float weight_sum = 1.0f; for (int t = 0; t < K; t++) { float w = powf(decay, (float)(t + 1)); - weighted_sum += w * isv_history[t * 8 + k]; + weighted_sum += w * isv_history[t * 12 + k]; weight_sum += w; } isv_temporal[k] = weighted_sum / weight_sum; } - /* Step 2: MLP — FC1 (8→16, ReLU) → FC2 (16→8) */ + /* Step 2: MLP — FC1 (12→16, ReLU) → FC2 (16→8) */ float hidden[16]; for (int j = 0; j < 16; j++) { float val = b_fc1[j]; - for (int k = 0; k < 8; k++) - val += w_fc1[j * 8 + k] * isv_temporal[k]; + for (int k = 0; k < 12; k++) + val += w_fc1[j * 12 + k] * isv_temporal[k]; hidden[j] = fmaxf(val, 0.0f); } float embedding[8]; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index d8fce3d9e..8beb874a5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -82,7 +82,8 @@ const MAMBA2_STATE_DIM: usize = 16; // SSM state dimension /// Introspective State Vector (ISV) configuration. const ISV_K: usize = 4; // Temporal ISV history length -const ISV_DIM: usize = 8; // ISV signal count +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. const FIRST_ISV_TENSOR: usize = 68; @@ -479,13 +480,13 @@ pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT cfg.adv_h, // [66] w_risk_out [AH] 1, // [67] b_risk_out [1] // ── Introspective State Vector (ISV) encoder + conditioning ── - 16 * ISV_DIM, // [68] w_isv_fc1 [16, 8] + 16 * ISV_DIM, // [68] w_isv_fc1 [16, ISV_DIM] 16, // [69] b_isv_fc1 [16] - ISV_DIM * 16, // [70] w_isv_fc2 [8, 16] - ISV_DIM, // [71] b_isv_fc2 [8] - 4 * ISV_DIM, // [72] w_isv_gate [4, 8] + ISV_EMB_DIM * 16, // [70] w_isv_fc2 [EMB, 16] + ISV_EMB_DIM, // [71] b_isv_fc2 [EMB] + 4 * ISV_EMB_DIM, // [72] w_isv_gate [4, EMB] 4, // [73] b_isv_gate [4] - ISV_DIM, // [74] w_isv_gamma [8] + ISV_EMB_DIM, // [74] w_isv_gamma [EMB] 1, // [75] b_isv_gamma [1] cfg.shared_h2, // [76] w_conf_fc [SH2] 1, // [77] b_conf_fc [1] @@ -1337,11 +1338,11 @@ pub struct GpuDqnTrainer { atom_util_scratch_dev_ptr: u64, // ── ISV core buffers (pinned device-mapped, GPU read/write) ── - isv_signals_pinned: *mut f32, // [ISV_DIM = 8] + isv_signals_pinned: *mut f32, // [ISV_DIM = 12] isv_signals_dev_ptr: u64, - isv_history_pinned: *mut f32, // [ISV_K * ISV_DIM = 32] + isv_history_pinned: *mut f32, // [ISV_K * ISV_DIM = 48] isv_history_dev_ptr: u64, - isv_decay_pinned: *mut f32, // [ISV_DIM = 8] learned decay weights + isv_decay_pinned: *mut f32, // [ISV_DIM = 12] learned decay weights isv_decay_dev_ptr: u64, lagged_td_error_pinned: *mut f32, // [1] recursive confidence target lagged_td_error_dev_ptr: u64, @@ -1350,7 +1351,7 @@ pub struct GpuDqnTrainer { // ── ISV forward outputs ── isv_forward_kernel: CudaFunction, fill_gamma_buf_kernel: CudaFunction, - isv_embedding_buf: CudaSlice, // [ISV_DIM] = [8] + isv_embedding_buf: CudaSlice, // [ISV_EMB_DIM] = [8] branch_gate_buf: CudaSlice, // [4] gamma_mod_buf: CudaSlice, // [1] gamma_buf: CudaSlice, // [B] per-sample effective gamma @@ -2017,10 +2018,12 @@ impl GpuDqnTrainer { unsafe { *self.q_mean_ema_pinned = new_ema; } } - /// Pure GPU ISV signal update — reads all pinned pointers, updates ISV [8] + history [K,8]. + /// Pure GPU ISV signal update — reads all pinned pointers, updates ISV [12] + history [K,12]. + /// Includes 4 regime awareness signals computed from ADX/CUSUM in states buffer. /// Call after graph_adam replay and stream sync, before next graph_forward. pub(crate) fn update_isv_signals(&self) -> Result<(), MLError> { let k = ISV_K as i32; + let state_dim = self.config.state_dim as i32; unsafe { self.stream.launch_builder(&self.isv_signal_update_kernel) .arg(&self.isv_signals_dev_ptr) @@ -2035,6 +2038,8 @@ impl GpuDqnTrainer { .arg(&self.atom_util_scratch_dev_ptr) .arg(&self.total_loss_dev_ptr) .arg(&k) + .arg(&self.ptrs.states_buf) // states for ADX/CUSUM regime signals + .arg(&state_dim) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -2047,7 +2052,7 @@ impl GpuDqnTrainer { /// ISV encoder MLP forward: temporal decay → MLP → branch gate + gamma mod. /// Reads ISV signals/history (pinned), ISV encoder weights from param buffer. - /// Writes isv_embedding_buf [8], branch_gate_buf [4], gamma_mod_buf [1]. + /// Writes isv_embedding_buf [ISV_EMB_DIM=8], branch_gate_buf [4], gamma_mod_buf [1]. pub(crate) fn launch_isv_forward(&self) -> Result<(), MLError> { let param_sizes = compute_param_sizes(&self.config); let w_fc1 = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 68); @@ -5351,7 +5356,7 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("risk_grad_buf alloc: {e}")))?; // ── ISV forward output buffers ── - let isv_embedding_buf = stream.alloc_zeros::(ISV_DIM) + let isv_embedding_buf = stream.alloc_zeros::(ISV_EMB_DIM) .map_err(|e| MLError::ModelError(format!("isv_embedding_buf alloc: {e}")))?; let branch_gate_buf = stream.alloc_zeros::(4) .map_err(|e| MLError::ModelError(format!("branch_gate_buf alloc: {e}")))?; @@ -9213,11 +9218,11 @@ impl GpuDqnTrainer { (0, 0), // [66] w_risk_out (special init below) (0, 0), // [67] b_risk_out (zero → sigmoid(0)=0.5) // ── ISV encoder + conditioning ── - (16, ISV_DIM), // [68] w_isv_fc1 (Xavier) + (16, ISV_DIM), // [68] w_isv_fc1 (Xavier, ISV_DIM input) (0, 0), // [69] b_isv_fc1 (zero) - (ISV_DIM, 16), // [70] w_isv_fc2 (Xavier) + (ISV_EMB_DIM, 16), // [70] w_isv_fc2 (Xavier, EMB output) (0, 0), // [71] b_isv_fc2 (zero) - (4, ISV_DIM), // [72] w_isv_gate (Xavier) + (4, ISV_EMB_DIM), // [72] w_isv_gate (Xavier, EMB input) (0, 0), // [73] b_isv_gate (zero) (0, 0), // [74] w_isv_gamma (zero → sigmoid(0)=0.5) (0, 0), // [75] b_isv_gamma (zero)