diff --git a/docs/superpowers/plans/2026-04-14-cross-pollinated-branch-intelligence.md b/docs/superpowers/plans/2026-04-14-cross-pollinated-branch-intelligence.md new file mode 100644 index 000000000..732c2600d --- /dev/null +++ b/docs/superpowers/plans/2026-04-14-cross-pollinated-branch-intelligence.md @@ -0,0 +1,1603 @@ +# Cross-Pollinated Branch Intelligence Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Combine Variable Selection (TFT), Gated Linear Units (TFT/Llama), Cross-Branch Q Attention, Liquid adaptive tau, and Mamba-2 selective replay into a unified 4-branch DQN architecture that selects features per-branch, gates activations smoothly, coordinates Q-values across branches, adapts learning rates per-branch, and prioritizes replay based on model state. + +**Architecture:** Five additive layers on the existing 4-branch DQN. (1) Per-branch softmax feature masks via two-GEMM bottleneck (R=16). (2) GLU replaces ReLU in branch FC via parallel gate GEMM + sigmoid multiply. (3) 2-head self-attention on 12 Q-values for cross-branch coordination. (4) Per-branch liquid ODE modulates gradient scales via pinned device-mapped floats. (5) Learned selectivity gate on trunk activations scales PER priorities. Dead code cleanup (11 items) precedes all construction. NUM_WEIGHT_TENSORS expands from 26 to 42, plus two separate small buffers (Q-attention: 624 params, selectivity: 257 params). + +**Tech Stack:** Rust (cudarc), CUDA (precompiled cubins via build.rs), cuBLAS (cublasLtMatmul TF32), pinned device-mapped memory for CPU-GPU scalar passing. + +**Spec:** `docs/superpowers/specs/2026-04-14-cross-pollinated-branch-intelligence-design.md` + +--- + +## Task 1: Dead Code Cleanup + +Remove 11 dead code items across 6 files. Must compile clean before adding new components. + +**Files:** +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` +- `crates/ml/src/cuda_pipeline/c51_loss_kernel.cu` +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- `crates/ml/src/cuda_pipeline/gpu_training_guard.rs` +- `crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs` +- `crates/ml/src/cuda_pipeline/decision_transformer.rs` + +- [ ] **Step 1: Remove `_reward_norm_kernel` field and load from `gpu_experience_collector.rs`** + +In `gpu_experience_collector.rs`: +- Delete the field at line 526: `_reward_norm_kernel: CudaFunction, // Kept for kernel compilation but never launched (reward v4)` +- Delete the assignment at line 1094: `_reward_norm_kernel: reward_norm_kernel,` +- Delete the `reward_norm_kernel` variable binding in the constructor (search for `let reward_norm_kernel =` or the `load_function("reward_normalize_kernel")` call and its result binding). +- Do NOT delete the kernel function from the `.cu` file (other kernels share the file). + +- [ ] **Step 2: Remove `cea_weight` parameter from `experience_kernels.cu` and all Rust launch sites** + +In `experience_kernels.cu` at line 1041: +- Remove the parameter `float cea_weight,` from `experience_env_step` function signature. +- Remove any usage of `cea_weight` inside the kernel body (search for `cea_weight` within the function). + +In `gpu_experience_collector.rs`: +- Delete the config field at line 181: `pub cea_weight: f32,` +- Delete the default value at line 333: `cea_weight: 0.0,` +- Delete the local variable at line 2147: `let cea_w = config.cea_weight as f32;` +- Remove `.arg(&cea_w)` from the `env_step_kernel` launch at line 2168. + +Grep for any other references to `cea_weight` across the crate and remove them. + +- [ ] **Step 3: Remove `mixup_alpha`, `mixup_seed`, `mixup_barrier` params from `c51_loss_kernel.cu` and Rust launch sites** + +In `c51_loss_kernel.cu`: +- Remove the 3 parameters from the main `c51_loss_batched` kernel signature at lines 223-225: + ``` + float mixup_alpha, + unsigned int mixup_seed, + int* __restrict__ mixup_barrier, + ``` +- Remove the 3 parameters from the `c51_mixup_ce` kernel signature at lines 701-702: + ``` + float mixup_alpha, + unsigned int mixup_seed, + ``` +- Remove all code inside these kernels that references `mixup_alpha`, `mixup_seed`, or `mixup_barrier`. + +In `gpu_dqn_trainer.rs`: +- Remove the config field at line 204: `pub mixup_alpha: f32,` +- Remove the default value at line 274: `mixup_alpha: 0.2,` +- Delete the `mixup_barrier_buf` field at line 817: `mixup_barrier_buf: CudaSlice,` +- Delete the allocation at line 2972-2973: + ```rust + let mixup_barrier_buf = stream.alloc_zeros::(3) + .map_err(|e| MLError::ModelError(format!("alloc mixup_barrier: {e}")))?; + ``` +- Delete the struct init at line 3329: `mixup_barrier_buf,` +- Delete the `mixup_seed` local at line 5432: `let mixup_seed: u32 = ...;` +- Remove the 3 `.arg(...)` lines at lines 5488-5490: + ```rust + .arg(&self.config.mixup_alpha) + .arg(&mixup_seed) + .arg(&self.mixup_barrier_buf) + ``` +- Delete the entire `launch_c51_mixup` method (around line 5506-5560) and the call to it at line 4805: `self.launch_c51_mixup()?;` +- Remove the `launch_loss_reduce` call after `launch_c51_mixup` at line 4806. +- Grep for any other `mixup_alpha`, `mixup_seed`, `mixup_barrier` references and remove them. + +- [ ] **Step 4: Remove `clipped_saxpy_kernel` field, `apply_cql_clipped_saxpy` method, and `clip_grad_buf_inplace` method from `gpu_dqn_trainer.rs`** + +In `gpu_dqn_trainer.rs`: +- Delete the field at lines 537-538: + ```rust + #[allow(dead_code)] + clipped_saxpy_kernel: CudaFunction, + ``` +- Delete the field at lines 519-522: + ```rust + /// Separate instance of grad_norm for non-graph launches (clip_grad_buf_inplace). + /// CUDA doesn't allow the same CUfunction to be launched both inside a captured + /// graph and outside on the same stream. + grad_norm_standalone: CudaFunction, + ``` + (Only if `grad_norm_standalone` is exclusively used by `clip_grad_buf_inplace`. Verify by grepping.) +- Delete the entire `apply_cql_clipped_saxpy` method (lines 2022-2083). +- Delete the entire `clip_grad_buf_inplace` method (lines 2196-2229). +- Remove the `clipped_saxpy` from the `compile_training_kernels` return tuple and function body (lines 6638-6639 and the return value). +- Remove the `clipped_saxpy_kernel` from the constructor struct init (line 3128). + +- [ ] **Step 5: Remove `debug_forward_no_graph` method from `gpu_dqn_trainer.rs`** + +Delete the entire `debug_forward_no_graph` method (lines 4717-4743): +```rust +#[allow(dead_code)] +fn debug_forward_no_graph(&mut self) -> Result<(), MLError> { + // ... (entire method body) +} +``` + +- [ ] **Step 6: Remove dead code blocks from `decision_transformer.rs` and `gpu_curiosity_trainer.rs`** + +In `decision_transformer.rs` at line 207: +- Remove the `#[allow(dead_code)]` attribute and the `LayerOffsets` struct that follows it (lines 207-213). If other dead code blocks exist with `#[allow(dead_code)]`, evaluate and remove if truly unused. + +In `gpu_curiosity_trainer.rs` at line 116: +- Remove the `#[allow(dead_code)]` attribute and the `launch_adam_step` free function that follows it (lines 116-153). + +- [ ] **Step 7: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +Fix any compilation errors from the removals. All references must be cleaned up. + +- [ ] **Step 8: Commit dead code cleanup** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \ + crates/ml/src/cuda_pipeline/experience_kernels.cu \ + crates/ml/src/cuda_pipeline/c51_loss_kernel.cu \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/cuda_pipeline/gpu_training_guard.rs \ + crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs \ + crates/ml/src/cuda_pipeline/decision_transformer.rs +git commit -m "cleanup: remove 11 dead code items (reward_norm, cea_weight, mixup, saxpy, debug_forward, DT/curiosity dead blocks)" +``` + +--- + +## Task 2: CUDA Kernels + +Write all 6 new CUDA kernels in `experience_kernels.cu`. These compile independently from the Rust integration and can be verified by checking cubin build output. + +**Files:** +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` + +- [ ] **Step 1: Add the `variable_select_bottleneck` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * Variable Selection Bottleneck (TFT-inspired per-branch feature mask). + * + * Per branch d: + * h_proj[B, R] = h_s2[B, SH2] @ W_vsn1[SH2, R] (bottleneck down) + * mask[B, SH2] = softmax(h_proj[B, R] @ W_vsn2[R, SH2]) (per-sample mask) + * h_masked[B, SH2] = h_s2[B, SH2] * mask[B, SH2] (element-wise gate) + * + * One block per sample. R=16 bottleneck. All in shared memory for small R. + * Grid: (B, 1, 1), Block: (256, 1, 1). + * + * NOTE: The two matmuls (SH2->R and R->SH2) are small enough to do in shared + * memory with thread-cooperative reduction. For production SH2=256, R=16: + * - Phase 1: 256 threads compute 16 outputs (each thread does partial dot of 16 elements) + * - Phase 2: 16 threads compute 256 outputs + softmax + * - Phase 3: 256 threads do element-wise multiply + */ +extern "C" __global__ void variable_select_bottleneck( + const float* __restrict__ h_s2, /* [B, SH2] input trunk activations */ + float* __restrict__ h_masked, /* [B, SH2] output masked activations */ + const float* __restrict__ W_vsn1, /* [SH2, R] bottleneck down weights */ + const float* __restrict__ W_vsn2, /* [R, SH2] bottleneck up weights */ + int B, int SH2, int R) +{ + int b = blockIdx.x; + if (b >= B) return; + int tid = threadIdx.x; + + extern __shared__ float smem[]; + /* smem layout: h_local[SH2] + proj[R] + mask[SH2] */ + float* h_local = smem; + float* proj = smem + SH2; + float* mask = proj + R; + + /* Load h_s2[b, :] into shared memory */ + for (int i = tid; i < SH2; i += blockDim.x) { + h_local[i] = h_s2[b * SH2 + i]; + } + __syncthreads(); + + /* Phase 1: proj[r] = sum_k h_local[k] * W_vsn1[k, r] for r in [0, R) */ + for (int r = tid; r < R; r += blockDim.x) { + float acc = 0.0f; + for (int k = 0; k < SH2; k++) { + acc += h_local[k] * W_vsn1[k * R + r]; + } + proj[r] = acc; + } + __syncthreads(); + + /* Phase 2: mask[i] = sum_r proj[r] * W_vsn2[r, i], then softmax */ + for (int i = tid; i < SH2; i += blockDim.x) { + float acc = 0.0f; + for (int r = 0; r < R; r++) { + acc += proj[r] * W_vsn2[r * SH2 + i]; + } + mask[i] = acc; + } + __syncthreads(); + + /* Softmax over mask[SH2] — find max, subtract, exp, normalize */ + /* Thread 0 computes max and sum (SH2=256 is small) */ + if (tid == 0) { + float mx = mask[0]; + for (int i = 1; i < SH2; i++) mx = fmaxf(mx, mask[i]); + float s = 0.0f; + for (int i = 0; i < SH2; i++) { + mask[i] = expf(mask[i] - mx); + s += mask[i]; + } + float inv_s = 1.0f / (s + 1e-8f); + for (int i = 0; i < SH2; i++) mask[i] *= inv_s; + } + __syncthreads(); + + /* Phase 3: h_masked[b, i] = h_local[i] * mask[i] */ + for (int i = tid; i < SH2; i += blockDim.x) { + h_masked[b * SH2 + i] = h_local[i] * mask[i]; + } +} +``` + +- [ ] **Step 2: Add the `glu_combine` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * GLU combine: output[i] = sigmoid(gate_pre[i]) * value[i] + * + * Replaces ReLU activation for branch FC layers. + * Grid: ceil(n / 256), Block: 256. + */ +extern "C" __global__ void glu_combine( + float* __restrict__ output, /* [n] output = sigmoid(gate) * value */ + const float* __restrict__ gate_pre, /* [n] pre-sigmoid gate activations */ + const float* __restrict__ value, /* [n] value path activations (W_bdf @ input + b_bdf) */ + int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n) return; + float g = 1.0f / (1.0f + expf(-gate_pre[i])); + output[i] = g * value[i]; +} +``` + +- [ ] **Step 3: Add the `glu_backward` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * GLU backward: compute d_gate_pre and d_value from d_output. + * + * Forward: out = sigmoid(gate_pre) * value + * Backward: + * d_value = d_out * sigmoid(gate_pre) + * d_gate_pre = d_out * value * sigmoid(gate_pre) * (1 - sigmoid(gate_pre)) + * + * Grid: ceil(n / 256), Block: 256. + */ +extern "C" __global__ void glu_backward( + const float* __restrict__ d_output, /* [n] upstream gradient */ + const float* __restrict__ gate_pre, /* [n] saved pre-sigmoid gate activations */ + const float* __restrict__ value, /* [n] saved value path activations */ + float* __restrict__ d_gate_pre, /* [n] gradient for gate path */ + float* __restrict__ d_value, /* [n] gradient for value path */ + int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n) return; + float sig = 1.0f / (1.0f + expf(-gate_pre[i])); + float dout = d_output[i]; + d_value[i] = dout * sig; + d_gate_pre[i] = dout * value[i] * sig * (1.0f - sig); +} +``` + +- [ ] **Step 4: Add the `cross_branch_q_attention` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * Cross-Branch Q-value Attention — 2-head self-attention on 12 Q-values. + * + * Input: Q_raw[B, 12] (4 branches x 3 actions each) + * Output: Q_coord[B, 12] = layernorm(Q_raw + W_O @ concat(head_0, head_1)) + * + * 2 heads, dim 6 each. All computation in registers (12-dim vectors are tiny). + * One thread per sample. Grid: ceil(B/256), Block: 256. + * + * Weight layout (624 floats total): + * W_Q[12, 12] (0..144) — query projection + * W_K[12, 12] (144..288) — key projection + * W_V[12, 12] (288..432) — value projection + * b_Q[12] (432..444) — query bias + * b_K[12] (444..456) — key bias + * b_V[12] (456..468) — value bias + * W_O[12, 12] (468..612) — output projection + * b_O[12] (612..624) — output bias + */ +extern "C" __global__ void cross_branch_q_attention( + const float* __restrict__ Q_raw, /* [B, 12] input Q-values */ + float* __restrict__ Q_coord, /* [B, 12] output coordinated Q-values */ + const float* __restrict__ attn_params, /* [624] weight buffer */ + int B) +{ + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= B) return; + + /* Pointers into weight buffer */ + const float* W_Q = attn_params; /* [12, 12] */ + const float* W_K = attn_params + 144; /* [12, 12] */ + const float* W_V = attn_params + 288; /* [12, 12] */ + const float* b_Q_w = attn_params + 432; /* [12] */ + const float* b_K_w = attn_params + 444; /* [12] */ + const float* b_V_w = attn_params + 456; /* [12] */ + const float* W_O = attn_params + 468; /* [12, 12] */ + const float* b_O_w = attn_params + 612; /* [12] */ + + /* Load Q_raw[b] into registers */ + float q_in[12]; + for (int i = 0; i < 12; i++) q_in[i] = Q_raw[b * 12 + i]; + + /* Project Q, K, V: out[i] = sum_j W[i,j]*q_in[j] + bias[i] */ + float Q_proj[12], K_proj[12], V_proj[12]; + for (int i = 0; i < 12; i++) { + float qsum = b_Q_w[i], ksum = b_K_w[i], vsum = b_V_w[i]; + for (int j = 0; j < 12; j++) { + qsum += W_Q[i * 12 + j] * q_in[j]; + ksum += W_K[i * 12 + j] * q_in[j]; + vsum += W_V[i * 12 + j] * q_in[j]; + } + Q_proj[i] = qsum; + K_proj[i] = ksum; + V_proj[i] = vsum; + } + + /* 2-head attention: head_dim = 6 */ + float head_out[12]; /* concat of head_0[6] and head_1[6] */ + float inv_sqrt6 = 0.40824829f; /* 1/sqrt(6) */ + + for (int h = 0; h < 2; h++) { + int off = h * 6; + /* Compute attention scores: attn[i,j] = Q[i] . K[j] / sqrt(6) */ + /* For 12 Q-values, we have 6 "tokens" per head */ + float attn[6 * 6]; /* 6x6 attention matrix */ + for (int i = 0; i < 6; i++) { + float row_max = -1e30f; + for (int j = 0; j < 6; j++) { + float dot = 0.0f; + /* Each "token" is 1-dim in this formulation, but we treat + * the 6 positions as 6 tokens with 1-dim key/query each. + * Actually: Q_proj[off+i] is a scalar, so attn = Q*K. */ + dot = Q_proj[off + i] * K_proj[off + j] * inv_sqrt6; + attn[i * 6 + j] = dot; + row_max = fmaxf(row_max, dot); + } + /* Softmax per row */ + float row_sum = 0.0f; + for (int j = 0; j < 6; j++) { + attn[i * 6 + j] = expf(attn[i * 6 + j] - row_max); + row_sum += attn[i * 6 + j]; + } + float inv_sum = 1.0f / (row_sum + 1e-8f); + for (int j = 0; j < 6; j++) attn[i * 6 + j] *= inv_sum; + + /* Weighted sum of values */ + float val = 0.0f; + for (int j = 0; j < 6; j++) { + val += attn[i * 6 + j] * V_proj[off + j]; + } + head_out[off + i] = val; + } + } + + /* Output projection: out[i] = sum_j W_O[i,j] * head_out[j] + b_O[i] */ + float proj_out[12]; + for (int i = 0; i < 12; i++) { + float s = b_O_w[i]; + for (int j = 0; j < 12; j++) s += W_O[i * 12 + j] * head_out[j]; + proj_out[i] = s; + } + + /* Residual + LayerNorm */ + float res[12]; + float mean = 0.0f; + for (int i = 0; i < 12; i++) { + res[i] = q_in[i] + proj_out[i]; + mean += res[i]; + } + mean /= 12.0f; + float var = 0.0f; + for (int i = 0; i < 12; i++) { + float d = res[i] - mean; + var += d * d; + } + var = var / 12.0f + 1e-5f; + float inv_std = rsqrtf(var); + for (int i = 0; i < 12; i++) { + Q_coord[b * 12 + i] = (res[i] - mean) * inv_std; + } +} +``` + +- [ ] **Step 5: Add the `selectivity_forward` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * Selectivity gate forward: sel[b] = sigmoid(dot(W_sel, h_s2[b]) + b_sel) + * + * One thread per sample. Grid: ceil(B/256), Block: 256. + * W_sel[SH2] + b_sel[1] = 257 params in sel_params buffer. + */ +extern "C" __global__ void selectivity_forward( + const float* __restrict__ h_s2, /* [B, SH2] trunk activations */ + float* __restrict__ sel_out, /* [B] selectivity values */ + const float* __restrict__ sel_params, /* [SH2 + 1]: W_sel[SH2] then b_sel[1] */ + int B, int SH2) +{ + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= B) return; + + float dot = sel_params[SH2]; /* b_sel */ + for (int k = 0; k < SH2; k++) { + dot += sel_params[k] * h_s2[b * SH2 + k]; + } + sel_out[b] = 1.0f / (1.0f + expf(-dot)); +} +``` + +- [ ] **Step 6: Add the `selectivity_backward` kernel** + +Append to `experience_kernels.cu`: + +```cuda +/** + * Selectivity gate backward: BCE gradient on (sel, normalized per_sample_loss). + * + * target[b] = clamp(per_sample_loss[b] / mean_loss, 0, 1) + * BCE gradient: d_sel = (sel - target) / (sel * (1-sel) + eps) [standard BCE] + * + * Actually, d/d(z) BCE(sigmoid(z), target) = sigmoid(z) - target + * where z is the pre-sigmoid logit. So we compute d_z = sel - target + * and accumulate weight gradients: + * dW[k] += d_z * h_s2[b, k] (atomicAdd across batch) + * db += d_z (atomicAdd across batch) + * + * Grid: ceil(B/256), Block: 256. + */ +extern "C" __global__ void selectivity_backward( + const float* __restrict__ h_s2, /* [B, SH2] trunk activations */ + const float* __restrict__ sel_out, /* [B] selectivity values (sigmoid output) */ + const float* __restrict__ per_sample_loss, /* [B] per-sample C51 loss */ + float* __restrict__ d_sel_params, /* [SH2 + 1] gradient accumulator */ + float mean_loss, /* scalar: mean of per_sample_loss */ + int B, int SH2) +{ + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= B) return; + + /* Normalized target: clamp(loss / mean_loss, 0, 1) */ + float inv_mean = (mean_loss > 1e-8f) ? (1.0f / mean_loss) : 1.0f; + float target = fminf(fmaxf(per_sample_loss[b] * inv_mean, 0.0f), 1.0f); + + /* BCE gradient w.r.t. pre-sigmoid logit: sel - target */ + float d_z = sel_out[b] - target; + + /* Accumulate weight gradients (atomicAdd for cross-batch accumulation) */ + for (int k = 0; k < SH2; k++) { + atomicAdd(&d_sel_params[k], d_z * h_s2[b * SH2 + k]); + } + /* Bias gradient */ + atomicAdd(&d_sel_params[SH2], d_z); +} +``` + +- [ ] **Step 7: Verify cubin builds** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +The `build.rs` should compile `experience_kernels.cu` into a cubin. If it fails, fix syntax errors in the CUDA code. The kernels only need to compile; Rust integration comes in later tasks. + +- [ ] **Step 8: Commit CUDA kernels** + +```bash +git add crates/ml/src/cuda_pipeline/experience_kernels.cu +git commit -m "feat: add 6 CUDA kernels for cross-pollinated branch intelligence (VSN, GLU, Q-attn, selectivity)" +``` + +--- + +## Task 3: Weight Layout Expansion + +Expand `NUM_WEIGHT_TENSORS` from 26 to 42, update `compute_param_sizes()`, update all array types, update `xavier_init_params_buf()`, and allocate the two separate small buffers (Q-attention: 624 params, selectivity: 257 params). + +**Files:** +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- `crates/ml/src/cuda_pipeline/batched_forward.rs` +- `crates/ml/src/cuda_pipeline/batched_backward.rs` +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` +- `crates/ml/src/cuda_pipeline/gpu_weights.rs` + +- [ ] **Step 1: Update `NUM_WEIGHT_TENSORS` constant** + +In `gpu_dqn_trainer.rs` at line 358, change: +```rust +pub(crate) const NUM_WEIGHT_TENSORS: usize = 26; +``` +to: +```rust +pub(crate) const NUM_WEIGHT_TENSORS: usize = 42; +``` + +- [ ] **Step 2: Expand `compute_param_sizes()` to include 16 new tensors** + +In `gpu_dqn_trainer.rs`, replace the `compute_param_sizes` function body. The existing 26 entries (indices 0-25) remain unchanged. Append 16 new entries (indices 26-41): + +```rust +pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT_TENSORS] { + let s1_input_dim = if cfg.bottleneck_dim > 0 { + cfg.bottleneck_dim + (cfg.state_dim.saturating_sub(cfg.market_dim)) + } else { + cfg.state_dim + }; + let bn_dim = cfg.bottleneck_dim; + let market_dim = cfg.market_dim; + let vsn_rank = 16_usize; // VSN bottleneck rank R + + [ + cfg.shared_h1 * s1_input_dim, // [0] w_s1 + cfg.shared_h1, // [1] b_s1 + cfg.shared_h2 * cfg.shared_h1, // [2] w_s2 + cfg.shared_h2, // [3] b_s2 + cfg.value_h * cfg.shared_h2, // [4] w_v1 + cfg.value_h, // [5] b_v1 + cfg.num_atoms * cfg.value_h, // [6] w_v2 + cfg.num_atoms, // [7] b_v2 + cfg.adv_h * cfg.shared_h2, // [8] w_b0fc + cfg.adv_h, // [9] b_b0fc + cfg.branch_0_size * cfg.num_atoms * cfg.adv_h, // [10] w_b0out + cfg.branch_0_size * cfg.num_atoms, // [11] b_b0out + cfg.adv_h * (cfg.shared_h2 + 3), // [12] w_b1fc (direction-conditioned) + cfg.adv_h, // [13] b_b1fc + cfg.branch_1_size * cfg.num_atoms * cfg.adv_h, // [14] w_b1out + cfg.branch_1_size * cfg.num_atoms, // [15] b_b1out + cfg.adv_h * cfg.shared_h2, // [16] w_b2fc + cfg.adv_h, // [17] b_b2fc + cfg.branch_2_size * cfg.num_atoms * cfg.adv_h, // [18] w_b2out + cfg.branch_2_size * cfg.num_atoms, // [19] b_b2out + cfg.adv_h * cfg.shared_h2, // [20] w_b3fc + cfg.adv_h, // [21] b_b3fc + cfg.branch_3_size * cfg.num_atoms * cfg.adv_h, // [22] w_b3out + cfg.branch_3_size * cfg.num_atoms, // [23] b_b3out + bn_dim * market_dim, // [24] w_bn + bn_dim, // [25] b_bn + // ── Variable Selection bottleneck (R=16) ── + cfg.shared_h2 * vsn_rank, // [26] w_vsn1_0 [SH2, R] + vsn_rank * cfg.shared_h2, // [27] w_vsn2_0 [R, SH2] + cfg.shared_h2 * vsn_rank, // [28] w_vsn1_1 [SH2, R] + vsn_rank * cfg.shared_h2, // [29] w_vsn2_1 [R, SH2] + cfg.shared_h2 * vsn_rank, // [30] w_vsn1_2 [SH2, R] + vsn_rank * cfg.shared_h2, // [31] w_vsn2_2 [R, SH2] + cfg.shared_h2 * vsn_rank, // [32] w_vsn1_3 [SH2, R] + vsn_rank * cfg.shared_h2, // [33] w_vsn2_3 [R, SH2] + // ── GLU gate weights (same shape as branch FC) ── + cfg.adv_h * cfg.shared_h2, // [34] w_gate_0 [AH, SH2] + cfg.adv_h, // [35] b_gate_0 [AH] + cfg.adv_h * (cfg.shared_h2 + 3), // [36] w_gate_1 [AH, SH2+3] (mag sees Q_dir) + cfg.adv_h, // [37] b_gate_1 [AH] + cfg.adv_h * cfg.shared_h2, // [38] w_gate_2 [AH, SH2] + cfg.adv_h, // [39] b_gate_2 [AH] + cfg.adv_h * cfg.shared_h2, // [40] w_gate_3 [AH, SH2] + cfg.adv_h, // [41] b_gate_3 [AH] + ] +} +``` + +- [ ] **Step 3: Update `xavier_init_params_buf()` with fan_dims for new tensors** + +Extend the `fan_dims` array from 26 to 42 entries. After existing `[25] b_bn` entry, append: + +```rust + (cfg.shared_h2, 16), // [26] w_vsn1_0 (Xavier) + (0, 0), // [27] w_vsn2_0 (ZERO init — uniform softmax) + (cfg.shared_h2, 16), // [28] w_vsn1_1 (Xavier) + (0, 0), // [29] w_vsn2_1 (ZERO init) + (cfg.shared_h2, 16), // [30] w_vsn1_2 (Xavier) + (0, 0), // [31] w_vsn2_2 (ZERO init) + (cfg.shared_h2, 16), // [32] w_vsn1_3 (Xavier) + (0, 0), // [33] w_vsn2_3 (ZERO init) + (0, 0), // [34] w_gate_0 (ZERO init — sigmoid=0.5) + (0, 0), // [35] b_gate_0 (ZERO init) + (0, 0), // [36] w_gate_1 (ZERO init) + (0, 0), // [37] b_gate_1 (ZERO init) + (0, 0), // [38] w_gate_2 (ZERO init) + (0, 0), // [39] b_gate_2 (ZERO init) + (0, 0), // [40] w_gate_3 (ZERO init) + (0, 0), // [41] b_gate_3 (ZERO init) +``` + +Key: VSN W_vsn2 tensors are ZERO initialized (not Xavier) so softmax starts uniform. GLU gate weights are all ZERO initialized so sigmoid starts at 0.5. + +- [ ] **Step 4: Add Q-attention and selectivity buffer fields to `GpuDqnTrainer`** + +Add these fields to the `GpuDqnTrainer` struct: + +```rust + // ── Cross-Branch Q Attention ─────────────────────────────────────── + /// Q-attention weight buffer [624] (W_QKV + W_O + biases). Separate Adam. + q_attn_params: CudaSlice, + /// Q-attention Adam first moment [624]. + q_attn_adam_m: CudaSlice, + /// Q-attention Adam second moment [624]. + q_attn_adam_v: CudaSlice, + /// Q-attention Adam step counter. + q_attn_adam_step: i32, + /// Q-attention output buffer [B, 12]. + q_coord_buf: CudaSlice, + /// Q-attention kernel function. + q_attn_kernel: CudaFunction, + + // ── Selectivity gate (Mamba-2 selective replay) ──────────────────── + /// Selectivity weight buffer [SH2 + 1] (W_sel + b_sel). Separate Adam. + sel_params: CudaSlice, + /// Selectivity Adam first moment [SH2 + 1]. + sel_adam_m: CudaSlice, + /// Selectivity Adam second moment [SH2 + 1]. + sel_adam_v: CudaSlice, + /// Selectivity gradient buffer [SH2 + 1]. + sel_grad: CudaSlice, + /// Selectivity Adam step counter. + sel_adam_step: i32, + /// Selectivity output buffer [B]. + sel_out_buf: CudaSlice, + /// Selectivity forward kernel function. + sel_fwd_kernel: CudaFunction, + /// Selectivity backward kernel function. + sel_bwd_kernel: CudaFunction, + + // ── Liquid tau per-branch dynamics ───────────────────────────────── + /// Per-branch Q-gap EMAs [4] (CPU-side state). + per_branch_q_gap_ema: [f32; 4], + /// Liquid modulation scalars — pinned device-mapped [4]. + liquid_mod_pinned: *mut f32, + /// Device pointer for liquid_mod_pinned. + liquid_mod_dev_ptr: u64, + + // ── VSN + GLU scratch buffers ────────────────────────────────────── + /// Per-branch VSN masked output [B, SH2] (one per branch, reused sequentially). + vsn_masked_buf: CudaSlice, + /// VSN kernel function. + vsn_kernel: CudaFunction, + /// Per-branch GLU gate pre-sigmoid activations [B, AH] (saved for backward). + glu_gate_pre_buf: [CudaSlice; 4], + /// Per-branch GLU value activations [B, AH] (saved for backward). + glu_value_buf: [CudaSlice; 4], + /// GLU combine kernel function. + glu_combine_kernel: CudaFunction, + /// GLU backward kernel function. + glu_backward_kernel: CudaFunction, +``` + +- [ ] **Step 5: Allocate new buffers and load kernel functions in the constructor** + +In the constructor, after loading the experience_kernels cubin module, load the 6 new kernels: + +```rust +let exp_module = context.load_cubin(EXPECTED_Q_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("experience_kernels cubin: {e}")))?; + +let vsn_kernel = exp_module.load_function("variable_select_bottleneck") + .map_err(|e| MLError::ModelError(format!("variable_select_bottleneck load: {e}")))?; +let glu_combine_kernel = exp_module.load_function("glu_combine") + .map_err(|e| MLError::ModelError(format!("glu_combine load: {e}")))?; +let glu_backward_kernel = exp_module.load_function("glu_backward") + .map_err(|e| MLError::ModelError(format!("glu_backward load: {e}")))?; +let q_attn_kernel = exp_module.load_function("cross_branch_q_attention") + .map_err(|e| MLError::ModelError(format!("cross_branch_q_attention load: {e}")))?; +let sel_fwd_kernel = exp_module.load_function("selectivity_forward") + .map_err(|e| MLError::ModelError(format!("selectivity_forward load: {e}")))?; +let sel_bwd_kernel = exp_module.load_function("selectivity_backward") + .map_err(|e| MLError::ModelError(format!("selectivity_backward load: {e}")))?; +``` + +Allocate buffers: + +```rust +// Q-attention: 624 params, zero-initialized (near-zero init) +let q_attn_params = stream.alloc_zeros::(624) + .map_err(|e| MLError::ModelError(format!("alloc q_attn_params: {e}")))?; +let q_attn_adam_m = stream.alloc_zeros::(624) + .map_err(|e| MLError::ModelError(format!("alloc q_attn_adam_m: {e}")))?; +let q_attn_adam_v = stream.alloc_zeros::(624) + .map_err(|e| MLError::ModelError(format!("alloc q_attn_adam_v: {e}")))?; +let q_coord_buf = stream.alloc_zeros::(b * 12) + .map_err(|e| MLError::ModelError(format!("alloc q_coord_buf: {e}")))?; + +// Selectivity: SH2 + 1 = 257 params, zero-initialized +let sel_param_count = config.shared_h2 + 1; +let sel_params = stream.alloc_zeros::(sel_param_count) + .map_err(|e| MLError::ModelError(format!("alloc sel_params: {e}")))?; +let sel_adam_m = stream.alloc_zeros::(sel_param_count) + .map_err(|e| MLError::ModelError(format!("alloc sel_adam_m: {e}")))?; +let sel_adam_v = stream.alloc_zeros::(sel_param_count) + .map_err(|e| MLError::ModelError(format!("alloc sel_adam_v: {e}")))?; +let sel_grad = stream.alloc_zeros::(sel_param_count) + .map_err(|e| MLError::ModelError(format!("alloc sel_grad: {e}")))?; +let sel_out_buf = stream.alloc_zeros::(b) + .map_err(|e| MLError::ModelError(format!("alloc sel_out_buf: {e}")))?; + +// VSN scratch: [B, SH2] +let vsn_masked_buf = stream.alloc_zeros::(b * config.shared_h2) + .map_err(|e| MLError::ModelError(format!("alloc vsn_masked_buf: {e}")))?; + +// GLU scratch: 4 x [B, AH] for gate_pre and value +let glu_gate_pre_buf = [ + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_gate_pre_0: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_gate_pre_1: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_gate_pre_2: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_gate_pre_3: {e}")))?, +]; +let glu_value_buf = [ + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_value_0: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_value_1: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_value_2: {e}")))?, + stream.alloc_zeros::(b * config.adv_h).map_err(|e| MLError::ModelError(format!("alloc glu_value_3: {e}")))?, +]; +``` + +Allocate liquid_mod pinned memory (same pattern as `spread_velocity_pinned`): + +```rust +// Liquid tau: 4 pinned device-mapped floats +let liquid_mod_pinned: *mut f32 = unsafe { + cudarc::driver::result::malloc_host(4 * std::mem::size_of::(), + cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP | cudarc::driver::sys::CU_MEMHOSTALLOC_PORTABLE) + .map_err(|e| MLError::ModelError(format!("pinned liquid_mod alloc: {e}")))? + as *mut f32 +}; +// Initialize to 1.0 (neutral modulation) +unsafe { + for i in 0..4 { *liquid_mod_pinned.add(i) = 1.0; } +} +let liquid_mod_dev_ptr = unsafe { + let mut dev_ptr = std::mem::MaybeUninit::uninit(); + cudarc::driver::sys::cuMemHostGetDevicePointer_v2( + dev_ptr.as_mut_ptr(), + liquid_mod_pinned.cast(), + 0, + ).result().map_err(|e| MLError::ModelError(format!("liquid_mod dev ptr: {e}")))?; + dev_ptr.assume_init() +}; +``` + +- [ ] **Step 6: Update all references to `[usize; NUM_WEIGHT_TENSORS]` arrays** + +The following files reference `NUM_WEIGHT_TENSORS` in array types and must compile with the new size of 42: + +- `batched_forward.rs`: `f32_weight_ptrs_from_base()` and all `w_ptrs: &[u64; NUM_WEIGHT_TENSORS]` parameters. These pass through existing indices 0-25 and will not touch 26-41 yet. The array is sized correctly by the constant. +- `batched_backward.rs`: Same pattern. +- `gpu_experience_collector.rs`: `param_sizes: [usize; NUM_WEIGHT_TENSORS]` field. +- `gpu_weights.rs`: `from_flat_buffer` methods. These only use indices 0-25 for existing weight sets, but the array size parameter must match. Verify they use `sizes[0..26]` only. + +No functional changes needed in these files beyond recompilation with the updated constant, but verify no hardcoded `26` appears in array initializations or loop bounds that should be `NUM_WEIGHT_TENSORS`. + +- [ ] **Step 7: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +Fix any compilation errors. The expanded arrays should be compatible with all existing code paths. + +- [ ] **Step 8: Commit weight layout expansion** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/cuda_pipeline/batched_forward.rs \ + crates/ml/src/cuda_pipeline/batched_backward.rs \ + crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \ + crates/ml/src/cuda_pipeline/gpu_weights.rs +git commit -m "feat: expand weight layout 26->42 tensors for VSN+GLU, add Q-attn/selectivity/liquid-tau buffers" +``` + +--- + +## Task 4: Variable Selection + GLU Forward Integration + +Wire VSN bottleneck and GLU gating into the forward pass, replacing ReLU activation for branch FC layers. + +**Files:** +- `crates/ml/src/cuda_pipeline/batched_forward.rs` +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` + +- [ ] **Step 1: Add VSN + GLU scratch buffer pointers to `CublasForward`** + +Add fields to `CublasGemmSet` (or `CublasForward`) for the new buffer pointers that the forward pass needs to access: + +```rust + /// VSN masked scratch [B, SH2] — output of variable_select_bottleneck. + vsn_masked_ptr: u64, + /// VSN kernel function handle. + vsn_kernel: CudaFunction, + /// GLU gate pre-sigmoid scratch [4][B, AH]. + glu_gate_pre_ptrs: [u64; 4], + /// GLU value path scratch [4][B, AH]. + glu_value_ptrs: [u64; 4], + /// GLU combine kernel function handle. + glu_combine_kernel: CudaFunction, + /// VSN bottleneck rank (R=16). + vsn_rank: usize, +``` + +Pass these from `GpuDqnTrainer` at construction or via a new method. + +- [ ] **Step 2: Modify `forward_online_raw()` branch loop to insert VSN before branch FC** + +In `batched_forward.rs`, in the branch loop within `forward_online_raw()` (both multi-stream and sequential paths), insert VSN before the branch FC GEMM. + +For each branch `d` in `0..4`: + +Before the existing branch FC GEMM, add: + +```rust +// VSN: variable_select_bottleneck(h_s2, vsn_masked, W_vsn1_d, W_vsn2_d, B, SH2, R) +let vsn1_idx = 26 + d * 2; // w_vsn1_d +let vsn2_idx = 26 + d * 2 + 1; // w_vsn2_d +let shmem_vsn = (self.shared_h2 + self.vsn_rank + self.shared_h2) * std::mem::size_of::(); +unsafe { + let current_stream = if distinct_branches { &self.branch_streams[d] } else { stream }; + current_stream + .launch_builder(&self.vsn_kernel) + .arg(&h_s2_ptr) + .arg(&self.vsn_masked_ptr) + .arg(&w_ptrs[vsn1_idx]) + .arg(&w_ptrs[vsn2_idx]) + .arg(&(b as i32)) + .arg(&(self.shared_h2 as i32)) + .arg(&(self.vsn_rank as i32)) + .launch(LaunchConfig { + grid_dim: (b as u32, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: shmem_vsn as u32, + }) + .map_err(|e| MLError::ModelError(format!("vsn branch {d}: {e}")))?; +} +// Use vsn_masked as input to branch FC (instead of h_s2) +let vsn_input = self.vsn_masked_ptr; +``` + +Then change the branch FC input from `h_s2_ptr` (or `fc_input`) to `vsn_input`. For branch 1 (magnitude), concatenate `[vsn_masked; Q_dir]` instead of `[h_s2; Q_dir]`. + +- [ ] **Step 3: Replace ReLU with GLU in branch FC** + +After the existing branch FC GEMM `W_bdf @ input + b_bdf -> branch_h`, instead of calling `launch_add_bias_relu_f32_raw`, use two parallel GEMMs + GLU combine: + +For each branch `d`: + +```rust +// Value path: GEMM already done → branch_h_ptrs[d] = W_bdf @ vsn_input +// Add bias to value path (no ReLU) +self.launch_add_bias_f32_raw(current_stream, branch_h_ptrs[d], w_ptrs[w_fc_idx + 1], self.adv_h, b)?; + +// Save value for backward +// d2d copy: glu_value_ptrs[d] <- branch_h_ptrs[d] [B * AH floats] +unsafe { + cudarc::driver::sys::cuMemcpyDtoDAsync_v2( + self.glu_value_ptrs[d], branch_h_ptrs[d], + (b * self.adv_h * std::mem::size_of::()) as u64, + current_stream.cu_stream(), + ); +} + +// Gate path: W_gate_d @ vsn_input + b_gate_d -> glu_gate_pre_ptrs[d] +let gate_w_idx = 34 + d * 2; // w_gate_d +let gate_b_idx = 34 + d * 2 + 1; // b_gate_d +let (gate_input, gate_k) = if d == 1 && mag_concat_ptr != 0 { + (mag_concat_ptr, self.mag_concat_dim) +} else { + (vsn_input, self.shared_h2) +}; +self.sgemm_f32(current_stream, w_ptrs[gate_w_idx], gate_input, + self.glu_gate_pre_ptrs[d], self.adv_h, b, gate_k, "gate_bd")?; +self.launch_add_bias_f32_raw(current_stream, self.glu_gate_pre_ptrs[d], + w_ptrs[gate_b_idx], self.adv_h, b)?; + +// GLU combine: branch_h[d] = sigmoid(gate_pre) * value +let n_glu = (b * self.adv_h) as i32; +let glu_blocks = ((n_glu as u32) + 255) / 256; +unsafe { + current_stream + .launch_builder(&self.glu_combine_kernel) + .arg(&branch_h_ptrs[d]) // output (overwrite) + .arg(&self.glu_gate_pre_ptrs[d]) // gate_pre + .arg(&self.glu_value_ptrs[d]) // value + .arg(&n_glu) + .launch(LaunchConfig { + grid_dim: (glu_blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("glu_combine branch {d}: {e}")))?; +} +``` + +Remove the `sgemm_f32_fused_relu_bias` call and `launch_add_bias_relu_f32_raw` fallback for branch FC layers. Keep fused ReLU+bias for trunk layers (s1, s2, value). + +- [ ] **Step 4: Update target forward and experience collector forward** + +The target network forward (`forward_target_raw`) and experience collector forward (`forward_online_f32`) must also apply VSN + GLU for consistent behavior. Apply the same changes to these forward paths, using the corresponding target weight pointers (indices 26-41 in target_params_buf). + +- [ ] **Step 5: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +- [ ] **Step 6: Commit VSN + GLU forward integration** + +```bash +git add crates/ml/src/cuda_pipeline/batched_forward.rs \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +git commit -m "feat: integrate VSN bottleneck + GLU gating into branch FC forward pass" +``` + +--- + +## Task 5: Variable Selection + GLU Backward Integration + +Wire VSN and GLU backward through the cuBLAS backward pass so gradients flow correctly to the 16 new weight tensors. + +**Files:** +- `crates/ml/src/cuda_pipeline/batched_backward.rs` +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` + +- [ ] **Step 1: Add GLU backward to each branch's backward path** + +In `batched_backward.rs`, in the branch backward path, after computing `d_h_bd` (gradient of branch hidden), apply GLU backward BEFORE the branch FC weight gradient GEMM: + +For each branch `d`, replace the existing ReLU mask with GLU backward: + +```rust +// GLU backward: d_h_bd is the upstream gradient into the GLU +// d_value = d_h_bd * sigmoid(gate_pre) +// d_gate_pre = d_h_bd * value * sigmoid'(gate_pre) +let n_glu = (b * adv_h) as i32; +let glu_blocks = ((n_glu as u32) + 255) / 256; +unsafe { + stream + .launch_builder(&self.glu_backward_kernel) + .arg(&d_h_bd_ptr) // d_output + .arg(&self.glu_gate_pre_ptrs[d]) // saved gate_pre + .arg(&self.glu_value_ptrs[d]) // saved value + .arg(&d_gate_pre_scratch) // output: d_gate_pre + .arg(&d_value_scratch) // output: d_value + .arg(&n_glu) + .launch(LaunchConfig { + grid_dim: (glu_blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("glu_backward branch {d}: {e}")))?; +} +``` + +Then compute weight gradients for both paths: + +```rust +// Value path weight gradient: dW_bdf += d_value^T @ vsn_input +// (same GEMM as before, but using d_value instead of d_h_bd after ReLU mask) +backward_dw_gemm(stream, d_value_scratch, vsn_input_ptr, grad_buf_w_bdf_ptr, ...)?; +backward_db_reduce(stream, d_value_scratch, grad_buf_b_bdf_ptr, ...)?; + +// Gate path weight gradient: dW_gate_d += d_gate_pre^T @ vsn_input +backward_dw_gemm(stream, d_gate_pre_scratch, vsn_input_ptr, grad_buf_w_gate_ptr, ...)?; +backward_db_reduce(stream, d_gate_pre_scratch, grad_buf_b_gate_ptr, ...)?; + +// Upstream gradient to VSN: d_vsn_input = W_bdf^T @ d_value + W_gate^T @ d_gate_pre +backward_dx_gemm(stream, w_bdf_ptr, d_value_scratch, d_vsn_input_scratch, ...)?; +// Add gate path contribution +backward_dx_gemm_accumulate(stream, w_gate_ptr, d_gate_pre_scratch, d_vsn_input_scratch, ...)?; +``` + +- [ ] **Step 2: Add VSN backward to produce gradients for W_vsn1 and W_vsn2** + +After computing `d_vsn_input` (the gradient flowing back through the branch FC into the VSN output), compute VSN weight gradients. This involves backpropagation through element-wise multiply, softmax, and two linear projections: + +```rust +// VSN backward for branch d: +// Forward was: mask = softmax(h_proj @ W_vsn2), h_masked = h_s2 * mask +// Backward: +// d_mask = d_vsn_input * h_s2 (element-wise) +// d_h_s2_from_vsn = d_vsn_input * mask (element-wise, accumulates to trunk gradient) +// d_pre_softmax = softmax_backward(d_mask, mask) +// d_W_vsn2 = h_proj^T @ d_pre_softmax +// d_h_proj = d_pre_softmax @ W_vsn2^T +// d_W_vsn1 = h_s2^T @ d_h_proj +// d_h_s2_from_proj += d_h_proj @ W_vsn1^T (accumulates to trunk gradient) + +// Implementation: These are small GEMMs ([SH2, R] and [R, SH2] with R=16). +// Use the existing sgemm infrastructure for the backward GEMMs. +// Softmax backward: d_pre_softmax[i] = mask[i] * (d_mask[i] - sum_j(mask[j] * d_mask[j])) +``` + +This is complex but follows standard neural network backward rules. Use existing `backward_dw_gemm` and `backward_dx_gemm` patterns for the GEMM operations. The softmax backward and element-wise operations can be done with a small custom kernel or as part of the existing flow. + +- [ ] **Step 3: Allocate backward scratch buffers for GLU and VSN** + +In the constructor, allocate: + +```rust +// GLU backward scratch +let d_gate_pre_scratch = stream.alloc_zeros::(b * config.adv_h) + .map_err(|e| MLError::ModelError(format!("alloc d_gate_pre_scratch: {e}")))?; +let d_value_scratch = stream.alloc_zeros::(b * config.adv_h) + .map_err(|e| MLError::ModelError(format!("alloc d_value_scratch: {e}")))?; + +// VSN backward scratch +let d_vsn_input_scratch = stream.alloc_zeros::(b * config.shared_h2) + .map_err(|e| MLError::ModelError(format!("alloc d_vsn_input: {e}")))?; +let d_vsn_mask_scratch = stream.alloc_zeros::(b * config.shared_h2) + .map_err(|e| MLError::ModelError(format!("alloc d_vsn_mask: {e}")))?; +let d_vsn_proj_scratch = stream.alloc_zeros::(b * 16) // R=16 + .map_err(|e| MLError::ModelError(format!("alloc d_vsn_proj: {e}")))?; +``` + +- [ ] **Step 4: Update the Adam graph to include indices 26-41 in gradient unflattening** + +The Adam optimizer kernel operates on the entire flat `grad_buf` which now includes tensors 26-41. Since `compute_total_params()` already sums all 42 tensors, the Adam kernel loop naturally covers the new parameters. Verify that the `graph_adam` CUDA graph is invalidated (set to `None`) when new tensors are added, forcing recapture. + +- [ ] **Step 5: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +- [ ] **Step 6: Commit VSN + GLU backward integration** + +```bash +git add crates/ml/src/cuda_pipeline/batched_backward.rs \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +git commit -m "feat: integrate VSN+GLU backward pass with gradient flow to 16 new weight tensors" +``` + +--- + +## Task 6: Cross-Branch Q Attention Integration + +Wire the `cross_branch_q_attention` kernel into the training pipeline and experience collection. Add separate Adam optimizer for the 624-param attention buffer. + +**Files:** +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` +- `crates/ml/src/trainers/dqn/fused_training.rs` + +- [ ] **Step 1: Add `launch_q_attention()` method to `GpuDqnTrainer`** + +```rust +/// Launch cross-branch Q-attention: Q_coord = attention(Q_raw). +/// +/// Reads q_out_buf [B, 12], writes q_coord_buf [B, 12]. +/// Uses q_attn_params [624] weight buffer. +/// Runs AFTER compute_expected_q, BEFORE action selection. +pub fn launch_q_attention(&mut self, batch_size: usize) -> Result<(), MLError> { + let _eg = EventTrackingGuard::new(self.stream.context()); + let b = batch_size as i32; + let blocks = ((batch_size as u32 + 255) / 256).max(1); + + unsafe { + self.stream + .launch_builder(&self.q_attn_kernel) + .arg(&self.q_out_buf) // Q_raw [B, 12] + .arg(&self.q_coord_buf) // Q_coord [B, 12] output + .arg(&self.q_attn_params) // weights [624] + .arg(&b) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("cross_branch_q_attention: {e}")))?; + } + Ok(()) +} +``` + +- [ ] **Step 2: Integrate Q-attention into experience collection** + +In `gpu_experience_collector.rs`, after `compute_expected_q` and before `experience_action_select`, add a call to the Q-attention kernel. The experience collector needs the `q_attn_params` device pointer and the `q_attn_kernel` function. Pass the Q-attention params pointer from the `GpuDqnTrainer` to the experience collector at construction or via a setter method: + +```rust +// After compute_expected_q populates q_out_buf with 12 Q-values per episode: +self.launch_q_attention(n_episodes)?; +// action_select now reads from q_coord_buf instead of q_out_buf +``` + +Update the action selection kernel to read from `q_coord_buf` (coordinated Q-values) instead of `q_out_buf` (raw Q-values) for the Boltzmann policy. + +- [ ] **Step 3: Integrate Q-attention into training forward (Expected SARSA target)** + +In the C51 loss kernel launch, the Q-values used for Expected SARSA target sampling should come from `q_coord_buf`. Either: +- Run Q-attention after `compute_expected_q` in the training forward, then pass `q_coord_buf` as the Q-value input to the C51 loss kernel. +- Or modify the C51 loss kernel to read Q_coord instead of computing Q-values internally. + +The first approach is cleaner. Add `launch_q_attention(batch_size)` call in `submit_forward_ops_main()` after the forward pass computes Q-values and before the C51 loss kernel. + +- [ ] **Step 4: Add separate Adam step for Q-attention parameters** + +Create a method `step_q_attn_adam()` that runs Adam on the 624-param buffer: + +```rust +/// Run one Adam step on the Q-attention parameters. +/// Uses the existing dqn_adam_update_kernel with the separate q_attn buffers. +pub fn step_q_attn_adam(&mut self) -> Result<(), MLError> { + self.q_attn_adam_step += 1; + let _eg = EventTrackingGuard::new(self.stream.context()); + let n = 624_i32; + let blocks = ((624_u32 + 255) / 256).max(1); + let lr = 1e-4_f32; // Lower LR for attention params + let beta1 = 0.9_f32; + let beta2 = 0.999_f32; + let eps = 1e-8_f32; + let wd = 0.0_f32; // No weight decay for attention + let max_norm = 1.0_f32; + let step = self.q_attn_adam_step; + + // Need gradient buffer for Q-attention — computed by attention backward + // For now, use the q_attn_grad buffer (to be populated by attention backward) + unsafe { + self.stream + .launch_builder(&self.adam_update_kernel) + .arg(&self.q_attn_params) + .arg(&self.q_attn_grad) // gradient (populated by backward) + .arg(&self.q_attn_adam_m) + .arg(&self.q_attn_adam_v) + .arg(&n) + .arg(&lr) + .arg(&beta1) + .arg(&beta2) + .arg(&eps) + .arg(&wd) + .arg(&max_norm) + .arg(&step) + .arg(&self.grad_norm_buf) // scalar output (reuse) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("q_attn adam: {e}")))?; + } + Ok(()) +} +``` + +Note: The Adam kernel signature may differ from the above. Match the exact parameter list of `dqn_adam_update_kernel` in `dqn_utility_kernels.cu`. + +- [ ] **Step 5: Call Q-attention Adam in the training loop** + +In `fused_training.rs` `run_full_step()`, after the main Adam step, call: + +```rust +self.trainer.step_q_attn_adam() + .map_err(|e| anyhow::anyhow!("q_attn adam: {e}"))?; +``` + +- [ ] **Step 6: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +- [ ] **Step 7: Commit Q-attention integration** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \ + crates/ml/src/trainers/dqn/fused_training.rs +git commit -m "feat: integrate cross-branch Q-attention with separate Adam optimizer" +``` + +--- + +## Task 7: Liquid Tau Dynamics + +Replace the single `spread_velocity` scalar with 4 per-branch `liquid_mod` values computed via continuous-time ODE. The c51_grad_kernel reads `liquid_mod[d]` instead of `spread_velocity[0]`. + +**Files:** +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` +- `crates/ml/src/trainers/dqn/fused_training.rs` + +- [ ] **Step 1: Modify c51_grad_kernel to read `liquid_mod[4]` instead of `spread_velocity[1]`** + +In `c51_grad_kernel.cu`, change the last parameter: + +```cuda +// OLD: +const float* __restrict__ spread_velocity) /* [1] pinned device-mapped modulator */ +// NEW: +const float* __restrict__ liquid_mod) /* [4] pinned device-mapped per-branch modulators */ +``` + +Inside the kernel, in the branch loop, replace: + +```cuda +// OLD: +float velocity_mod = spread_velocity[0]; +// NEW: +float velocity_mod = liquid_mod[d]; +``` + +- [ ] **Step 2: Update c51_grad_kernel launch in `gpu_dqn_trainer.rs`** + +In `launch_c51_grad()`, replace the `spread_velocity_dev_ptr` argument: + +```rust +// OLD: +.arg(&self.spread_velocity_dev_ptr) +// NEW: +.arg(&self.liquid_mod_dev_ptr) +``` + +- [ ] **Step 3: Remove old `spread_velocity` infrastructure** + +In `gpu_dqn_trainer.rs`: +- Remove the `spread_velocity_pinned` field (line 669). +- Remove the `spread_velocity_dev_ptr` field (line 671). +- Remove the allocation of `spread_velocity_pinned` in the constructor (lines 2588-2599). +- Remove the struct init lines for `spread_velocity_pinned` and `spread_velocity_dev_ptr` (lines 3209-3210). +- Remove the Drop impl line that frees `spread_velocity_pinned` (line 1019-1020). +- Remove the write in `update_eval_v_range` at line 1072: `unsafe { *self.spread_velocity_pinned = modulator; }` + +- [ ] **Step 4: Add `update_liquid_tau()` method to compute per-branch ODE** + +Add a new method to `GpuDqnTrainer`: + +```rust +/// Update per-branch liquid tau modulation from per-branch Q-gaps. +/// +/// Liquid ODE per branch d: +/// q_gap_d = max(Q_d) - mean(Q_d) +/// velocity_d = q_gap_d - q_gap_ema_d +/// tau_d = tau_min + (tau_max - tau_min) * sigmoid(velocity_d / delta_z) +/// liquid_mod_d += (1/tau_d) * (1.0 - liquid_mod_d) * dt +/// +/// Called from `update_eval_v_range` at epoch boundary. +pub fn update_liquid_tau(&mut self, per_branch_q_gaps: [f32; 4]) { + let tau_min = 0.01_f32; + let tau_max = 1.0_f32; + let dt = 1.0_f32; + let delta_z_approx = self.eval_q_std_ema.max(0.01); + + for d in 0..4 { + let q_gap_d = per_branch_q_gaps[d]; + + // Adaptive-rate EMA for per-branch Q-gap + if self.per_branch_q_gap_ema[d] < 1e-12 { + self.per_branch_q_gap_ema[d] = q_gap_d; + } else { + let err = (q_gap_d - self.per_branch_q_gap_ema[d]).abs(); + let alpha = (err / (err + self.per_branch_q_gap_ema[d].max(0.001))).clamp(0.01, 0.3); + self.per_branch_q_gap_ema[d] = (1.0 - alpha) * self.per_branch_q_gap_ema[d] + alpha * q_gap_d; + } + + let velocity_d = q_gap_d - self.per_branch_q_gap_ema[d]; + + // Liquid ODE: tau controls adaptation speed + let sigmoid_arg = velocity_d / delta_z_approx; + let sig = 1.0 / (1.0 + (-sigmoid_arg).exp()); + let tau_d = tau_min + (tau_max - tau_min) * sig; + + // Branch scale evolves toward 1.0 at rate 1/tau_d + let current = unsafe { *self.liquid_mod_pinned.add(d) }; + let updated = current + (1.0 / tau_d) * (1.0 - current) * dt; + unsafe { *self.liquid_mod_pinned.add(d) = updated.clamp(0.1, 2.0); } + } +} +``` + +- [ ] **Step 5: Compute per-branch Q-gaps from 12 Q-values in `reduce_current_q_stats`** + +Modify `reduce_current_q_stats` (or `update_eval_v_range`) to extract per-branch Q-gaps from the 12-element Q-value vector: + +```rust +// After downloading q_stats (q_min, q_max, q_mean, q_gap): +// Compute per-branch Q-gaps from q_out_buf [B, 12] +// Branch 0: Q[0..3], Branch 1: Q[3..6], Branch 2: Q[6..9], Branch 3: Q[9..12] +// For each branch: gap = max(Q_d) - mean(Q_d) +// This can be done on CPU from the downloaded Q-stats, or via a small GPU kernel. +// For simplicity, compute during q_stats readback. +``` + +Add a GPU kernel or CPU-side computation that extracts per-branch Q-gaps. The simplest approach is to read back a small buffer and compute on CPU, since this only runs every 50 training steps. + +- [ ] **Step 6: Call `update_liquid_tau` from `update_eval_v_range`** + +Modify `update_eval_v_range` to call `update_liquid_tau` with the per-branch Q-gaps: + +```rust +pub fn update_eval_v_range(&mut self, q_mean: f32, q_std: f32, q_gap: f32, per_branch_q_gaps: [f32; 4]) { + // ... existing EMA and v_range code ... + + // Replace old spread_velocity write with liquid tau update + self.update_liquid_tau(per_branch_q_gaps); +} +``` + +Update all callers of `update_eval_v_range` to pass the per-branch Q-gaps. + +- [ ] **Step 7: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +- [ ] **Step 8: Commit liquid tau dynamics** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/cuda_pipeline/c51_grad_kernel.cu \ + crates/ml/src/trainers/dqn/fused_training.rs +git commit -m "feat: per-branch liquid tau ODE replaces global spread_velocity" +``` + +--- + +## Task 8: Mamba-2 Selective Replay + +Wire the selectivity gate into the training loop. The gate learns which trunk activations predict high-loss transitions and scales PER priorities accordingly. + +**Files:** +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- `crates/ml/src/trainers/dqn/fused_training.rs` +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` + +- [ ] **Step 1: Add `launch_selectivity_forward()` method** + +```rust +/// Selectivity gate forward: sel[b] = sigmoid(dot(W_sel, h_s2[b]) + b_sel) +/// +/// Reads save_h_s2 [B, SH2] (populated by training forward pass). +/// Writes sel_out_buf [B]. +/// Runs AFTER the training forward pass graph, BEFORE PER priority update. +pub fn launch_selectivity_forward(&mut self, batch_size: usize) -> Result<(), MLError> { + let _eg = EventTrackingGuard::new(self.stream.context()); + let b = batch_size as i32; + let sh2 = self.config.shared_h2 as i32; + let blocks = ((batch_size as u32 + 255) / 256).max(1); + + unsafe { + self.stream + .launch_builder(&self.sel_fwd_kernel) + .arg(&self.save_h_s2) // h_s2 [B, SH2] + .arg(&self.sel_out_buf) // sel [B] output + .arg(&self.sel_params) // W_sel + b_sel [SH2+1] + .arg(&b) + .arg(&sh2) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("selectivity_forward: {e}")))?; + } + Ok(()) +} +``` + +- [ ] **Step 2: Add `launch_selectivity_backward()` method** + +```rust +/// Selectivity gate backward: BCE gradient on (sel, normalized per_sample_loss). +/// +/// Reads save_h_s2, sel_out_buf, per_sample_loss_buf. +/// Writes sel_grad [SH2+1]. +/// Must be called AFTER C51 loss (per_sample_loss_buf populated). +pub fn launch_selectivity_backward(&mut self, batch_size: usize) -> Result<(), MLError> { + let _eg = EventTrackingGuard::new(self.stream.context()); + + // Compute mean loss on CPU from per_sample_loss_buf (sync required) + // Alternative: use a small GPU reduction. For now, use the total_loss + // scalar (already computed) divided by batch_size. + let mean_loss = self.last_total_loss / batch_size as f32; + + // Zero gradient accumulator + let sel_param_count = self.config.shared_h2 + 1; + self.stream.memset_zeros(&mut self.sel_grad) + .map_err(|e| MLError::ModelError(format!("zero sel_grad: {e}")))?; + + let b = batch_size as i32; + let sh2 = self.config.shared_h2 as i32; + let blocks = ((batch_size as u32 + 255) / 256).max(1); + + unsafe { + self.stream + .launch_builder(&self.sel_bwd_kernel) + .arg(&self.save_h_s2) + .arg(&self.sel_out_buf) + .arg(&self.per_sample_loss_buf) + .arg(&self.sel_grad) + .arg(&mean_loss) + .arg(&b) + .arg(&sh2) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("selectivity_backward: {e}")))?; + } + Ok(()) +} +``` + +- [ ] **Step 3: Add `step_selectivity_adam()` method** + +```rust +/// Run one Adam step on the selectivity gate parameters. +pub fn step_selectivity_adam(&mut self) -> Result<(), MLError> { + self.sel_adam_step += 1; + let _eg = EventTrackingGuard::new(self.stream.context()); + let n = (self.config.shared_h2 + 1) as i32; + let blocks = ((n as u32 + 255) / 256).max(1); + let lr = 1e-4_f32; + let beta1 = 0.9_f32; + let beta2 = 0.999_f32; + let eps = 1e-8_f32; + let wd = 0.0_f32; + let max_norm = 1.0_f32; + let step = self.sel_adam_step; + + unsafe { + self.stream + .launch_builder(&self.adam_update_kernel) + .arg(&self.sel_params) + .arg(&self.sel_grad) + .arg(&self.sel_adam_m) + .arg(&self.sel_adam_v) + .arg(&n) + .arg(&lr) + .arg(&beta1) + .arg(&beta2) + .arg(&eps) + .arg(&wd) + .arg(&max_norm) + .arg(&step) + .arg(&self.aux_norm_scratch) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("selectivity adam: {e}")))?; + } + Ok(()) +} +``` + +- [ ] **Step 4: Add `read_selectivity_for_per()` method** + +```rust +/// Read selectivity values and return for PER priority scaling. +/// +/// Downloads sel_out_buf [B] from GPU. Each value in [0, 1]. +/// Priority scaling: new_priority[i] = old_priority[i] * (1 + sel[i]) +pub fn read_selectivity_for_per(&self, batch_size: usize) -> Result, MLError> { + // Sync stream before readback + unsafe { cudarc::driver::sys::cuStreamSynchronize(self.stream.cu_stream()); } + let mut host_buf = vec![0.0_f32; batch_size]; + unsafe { + cudarc::driver::sys::cuMemcpyDtoH_v2( + host_buf.as_mut_ptr().cast(), + self.sel_out_buf.raw_ptr(), + (batch_size * std::mem::size_of::()) as u64, + ); + } + Ok(host_buf) +} +``` + +- [ ] **Step 5: Integrate selectivity into `run_full_step` in `fused_training.rs`** + +After the main training graph replay and Adam step, add: + +```rust +// ── Selectivity gate: state-conditioned PER priority scaling ── +self.trainer.launch_selectivity_forward(self.batch_size) + .map_err(|e| anyhow::anyhow!("selectivity fwd: {e}"))?; +self.trainer.launch_selectivity_backward(self.batch_size) + .map_err(|e| anyhow::anyhow!("selectivity bwd: {e}"))?; +self.trainer.step_selectivity_adam() + .map_err(|e| anyhow::anyhow!("selectivity adam: {e}"))?; +``` + +- [ ] **Step 6: Integrate selectivity into PER priority update** + +In the PER priority update section of the training loop (where `td_errors` are used to update priorities), multiply by selectivity: + +```rust +// After reading td_errors for PER update: +let sel_values = self.trainer.read_selectivity_for_per(self.batch_size) + .map_err(|e| anyhow::anyhow!("read selectivity: {e}"))?; + +// Scale priorities: priority[i] = |td_error[i]| * (1 + sel[i]) + epsilon +for i in 0..td_errors.len() { + td_errors[i] = td_errors[i].abs() * (1.0 + sel_values[i]); +} +``` + +Locate the exact PER priority update code in `fused_training.rs` or `training_loop.rs` and modify the priority computation. + +- [ ] **Step 7: Verify clean build** + +```bash +SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -20 +``` + +- [ ] **Step 8: Commit selective replay** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/trainers/dqn/fused_training.rs \ + crates/ml/src/trainers/dqn/trainer/training_loop.rs +git commit -m "feat: Mamba-2 selective replay — learned state-conditioned PER priority gate" +``` + +--- + +## Task 9: Build Verification + Smoke Tests + +Full build + smoke test to confirm everything compiles and runs without crashing. + +**Files:** +- No new files. Verification only. + +- [ ] **Step 1: Full workspace check** + +```bash +SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -30 +``` + +Fix any compilation errors across the workspace. + +- [ ] **Step 2: Run ML crate unit tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | tail -30 +``` + +Fix any test failures. Tests that reference old APIs (e.g., `mixup_alpha`, `spread_velocity`, `cea_weight`) must be updated to use the new APIs. + +- [ ] **Step 3: Run smoke tests (requires test data)** + +```bash +SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -50 +``` + +The smoke tests exercise the full training pipeline on a small dataset. If they pass, the integration is functionally correct. + +- [ ] **Step 4: Verify parameter count logging** + +Check that the training startup log reports the expanded parameter count. With SH2=256, AH=128, R=16: +- VSN: 8 x (256 x 16) = 32,768 +- GLU: 4 x (128 x 256 + 128) + (128 x 259 + 128) = ~263,168 (branch 1 wider) +- Total new in params_buf: ~296,000 additional parameters + +The log should show `total_params` increased by approximately this amount. + +- [ ] **Step 5: Final commit with verification note** + +```bash +git add -A +git status +# Only commit if there are meaningful changes (test fixes, etc.) +git commit -m "fix: resolve compilation and test issues from cross-pollinated branch intelligence integration" +```