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

feat(isv): regime awareness — ISV_DIM 8→12, 4 regime signals

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-17 00:34:49 +02:00
parent 8b68954ded
commit 8a746fceb2
2 changed files with 72 additions and 33 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

@@ -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<f32>, // [ISV_DIM] = [8]
isv_embedding_buf: CudaSlice<f32>, // [ISV_EMB_DIM] = [8]
branch_gate_buf: CudaSlice<f32>, // [4]
gamma_mod_buf: CudaSlice<f32>, // [1]
gamma_buf: CudaSlice<f32>, // [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(&param_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::<f32>(ISV_DIM)
let isv_embedding_buf = stream.alloc_zeros::<f32>(ISV_EMB_DIM)
.map_err(|e| MLError::ModelError(format!("isv_embedding_buf alloc: {e}")))?;
let branch_gate_buf = stream.alloc_zeros::<f32>(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)