diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 0486af8af..558c1dbf9 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -4944,3 +4944,43 @@ extern "C" __global__ void fill_gamma_buf( if (i >= B) return; gamma_buf[i] = base_gamma * gamma_mod[0]; } + +/* ================================================================== */ +/* Kernel: branch_confidence_routing — ISV gate × Q-value confidence */ +/* ================================================================== */ +extern "C" __global__ void branch_confidence_routing( + float* __restrict__ q_values, /* [B, total_actions] in-place */ + const float* __restrict__ branch_gate, /* [4] shared from isv_forward */ + int B, + int b0_size, int b1_size, int b2_size, int b3_size +) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= B) return; + + int branch_sizes[4] = { b0_size, b1_size, b2_size, b3_size }; + int total_actions = b0_size + b1_size + b2_size + b3_size; + int branch_offset = 0; + + for (int d = 0; d < 4; d++) { + int A_d = branch_sizes[d]; + + float q_max = -1e9f; + float q_sum = 0.0f; + for (int a = 0; a < A_d; a++) { + float q = q_values[(long long)i * total_actions + branch_offset + a]; + q_max = fmaxf(q_max, q); + q_sum += q; + } + float q_mean_branch = q_sum / (float)A_d; + float separation = q_max - q_mean_branch; + float confidence = 1.0f / (1.0f + expf(-5.0f * separation)); + + float isv_gate = branch_gate[d]; + float effective_weight = isv_gate * fmaxf(confidence, 0.3f); + + for (int a = 0; a < A_d; a++) { + q_values[(long long)i * total_actions + branch_offset + a] *= effective_weight; + } + branch_offset += A_d; + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 6b8af79a5..895ceb5b7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -749,8 +749,8 @@ pub struct GpuDqnTrainer { kan_gate_combine_kernel: CudaFunction, kan_gate_backward_kernel: CudaFunction, - // ── Regime branch gate ── - regime_gate_kernel: CudaFunction, + // ── Branch confidence routing (ISV gate × Q-value confidence) ── + branch_confidence_routing_kernel: CudaFunction, /// Per-sample Q-gap for regime gate input [B]. regime_q_gap_buf: CudaSlice, /// Atom utilization scalar for regime gate [1] — pinned device-mapped. @@ -2126,48 +2126,34 @@ impl GpuDqnTrainer { Ok(()) } - /// Apply regime branch gate to Q-values (after Q-mean centering, before Q-attention). + /// Apply branch confidence routing: ISV gate × Q-value separation confidence. /// - /// Uses W_regime[4,4] + b_regime[4] to produce softmax importance weights from - /// [ADX, CUSUM, Q_gap, atom_utilization]. Scales per-branch Q-values so direction - /// matters more in trends, magnitude matters more in ranges. - pub(crate) fn apply_regime_gate(&self, batch_size: usize) -> Result<(), MLError> { - let param_sizes = compute_param_sizes(&self.config); - let w_regime_ptr = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 50); - let b_regime_ptr = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 51); - + /// For each branch, computes confidence = sigmoid(5 * (Q_max - Q_mean)) and + /// scales Q-values by isv_gate[d] * max(confidence, 0.3). Market regime info + /// (ADX, CUSUM) still flows through the trunk → branches; ISV adds training + /// dynamics awareness on top. + pub(crate) fn apply_branch_confidence_routing(&self, batch_size: usize) -> Result<(), MLError> { let blocks = ((batch_size as u32 + 255) / 256).max(1); let q_out_ptr = self.q_out_buf.raw_ptr(); - let states_ptr = self.ptrs.states_buf; - let q_gap_ptr = self.regime_q_gap_buf.raw_ptr(); - let util_ptr = self.regime_util_dev_ptr; + let gate_ptr = self.branch_gate_buf.raw_ptr(); let b_i32 = batch_size as i32; - let sd = self.config.state_dim as i32; let b0 = self.config.branch_0_size as i32; let b1 = self.config.branch_1_size as i32; let b2 = self.config.branch_2_size as i32; let b3 = self.config.branch_3_size as i32; unsafe { - self.stream.launch_builder(&self.regime_gate_kernel) + self.stream.launch_builder(&self.branch_confidence_routing_kernel) .arg(&q_out_ptr) - .arg(&states_ptr) - .arg(&q_gap_ptr) - .arg(&util_ptr) - .arg(&w_regime_ptr) - .arg(&b_regime_ptr) + .arg(&gate_ptr) .arg(&b_i32) - .arg(&sd) - .arg(&b0) - .arg(&b1) - .arg(&b2) - .arg(&b3) + .arg(&b0).arg(&b1).arg(&b2).arg(&b3) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: 0, }) - .map_err(|e| MLError::ModelError(format!("regime_branch_gate launch: {e}")))?; + .map_err(|e| MLError::ModelError(format!("branch_confidence_routing: {e}")))?; } Ok(()) } @@ -2305,7 +2291,7 @@ impl GpuDqnTrainer { /// G5: Apply epistemic gate to magnitude branch Q-values. /// Scales magnitude Q-values by sigmoid(5 * (var_mean - threshold)): /// high ensemble variance → conservative (small) magnitude bias. - /// Must be called AFTER apply_regime_gate and BEFORE launch_q_attention. + /// Must be called AFTER apply_branch_confidence_routing and BEFORE launch_q_attention. pub(crate) fn apply_epistemic_gate(&self, batch_size: usize) -> Result<(), MLError> { let ta = self.total_actions() as i32; unsafe { @@ -4155,8 +4141,8 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("strided_scatter load: {e}")))?; let concat_ofi_kernel = exp_module_for_mag.load_function("concat_ofi_features") .map_err(|e| MLError::ModelError(format!("concat_ofi_features load: {e}")))?; - let regime_gate_kernel = exp_module_for_mag.load_function("regime_branch_gate") - .map_err(|e| MLError::ModelError(format!("regime_branch_gate load: {e}")))?; + let branch_confidence_routing_kernel = exp_module_for_mag.load_function("branch_confidence_routing") + .map_err(|e| MLError::ModelError(format!("branch_confidence_routing load: {e}")))?; let adaptive_atom_kernel = exp_module_for_mag.load_function("adaptive_atom_positions") .map_err(|e| MLError::ModelError(format!("adaptive_atom_positions load: {e}")))?; let atom_position_grad_kernel = exp_module_for_mag.load_function("atom_position_gradient") @@ -5390,7 +5376,7 @@ impl GpuDqnTrainer { glu_backward_kernel, kan_gate_combine_kernel, kan_gate_backward_kernel, - regime_gate_kernel, + branch_confidence_routing_kernel, regime_q_gap_buf, regime_util_pinned, regime_util_dev_ptr, @@ -6965,8 +6951,8 @@ impl GpuDqnTrainer { // Apply risk budget: scale magnitude Q-values, produce CVaR alpha + commitment lambda self.apply_risk_budget(batch_size)?; - // Regime branch gate: scale Q-values by learned per-branch importance - self.apply_regime_gate(batch_size)?; + // Branch confidence routing: ISV gate × Q-value separation confidence + self.apply_branch_confidence_routing(batch_size)?; // G5: Epistemic gate — high variance → conservative magnitude self.apply_epistemic_gate(batch_size)?;