diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 408e5ec74..c2baba98e 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -767,6 +767,25 @@ fn main() { // producer chain has fresh labels every rollout step. Same per-thread // O(1) map; same lookahead window semantics; same skip sentinel. "aux_sign_label_per_step_kernel.cu", + // SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction + // horizon producer. Single-thread kernel writing + // `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from + // `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]`. Pearl-A first-observation + // bootstrap (sentinel 60.0 → direct replacement); Wiener-α slow + // EMA blend thereafter (α=0.01 fixed; no target-variance EMA + // available in C.4b). Per-epoch boundary launch — H is + // slow-moving. Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.4b. + "aux_horizon_update_kernel.cu", + // SP14 Layer C Phase C.4b (2026-05-08): avg winning hold time + // EMA producer. Single-block kernel that sweeps the per-epoch + // `hold_at_exit_per_sample` + `trade_profitable_per_sample` + // buffers (populated by `unified_env_step_core` in + // `experience_kernels.cu`) and writes the EMA-blended mean to + // `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` for downstream + // consumption by `aux_horizon_update_kernel`. Block-tree-reduce + // (no atomicAdd) per `feedback_no_atomicadd.md`. Pearl-A + // first-observation bootstrap; α=0.05 EMA thereafter. Plan: §C.4b. + "avg_win_hold_time_update_kernel.cu", // SP14 Layer B Task B.3 (2026-05-05): Earned Gradient Flow // producer kernel. Per-step computes the K=4↔K=2 mapped argmax // mismatch between the Q-head's 4-way direction action diff --git a/crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu b/crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu new file mode 100644 index 000000000..ee27492da --- /dev/null +++ b/crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu @@ -0,0 +1,118 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP14 Layer C Phase C.4b — adaptive aux prediction horizon producer + * (2026-05-08). + * + * Drives `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` (the H consumed by the + * aux next-bar sign label kernels) from observed avg winning hold time + * `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` (produced upstream by + * `avg_win_hold_time_update_kernel`). + * + * Algorithm (per `pearl_first_observation_bootstrap.md` + + * `pearl_wiener_optimal_adaptive_alpha.md`): + * + * 1. If `h_target <= 0` → no winning trade observed yet. Keep + * sentinel H=60.0 (≈1hr @ 1-min bars). This is the cold-start guard + * per `pearl_cold_start_exit_signal_or.md` — without it, a sentinel + * `h_target == 0.0` would clamp to h_min=1 on the bootstrap branch. + * + * 2. If `|h_current - sentinel_h| < 1e-6f` → first valid observation. + * REPLACE H with `clamp(h_target, h_min, h_max)` directly (Pearl-A + * first-observation bootstrap). NO blend — the sentinel must not + * contaminate the EMA. + * + * 3. Otherwise → steady-state Wiener-α EMA blend. Since no variance + * EMA exists for the target signal, falls back to fixed slow blend + * α=0.01 (per `pearl_wiener_optimal_adaptive_alpha.md` fallback). + * H is slow-moving (avg winning hold time changes on epoch scale, + * not step scale), so a slow EMA is appropriate — would otherwise + * track noise. + * + * Pearls + invariants: + * + * - `pearl_no_host_branches_in_captured_graph.md` — single-thread + * kernel (one block, one thread). The launcher in `gpu_aux_trunk.rs` + * pre-loads the `CudaFunction` once at construction; no runtime + * branch on host. Per-epoch boundary launch (horizon is slow-moving; + * per-step would be wasteful). + * + * - `feedback_no_atomicadd.md` — single-thread compute, no atomics. + * + * - `feedback_isv_for_adaptive_bounds.md` — `h_min=1, h_max=240` are + * fundamental floors/ceilings (cannot predict the past; rollout + * buffer guard), NOT tuning parameters. Acceptable as constants. + * + * - `feedback_no_stubs.md` — full body, no placeholder return. + * + * Args: + * isv — `[ISV_TOTAL_DIM]` mapped-pinned f32, modified + * in place at index `isv_h_idx`. + * isv_h_idx — AUX_PRED_HORIZON_BARS_INDEX (450). + * isv_h_target_idx — AVG_WIN_HOLD_TIME_BARS_INDEX (451) — read-only + * source slot. + * isv_h_target_var_idx — variance EMA slot for the target signal, or + * `-1` if no variance EMA exists (uses fixed + * α=0.01). Currently `-1` — Phase C.4b does + * not introduce a target-variance EMA. + * h_min, h_max — fundamental bounds (1.0, 240.0). + * sentinel_h — sentinel value (60.0) used to detect first + * observation. Must match + * `SENTINEL_AUX_PRED_HORIZON_BARS` in + * `sp14_isv_slots.rs`. + * ══════════════════════════════════════════════════════════════════════════ */ +extern "C" __global__ +void aux_horizon_update( + float* __restrict__ isv, + int isv_h_idx, + int isv_h_target_idx, + int isv_h_target_var_idx, /* -1 if not available */ + float h_min, + float h_max, + float sentinel_h) +{ + /* Single-thread kernel: only thread (0, 0) does the compute. The + * extra threads early-return so the launch dim can stay (1, 1, 1) + * without branching on the host side. */ + if (blockIdx.x != 0 || threadIdx.x != 0) return; + + float h_target = isv[isv_h_target_idx]; + float h_current = isv[isv_h_idx]; + + /* Guard 1: no winning trade observed yet. Keep sentinel. */ + if (h_target <= 0.0f) { + isv[isv_h_idx] = sentinel_h; + return; + } + + /* Guard 2: Pearl-A first-observation bootstrap. The current value is + * exactly the sentinel ⇒ this is the first valid h_target observation; + * REPLACE directly (no blend) per + * `pearl_first_observation_bootstrap.md`. */ + if (fabsf(h_current - sentinel_h) < 1e-6f) { + float h_clamped = fmaxf(h_min, fminf(h_max, h_target)); + isv[isv_h_idx] = h_clamped; + return; + } + + /* Steady-state EMA blend. */ + float alpha; + if (isv_h_target_var_idx >= 0) { + /* Wiener-α: α = sample_var / (sample_var + diff_var + ε). + * Bounded into [0.001, 0.1] so a single noisy sample can never + * fully overwrite the EMA, and a perfectly steady target can + * still adapt. Currently unreachable — Phase C.4b passes -1. */ + float sample_var = isv[isv_h_target_var_idx]; + float diff = h_target - h_current; + float diff_var = diff * diff; /* single-sample approximation */ + alpha = sample_var / (sample_var + diff_var + 1e-6f); + alpha = fmaxf(0.001f, fminf(0.1f, alpha)); + } else { + /* No variance EMA available — fixed slow blend. The horizon + * tracks avg winning hold time which itself is an EMA, so a + * second-order slow blend (α=0.01) suffices to filter + * remaining sample-to-sample noise. */ + alpha = 0.01f; + } + + float h_blended = (1.0f - alpha) * h_current + alpha * h_target; + isv[isv_h_idx] = fmaxf(h_min, fminf(h_max, h_blended)); +} diff --git a/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu b/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu index ff16cfdfc..47b7d2b4e 100644 --- a/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu @@ -1,6 +1,16 @@ /* ══════════════════════════════════════════════════════════════════════════ * SP13 Layer B — Commit B1.1b producer kernel for aux next-bar sign label - * (2026-05-05). + * (2026-05-05). SP14 Layer C Phase C.4b (2026-05-08): migrated from + * fixed-`lookahead` HFT-scale label `(p_{t+1} > p_t)` to ISV-driven + * multi-bar horizon `(p_{t+H} > p_t)` where H is read from + * `ISV[isv_h_idx = AUX_PRED_HORIZON_BARS_INDEX]` (slot 450). + * + * Rationale: aux's original next-bar prediction is HFT-scale + * microstructure noise — unlearnable at our HFT-MFT trading frequency + * (multi-bar holds). With H read from ISV, the adaptive producer + * (`aux_horizon_update_kernel`) drives H from observed avg winning hold + * time so the label semantic matches actual position lifetime. See + * `2026-05-07-sp14-layer-c-separate-aux-trunk.md` §C.4b for the design. * * Replaces the B0 `alloc_zeros::(total)` placeholder for * `aux_sign_labels` in `GpuExperienceCollector::collect_experiences_gpu`. @@ -13,17 +23,21 @@ * Per-thread O(1) map: read `targets[bar*6+2]` (raw_close column — see * `dt_kernels.cu:803` and `scripted_policy_kernel.cu:56` for the canonical * column layout: 0=preproc_close, 1=preproc_next, 2=raw_close, 3=raw_next, - * 4=raw_open, 5=mid_open) at `bar` and `bar+lookahead`, classify by + * 4=raw_open, 5=mid_open) at `bar` and `bar+H`, classify by * strict greater-than — flat/down both map to label 0 because the K=2 * softmax can only distinguish "up" from "not-up". The skip sentinel -1 * fires when the lookahead window would run past `total_bars`; the CE * loss reduce kernel masks it (`label < 0` → row excluded from * mean + B_valid count per `aux_heads_kernel.cu`). * + * H clamping: `H = clamp(round(isv[isv_h_idx]), 1, 240)` — fundamental + * floor (cannot predict the past) and ceiling (rollout buffer guard). + * Clamping is on-device per `feedback_isv_for_adaptive_bounds.md`. + * * Pure map; no atomicAdd, no reduction, no shared memory — block * tree-reduce N/A here per `feedback_no_atomicadd.md`. GPU-only data * read/write per `feedback_cpu_is_read_only.md`. Inputs are mapped-pinned - * (targets) or device-allocated (bar_indices, out_labels) per + * (targets, isv) or device-allocated (bar_indices, out_labels) per * `feedback_no_htod_htoh_only_mapped_pinned.md`. * * Launch config: grid=(ceil(total / 256)), block=(256). Mirrors @@ -38,36 +52,43 @@ * experience (constructed by host in * `collect_experiences_gpu`'s pre-existing * `bar_indices_pinned` fill loop). + * isv — `[ISV_TOTAL_DIM]` mapped-pinned f32. The kernel reads + * `isv[isv_h_idx]` once per thread (broadcast value). + * isv_h_idx — index of the adaptive horizon slot + * (AUX_PRED_HORIZON_BARS_INDEX = 450). * out_labels — `[total]` device i32; receives -1/0/1 per experience. * total — number of experiences. * total_bars — total bars in the data series (for skip-sentinel - * bound check: bar+lookahead >= total_bars ⇒ -1). - * lookahead — number of bars to look ahead (structural design - * constant, value passed in by the host launcher; - * must match the trainer's expectation since the K=2 - * head's "up" label semantics encode the same - * lookahead window). + * bound check: bar+H >= total_bars ⇒ -1). * ══════════════════════════════════════════════════════════════════════════ */ extern "C" __global__ void aux_sign_label_kernel( const float* __restrict__ targets, const int* __restrict__ bar_indices, + const float* __restrict__ isv, + int isv_h_idx, int* __restrict__ out_labels, int total, - int total_bars, - int lookahead) + int total_bars) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid >= total) return; + /* Read H from ISV (broadcast value — every thread reads the same + * cell). Round-to-nearest then clamp into [1, 240] per + * `feedback_isv_for_adaptive_bounds.md`. */ + int H = (int)__float2int_rn(isv[isv_h_idx]); + if (H < 1) H = 1; + if (H > 240) H = 240; + int bar = bar_indices[tid]; - if (bar < 0 || bar + lookahead >= total_bars) { + if (bar < 0 || bar + H >= total_bars) { out_labels[tid] = -1; // skip sentinel — no lookahead window available return; } float p_now = targets[bar * 6 + 2]; - float p_fut = targets[(bar + lookahead) * 6 + 2]; + float p_fut = targets[(bar + H) * 6 + 2]; // Strict greater-than tie-break: ties (p_fut == p_now, e.g. flat) // map to "down"=0 because the K=2 head can only distinguish "up" diff --git a/crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu b/crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu index 430641db6..2e90876f9 100644 --- a/crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu @@ -1,6 +1,11 @@ /* ══════════════════════════════════════════════════════════════════════════ * SP14 β-migration step 3 — per-rollout-step variant of aux_sign_label_kernel - * (2026-05-07). + * (2026-05-07). SP14 Layer C Phase C.4b (2026-05-08): migrated to + * ISV-driven multi-bar horizon `(p_{t+H} > p_t)` where H is read from + * `ISV[isv_h_idx = AUX_PRED_HORIZON_BARS_INDEX]` (slot 450). Migration + * MUST happen atomically with the trajectory kernel + * (`aux_sign_label_kernel.cu`) per `feedback_no_partial_refactor.md` — + * if only one is migrated, training-time and eval-time labels diverge. * * The full collector-end `aux_sign_label_kernel` (this directory) writes * the entire trajectory's labels at the end of `collect_experiences_gpu` @@ -12,14 +17,17 @@ * Same per-thread O(1) map as the trajectory kernel: read `targets[bar*6+2]` * (raw_close column — see `dt_kernels.cu:803` for the canonical column * layout: 0=preproc_close, 1=preproc_next, 2=raw_close, 3=raw_next, - * 4=raw_open, 5=mid_open) at `bar` and `bar+lookahead`, classify by + * 4=raw_open, 5=mid_open) at `bar` and `bar+H`, classify by * strict greater-than. Skip sentinel -1 fires when the lookahead window * runs past `total_bars`. * + * H clamping: `H = clamp(round(isv[isv_h_idx]), 1, 240)` — same as + * trajectory variant per `feedback_isv_for_adaptive_bounds.md`. + * * Pure map; no atomicAdd, no reduction, no shared memory — block * tree-reduce N/A here per `feedback_no_atomicadd.md`. GPU-only data * read/write per `feedback_cpu_is_read_only.md`. Inputs are mapped-pinned - * (targets) or device-allocated (episode_starts, out_labels) per + * (targets, isv) or device-allocated (episode_starts, out_labels) per * `feedback_no_htod_htoh_only_mapped_pinned.md`. * * Launch config: grid=(ceil(n_episodes / 256)), block=(256). Mirrors the @@ -29,34 +37,39 @@ * targets — `[total_bars, 6]` mapped-pinned f32. * episode_starts — `[n_episodes]` device i32, base bar index per env. * t — current rollout step (0..timesteps). + * isv — `[ISV_TOTAL_DIM]` mapped-pinned f32. + * isv_h_idx — AUX_PRED_HORIZON_BARS_INDEX (450). * out_labels — `[n_episodes]` device i32; receives -1/0/1 per env. * n_episodes — number of envs in this rollout slice. * total_bars — total bars in the data series. - * lookahead — bars to look ahead (must match the trajectory - * kernel's lookahead — both encode the same K=2 - * "up" boundary). * ══════════════════════════════════════════════════════════════════════════ */ extern "C" __global__ void aux_sign_label_per_step_kernel( const float* __restrict__ targets, const int* __restrict__ episode_starts, int t, + const float* __restrict__ isv, + int isv_h_idx, int* __restrict__ out_labels, int n_episodes, - int total_bars, - int lookahead) + int total_bars) { int ep = blockIdx.x * blockDim.x + threadIdx.x; if (ep >= n_episodes) return; + /* Read H from ISV (broadcast value); clamp [1, 240]. */ + int H = (int)__float2int_rn(isv[isv_h_idx]); + if (H < 1) H = 1; + if (H > 240) H = 240; + int bar = episode_starts[ep] + t; - if (bar < 0 || bar + lookahead >= total_bars) { + if (bar < 0 || bar + H >= total_bars) { out_labels[ep] = -1; // skip sentinel — no lookahead window available return; } float p_now = targets[bar * 6 + 2]; - float p_fut = targets[(bar + lookahead) * 6 + 2]; + float p_fut = targets[(bar + H) * 6 + 2]; // Strict greater-than tie-break: matches trajectory kernel — flat/down // both map to "down"=0 because the K=2 head can only distinguish diff --git a/crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu b/crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu new file mode 100644 index 000000000..1ae2bbc9c --- /dev/null +++ b/crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu @@ -0,0 +1,137 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP14 Layer C Phase C.4b — avg winning hold time EMA producer (2026-05-08). + * + * Aggregates per-sample trade-close events from `experience_kernels.cu`: + * + * - `hold_at_exit_per_sample[N*L]` — `saved_hold_time` on + * exit/reverse, else 0.0 + * - `trade_profitable_per_sample[N*L]` — 1 iff (trade_close && step_ret>0), + * else 0 + * + * Computes mean(hold_at_exit | profitable) over the per-epoch sample + * buffer, then EMA-blends into `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` + * for downstream consumption by `aux_horizon_update_kernel`. + * + * Algorithm (Pearl-A first-observation bootstrap + + * `pearl_wiener_optimal_adaptive_alpha.md`): + * + * Phase 1 (block-tree-reduce): each block of `BLK_DIM` threads sweeps + * its tile of N=B*T samples, summing `hold_at_exit_per_sample[i]` and + * counting `trade_profitable_per_sample[i]==1`. Local block reduce in + * shared memory; block-0-thread-0 finalises after grid sync (single + * block — see launch config below). + * + * Phase 2 (Pearl-A + EMA): if sample mean is invalid (count==0), keep + * ISV slot unchanged. If current ISV value is sentinel 0.0 (or near + * 0.0), REPLACE with sample mean (Pearl-A first-observation bootstrap + * per `pearl_first_observation_bootstrap.md`). Otherwise EMA-blend + * with α=0.05 (slow blend — avg winning hold time is a per-epoch + * aggregate that varies on policy-evolution timescales, faster than + * the horizon EMA but still slow enough to filter + * sample-distribution noise). + * + * Pearls + invariants: + * + * - `feedback_no_atomicadd.md` — single-block kernel (the per-epoch + * sample buffer is ≤ N*L≈64k slots; one block of 256 threads sweeps + * it via stride-based for-loop, block-tree-reduce in shmem, no + * atomics across blocks). + * + * - `pearl_no_host_branches_in_captured_graph.md` — pre-loaded + * `CudaFunction` at construction; no per-launch host branches. The + * guard for `count == 0` is on-device. + * + * - `pearl_first_observation_bootstrap.md` — first valid sample + * replaces sentinel directly; no blend. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * Args: + * hold_at_exit_per_sample — `[total_samples]` device f32. + * trade_profitable_per_sample — `[total_samples]` device i32. + * total_samples — N*L (B*T total samples in the epoch). + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at index `isv_target_idx`. + * isv_target_idx — AVG_WIN_HOLD_TIME_BARS_INDEX (451). + * sentinel — SENTINEL_AVG_WIN_HOLD_TIME_BARS (0.0). + * alpha — EMA blend rate (0.05 — slow). + * + * Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1). + * Shared memory: 2 * BLK_DIM floats (sum_hold + count_profitable). The + * count is summed as float for unified shmem layout — values are + * bounded by total_samples ≤ 2^16 so f32 representation is exact for + * counts up to 2^24. + * ══════════════════════════════════════════════════════════════════════════ */ + +#define BLK_DIM 256 + +extern "C" __global__ +void avg_win_hold_time_update( + const float* __restrict__ hold_at_exit_per_sample, + const int* __restrict__ trade_profitable_per_sample, + int total_samples, + float* __restrict__ isv, + int isv_target_idx, + float sentinel, + float alpha) +{ + /* Single-block kernel — only block 0 does anything. */ + if (blockIdx.x != 0) return; + + extern __shared__ float shmem[]; + float* sh_sum = shmem; /* [BLK_DIM] */ + float* sh_count = shmem + BLK_DIM; /* [BLK_DIM] */ + + int tid = threadIdx.x; + + /* Phase 1a: stride-based sweep — accumulate sum + count locally. */ + float local_sum = 0.0f; + float local_count = 0.0f; + for (int i = tid; i < total_samples; i += BLK_DIM) { + int profit = trade_profitable_per_sample[i]; /* 0 or 1 */ + float hold = hold_at_exit_per_sample[i]; /* 0 if not exit/reverse */ + if (profit != 0) { + local_sum += hold; + local_count += 1.0f; + } + } + sh_sum[tid] = local_sum; + sh_count[tid] = local_count; + __syncthreads(); + + /* Phase 1b: block-tree-reduce (no atomicAdd) — halve every iter. */ + for (int s = BLK_DIM >> 1; s > 0; s >>= 1) { + if (tid < s) { + sh_sum[tid] += sh_sum[tid + s]; + sh_count[tid] += sh_count[tid + s]; + } + __syncthreads(); + } + + /* Phase 2: thread 0 finalises EMA. */ + if (tid != 0) return; + + float total_sum = sh_sum[0]; + float total_count = sh_count[0]; + + /* Guard: no winning trades observed this epoch — keep ISV unchanged. */ + if (total_count <= 0.0f) { + return; + } + + float sample_mean = total_sum / total_count; + float current = isv[isv_target_idx]; + + /* Pearl-A first-observation bootstrap. The sentinel is exactly 0.0; + * any value within 1e-6 is treated as sentinel. (`current <= 0.0f` + * is a stronger condition that also handles any out-of-spec + * negative writes.) */ + if (current <= 1e-6f) { + isv[isv_target_idx] = sample_mean; + return; + } + + /* Steady-state EMA blend. */ + float blended = (1.0f - alpha) * current + alpha * sample_mean; + isv[isv_target_idx] = blended; +} diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index e8b5459e3..1a7854ae5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -50,7 +50,10 @@ use cudarc::driver::{CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; use crate::MLError; -use super::gpu_dqn_trainer::{AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN}; +use super::gpu_dqn_trainer::{ + AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN, + AVG_WIN_HOLD_TIME_UPDATE_CUBIN, +}; /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU /// output). Mirrors the encoder output dimension so Layer-1 acts as a @@ -450,3 +453,167 @@ impl AuxTrunkBackwardOps { Ok(()) } } + +/// SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction horizon +/// producer. +/// +/// Drives `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from +/// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` via Pearl-A first-observation +/// bootstrap + slow Wiener-α EMA. Single-thread, single-block kernel — +/// horizon is slow-moving so per-epoch boundary launch is appropriate +/// (per-step would be wasteful and would track sample noise). +/// +/// # Pearls applied +/// - `pearl_first_observation_bootstrap.md` — sentinel = 60.0; first +/// valid observation REPLACES (no blend) so the cold-start sentinel +/// never contaminates the EMA. +/// - `pearl_wiener_optimal_adaptive_alpha.md` — fixed α=0.01 fallback +/// when no target-variance EMA exists. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded once at construction. +/// - `feedback_isv_for_adaptive_bounds.md` — h_min=1, h_max=240 are +/// fundamental floors/ceilings, not tuning parameters. +#[allow(missing_debug_implementations)] +pub(crate) struct AuxHorizonUpdateOps { + update_kernel: CudaFunction, +} + +impl AuxHorizonUpdateOps { + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(AUX_HORIZON_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("aux_horizon_update cubin load: {e}")))?; + let update_kernel = module + .load_function("aux_horizon_update") + .map_err(|e| MLError::ModelError(format!("aux_horizon_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the horizon producer. + /// + /// Args: + /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). + /// - `isv_h_idx`: AUX_PRED_HORIZON_BARS_INDEX (450). + /// - `isv_h_target_idx`: AVG_WIN_HOLD_TIME_BARS_INDEX (451). + /// - `isv_h_target_var_idx`: variance EMA slot for target, or `-1` + /// if none (uses fixed α=0.01). + /// - `h_min`, `h_max`: fundamental bounds (1.0, 240.0). + /// - `sentinel_h`: SENTINEL_AUX_PRED_HORIZON_BARS (60.0). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + isv_ptr: u64, + isv_h_idx: i32, + isv_h_target_idx: i32, + isv_h_target_var_idx: i32, + h_min: f32, + h_max: f32, + sentinel_h: f32, + ) -> Result<(), MLError> { + unsafe { + stream + .launch_builder(&self.update_kernel) + .arg(&isv_ptr) + .arg(&isv_h_idx) + .arg(&isv_h_target_idx) + .arg(&isv_h_target_var_idx) + .arg(&h_min) + .arg(&h_max) + .arg(&sentinel_h) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("aux_horizon_update: {e}")))?; + } + Ok(()) + } +} + +/// SP14 Layer C Phase C.4b (2026-05-08): avg winning hold time EMA producer. +/// +/// Sweeps the per-epoch `hold_at_exit_per_sample` + +/// `trade_profitable_per_sample` buffers (populated by +/// `unified_env_step_core` in `experience_kernels.cu`) and writes the +/// EMA-blended mean to `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` for +/// downstream consumption by `AuxHorizonUpdateOps`. +/// +/// Single-block kernel; block-tree-reduce over the per-sample buffer. +/// Pearl-A first-observation bootstrap; α=0.05 EMA thereafter. +/// +/// # Pearls applied +/// - `feedback_no_atomicadd.md` — single block, block-tree-reduce in shmem. +/// - `pearl_first_observation_bootstrap.md` — sentinel = 0.0 → REPLACE. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded at construction; on-device `count == 0` guard. +#[allow(missing_debug_implementations)] +pub(crate) struct AvgWinHoldTimeUpdateOps { + update_kernel: CudaFunction, +} + +impl AvgWinHoldTimeUpdateOps { + /// Block dim used by the producer kernel — must match `BLK_DIM` in + /// `avg_win_hold_time_update_kernel.cu`. + const BLK_DIM: u32 = 256; + /// Steady-state EMA blend rate. Slow enough to filter + /// sample-distribution noise, fast enough to track policy evolution + /// across epochs. + pub(crate) const ALPHA: f32 = 0.05; + + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(AVG_WIN_HOLD_TIME_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("avg_win_hold_time_update cubin load: {e}")))?; + let update_kernel = module + .load_function("avg_win_hold_time_update") + .map_err(|e| MLError::ModelError(format!("avg_win_hold_time_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the producer. + /// + /// Args: + /// - `hold_at_exit_ptr`: f32 device ptr `[total_samples]`. + /// - `trade_profitable_ptr`: i32 device ptr `[total_samples]`. + /// - `total_samples`: N*L (B*T total samples this epoch). + /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer. + /// - `isv_target_idx`: AVG_WIN_HOLD_TIME_BARS_INDEX (451). + /// - `sentinel`: SENTINEL_AVG_WIN_HOLD_TIME_BARS (0.0). + /// - `alpha`: EMA blend rate (use `Self::ALPHA = 0.05`). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + hold_at_exit_ptr: u64, + trade_profitable_ptr: u64, + total_samples: i32, + isv_ptr: u64, + isv_target_idx: i32, + sentinel: f32, + alpha: f32, + ) -> Result<(), MLError> { + let smem_bytes = 2 * Self::BLK_DIM * std::mem::size_of::() as u32; + unsafe { + stream + .launch_builder(&self.update_kernel) + .arg(&hold_at_exit_ptr) + .arg(&trade_profitable_ptr) + .arg(&total_samples) + .arg(&isv_ptr) + .arg(&isv_target_idx) + .arg(&sentinel) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (Self::BLK_DIM, 1, 1), + shared_mem_bytes: smem_bytes, + }) + .map_err(|e| MLError::ModelError(format!("avg_win_hold_time_update: {e}")))?; + } + Ok(()) + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 93163d089..28b4cc5bb 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2082,6 +2082,23 @@ pub(crate) static AUX_TRUNK_FORWARD_CUBIN: &[u8] = include_bytes!(concat!(env!(" #[allow(dead_code)] pub(crate) static AUX_TRUNK_BACKWARD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/aux_trunk_backward_kernel.cubin")); +/// SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction horizon +/// producer cubin. Single-thread kernel that drives +/// `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from +/// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` via Pearl-A first-observation +/// bootstrap + slow Wiener-α EMA. Loaded by +/// `gpu_aux_trunk::AuxHorizonUpdateOps`. Per-epoch boundary launch. See +/// `aux_horizon_update_kernel.cu` for the algorithm + plan §C.4b. +pub(crate) static AUX_HORIZON_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/aux_horizon_update_kernel.cubin")); + +/// SP14 Layer C Phase C.4b (2026-05-08): avg winning hold time EMA +/// producer cubin. Single-block kernel that sweeps the per-epoch +/// `hold_at_exit_per_sample` + `trade_profitable_per_sample` buffers and +/// writes `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` for downstream +/// consumption by `aux_horizon_update_kernel`. Loaded by +/// `gpu_aux_trunk::AvgWinHoldTimeUpdateOps`. Per-epoch boundary launch. +pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/avg_win_hold_time_update_kernel.cubin")); + /// Plan C Phase 2 follow-up A.2 (2026-04-29): q-drift rate ISV producer. /// Single-thread single-block cold-path kernel mirroring /// `moe_lambda_eff_kernel.cu` / `kelly_cap_update_kernel.cu`. Reads @@ -2399,7 +2416,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 450; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) +pub(crate) const ISV_TOTAL_DIM: usize = 452; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -7336,6 +7353,20 @@ pub struct GpuDqnTrainer { #[allow(dead_code)] aux_trunk_backward_ops: super::gpu_aux_trunk::AuxTrunkBackwardOps, + /// SP14 Layer C Phase C.4b (2026-05-08): adaptive aux prediction + /// horizon producer. Drives `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` + /// from `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` via Pearl-A + /// first-observation bootstrap + slow Wiener-α EMA. Single-thread + /// kernel; per-epoch boundary launch. See `aux_horizon_update_kernel.cu`. + aux_horizon_update_ops: super::gpu_aux_trunk::AuxHorizonUpdateOps, + + /// SP14 Layer C Phase C.4b (2026-05-08): avg winning hold time EMA + /// producer. Sweeps the per-epoch `hold_at_exit_per_sample` + + /// `trade_profitable_per_sample` buffers and writes + /// `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]`. Block-tree-reduce + /// (no atomicAdd); per-epoch boundary launch. + avg_win_hold_time_update_ops: super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps, + // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. speculative_h_s2: CudaSlice, @@ -14049,6 +14080,75 @@ impl GpuDqnTrainer { Ok(()) } + /// SP14 Layer C Phase C.4b (2026-05-08): launch the aux prediction + /// horizon producer chain at per-epoch boundary. + /// + /// Two sequential single-launch kernels: + /// 1. `avg_win_hold_time_update` — sweeps per-sample + /// `hold_at_exit_per_sample` + `trade_profitable_per_sample`, + /// block-tree-reduces winning-trade hold-time mean, EMA-blends + /// into `ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]` (Pearl-A + /// bootstrap then α=0.05). + /// 2. `aux_horizon_update` — single-thread Pearl-A bootstrap + + /// slow Wiener-α (α=0.01 fixed) drives + /// `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` from the slot above. + /// + /// Per-epoch boundary launch — both signals are slow-moving and + /// per-step launches would track sample noise. Caller passes the + /// per-sample buffer pointers + length from the collector's + /// `hold_at_exit_per_sample` / `trade_profitable_per_sample` slabs. + /// + /// Pearls applied: + /// - `pearl_first_observation_bootstrap.md` + /// - `pearl_wiener_optimal_adaptive_alpha.md` + /// - `pearl_no_host_branches_in_captured_graph.md` + /// - `feedback_no_atomicadd.md` (block-tree-reduce in shmem only). + pub(crate) fn launch_aux_horizon_chain( + &self, + hold_at_exit_dev_ptr: u64, + trade_profitable_dev_ptr: u64, + total_samples: i32, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp14_isv_slots::{ + AUX_PRED_HORIZON_BARS_INDEX, AUX_PRED_HORIZON_BARS_MAX, + AUX_PRED_HORIZON_BARS_MIN, AVG_WIN_HOLD_TIME_BARS_INDEX, + SENTINEL_AUX_PRED_HORIZON_BARS, SENTINEL_AVG_WIN_HOLD_TIME_BARS, + }; + + debug_assert!( + self.isv_signals_dev_ptr != 0, + "launch_aux_horizon_chain: isv_signals_dev_ptr must be allocated" + ); + + // Step 1: avg winning hold time EMA producer. + self.avg_win_hold_time_update_ops.launch( + &self.stream, + hold_at_exit_dev_ptr, + trade_profitable_dev_ptr, + total_samples, + self.isv_signals_dev_ptr, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + SENTINEL_AVG_WIN_HOLD_TIME_BARS, + super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps::ALPHA, + )?; + + // Step 2: horizon producer (reads slot 451 from step 1). + // No target-variance EMA in C.4b → pass -1 for var idx (kernel + // falls back to fixed α=0.01). + self.aux_horizon_update_ops.launch( + &self.stream, + self.isv_signals_dev_ptr, + AUX_PRED_HORIZON_BARS_INDEX as i32, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + -1, + AUX_PRED_HORIZON_BARS_MIN, + AUX_PRED_HORIZON_BARS_MAX, + SENTINEL_AUX_PRED_HORIZON_BARS, + )?; + + Ok(()) + } + /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. /// @@ -22200,6 +22300,17 @@ impl GpuDqnTrainer { let aux_trunk_backward_ops = super::gpu_aux_trunk::AuxTrunkBackwardOps::new(&stream)?; + // SP14 Layer C Phase C.4b (2026-05-08): aux prediction horizon + // producer + winning-hold-time EMA producer. Both pre-load their + // `CudaFunction` handle at construction per + // `pearl_no_host_branches_in_captured_graph.md`. Per-epoch + // boundary launch (horizon is slow-moving — per-step would track + // sample noise). + let aux_horizon_update_ops = + super::gpu_aux_trunk::AuxHorizonUpdateOps::new(&stream)?; + let avg_win_hold_time_update_ops = + super::gpu_aux_trunk::AvgWinHoldTimeUpdateOps::new(&stream)?; + // ── Speculative inference cache ── let speculative_h_s2 = stream.alloc_zeros::(b * sh2) .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?; @@ -23004,6 +23115,9 @@ impl GpuDqnTrainer { aux_trunk_forward_ops, // SP14 Layer C Phase C.4 (2026-05-08): aux trunk backward orchestrator. aux_trunk_backward_ops, + // SP14 Layer C Phase C.4b (2026-05-08): aux horizon + winning-hold-time producers. + aux_horizon_update_ops, + avg_win_hold_time_update_ops, speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 5c62eb999..903e42177 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -4090,19 +4090,39 @@ impl GpuExperienceCollector { .alloc_zeros::(total) .map_err(|e| MLError::ModelError(format!("alloc aux_sign_labels: {e}")))?; { + // SP14 Layer C Phase C.4b (2026-05-08): horizon H is no longer + // a host-passed scalar derived from `config.hindsight_lookahead` + // — it is read from `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` on + // device. The slot is initialized to sentinel 60.0 by the fold- + // boundary registry and adaptively driven by + // `aux_horizon_update_kernel` from observed avg winning hold + // time. See plan §C.4b. + // + // The collector's `isv_signals_dev_ptr` is set by the trainer + // post-construction; if NULL (uncommon — should be set before + // any aux label launch in production), the kernel would read + // garbage. Assert non-zero. + use crate::cuda_pipeline::sp14_isv_slots::AUX_PRED_HORIZON_BARS_INDEX; + debug_assert!( + self.isv_signals_dev_ptr != 0, + "aux_sign_label_kernel: isv_signals_dev_ptr must be set before \ + collect_experiences_gpu — H is read from ISV[450] on device." + ); let total_i32 = total as i32; let total_bars_i32 = config.total_bars; - let lookahead_i32 = config.hindsight_lookahead.max(1); + let isv_h_idx_i32 = AUX_PRED_HORIZON_BARS_INDEX as i32; + let isv_dev_ptr = self.isv_signals_dev_ptr; let blocks = ((total + 255) / 256) as u32; unsafe { self.stream .launch_builder(&self.aux_sign_label_kernel) .arg(&targets_buf.dev_ptr) // const float* targets (mapped pinned dev_ptr) .arg(&bar_indices_gpu_ptr) // const int* bar_indices (mapped pinned dev_ptr) + .arg(&isv_dev_ptr) // const float* isv + .arg(&isv_h_idx_i32) // int isv_h_idx .arg(&aux_sign_labels) // int* out_labels .arg(&total_i32) // int total .arg(&total_bars_i32) // int total_bars - .arg(&lookahead_i32) // int lookahead .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), @@ -4563,7 +4583,16 @@ impl GpuExperienceCollector { // `aux_sign_label_kernel` launch (line ~3776 below) so // both encode the same K=2 "up vs not-up" boundary. { - let lookahead_i32 = config.hindsight_lookahead.max(1); + // SP14 Layer C Phase C.4b (2026-05-08): H is read from + // ISV[AUX_PRED_HORIZON_BARS_INDEX=450] on device, + // mirrors trajectory kernel migration. See plan §C.4b. + use crate::cuda_pipeline::sp14_isv_slots::AUX_PRED_HORIZON_BARS_INDEX; + debug_assert!( + self.isv_signals_dev_ptr != 0, + "aux_sign_label_per_step_kernel: isv_signals_dev_ptr must be set." + ); + let isv_h_idx_i32 = AUX_PRED_HORIZON_BARS_INDEX as i32; + let isv_dev_ptr = self.isv_signals_dev_ptr; let labels_dev = self.exp_aux_nb_label_buf.raw_ptr(); let targets_dev = targets_buf.dev_ptr; let starts_dev = self.episode_starts_buf.raw_ptr(); @@ -4575,10 +4604,11 @@ impl GpuExperienceCollector { .arg(&targets_dev) .arg(&starts_dev) .arg(&t_i32) + .arg(&isv_dev_ptr) + .arg(&isv_h_idx_i32) .arg(&labels_dev) .arg(&n_i32) .arg(&total_bars) - .arg(&lookahead_i32) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), @@ -5798,6 +5828,26 @@ impl GpuExperienceCollector { ptr } + /// SP14 Layer C Phase C.4b (2026-05-08): raw device pointer to + /// `hold_at_exit_per_sample` `[alloc_episodes * alloc_timesteps]` — + /// `saved_hold_time` on exit/reverse, else 0.0. Consumed by + /// `avg_win_hold_time_update_kernel` at the per-epoch boundary to + /// compute the EMA of winning-trade hold time. + pub fn hold_at_exit_per_sample_dev_ptr(&self) -> u64 { + let (ptr, _guard) = self.hold_at_exit_per_sample.device_ptr(&self.stream); + ptr + } + + /// SP14 Layer C Phase C.4b (2026-05-08): raw device pointer to + /// `trade_profitable_per_sample` `[alloc_episodes * alloc_timesteps]` + /// — 1 iff (trade_close && step_ret > 0). Consumed by + /// `avg_win_hold_time_update_kernel` to mask the hold-time mean to + /// winning trades only. + pub fn trade_profitable_per_sample_dev_ptr(&self) -> u64 { + let (ptr, _guard) = self.trade_profitable_per_sample.device_ptr(&self.stream); + ptr + } + /// Last experience count from the most recent `collect_experiences_gpu` call. pub fn last_experience_count(&self) -> usize { self.last_experience_count } diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 9eef639a1..5b4c433fa 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -71,6 +71,40 @@ pub const SENTINEL_H_S2_AUX_RMS: f32 = 0.0; // Pearl-A bootstrap pub const SP14_C_SLOT_BASE: usize = 444; pub const SP14_C_SLOT_END: usize = 450; +// ── SP14 Layer C Phase C.4b — aux prediction horizon (added 2026-05-08) ── +// Two slots driving the new ISV-driven aux label horizon: +// +// - AUX_PRED_HORIZON_BARS_INDEX (450): adaptive horizon H read by both +// `aux_sign_label_kernel.cu` and `aux_sign_label_per_step_kernel.cu` +// to compute label = (p[bar+H] > p[bar]). FoldReset sentinel +// SENTINEL_AUX_PRED_HORIZON_BARS=60.0 (≈1hr @ 1-min bars). Pearl-A +// first-observation bootstrap on first valid AVG_WIN_HOLD_TIME_BARS +// observation; Wiener-α EMA blend thereafter (slow α=0.01 — no +// variance EMA available for the target). +// +// - AVG_WIN_HOLD_TIME_BARS_INDEX (451): aggregate EMA of winning-trade +// hold time (bars) read by the horizon producer. Computed by +// `avg_win_hold_time_update_kernel.cu` from the per-sample buffers +// `hold_at_exit_per_sample` + `trade_profitable_per_sample` populated +// by `unified_env_step_core` in `experience_kernels.cu`. FoldReset +// sentinel SENTINEL_AVG_WIN_HOLD_TIME_BARS=0.0 (Pearl-A bootstrap). +// +// Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.4b +// Case B (Step 5b): existing trade-close infrastructure has per-sample +// `hold_at_exit_per_sample[N*L]` + `trade_profitable_per_sample[N*L]`, +// but no aggregate slot — added new slot 451 to expose the EMA. +pub const AUX_PRED_HORIZON_BARS_INDEX: usize = 450; +pub const AVG_WIN_HOLD_TIME_BARS_INDEX: usize = 451; + +// ── Aux horizon producer constants ──────────────────────────────────── +pub const SENTINEL_AUX_PRED_HORIZON_BARS: f32 = 60.0; // Pearl-A bootstrap +pub const SENTINEL_AVG_WIN_HOLD_TIME_BARS: f32 = 0.0; // Pearl-A bootstrap +pub const AUX_PRED_HORIZON_BARS_MIN: f32 = 1.0; // floor (cannot predict past) +pub const AUX_PRED_HORIZON_BARS_MAX: f32 = 240.0; // ceiling (rollout buffer guard) + +pub const SP14_C4B_SLOT_BASE: usize = 450; +pub const SP14_C4B_SLOT_END: usize = 452; + #[cfg(test)] mod tests { use super::*; @@ -125,4 +159,26 @@ mod tests { SP14_C_SLOT_END, ISV_TOTAL_DIM, ); } + + /// Lock SP14 Layer C Phase C.4b aux-prediction-horizon slot layout. + /// Two contiguous slots [450..452) — adaptive horizon H + winning + /// hold-time EMA target. + #[test] + fn sp14_c4b_aux_horizon_slot_layout_locked() { + assert_eq!(SP14_C4B_SLOT_BASE, 450); + assert_eq!(SP14_C4B_SLOT_END, 452); + assert_eq!(AUX_PRED_HORIZON_BARS_INDEX, 450); + assert_eq!(AVG_WIN_HOLD_TIME_BARS_INDEX, 451); + } + + #[test] + fn all_sp14_c4b_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP14_C4B_SLOT_END <= ISV_TOTAL_DIM, + "SP14_C4B_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for SP14 Layer C Phase C.4b aux horizon slots; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP14_C4B_SLOT_END, ISV_TOTAL_DIM, + ); + } } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index edfdfb1fd..bb1425257 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -999,6 +999,28 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[H_S2_AUX_RMS_EMA_INDEX=449] — SP14 Layer C aux-trunk output RMS EMA (mirrors GRN trunk's H_S2_RMS_EMA_INDEX=96 for the new aux trunk's `h_s2_aux` output). Stateful kernel output produced in Phase C.6 by `h_s2_aux_rms_ema_kernel`. FoldReset sentinel SENTINEL_H_S2_AUX_RMS=0.0 — Pearl-A first-observation bootstrap per `pearl_first_observation_bootstrap.md`. Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.1.", }, + // ── SP14 Layer C Phase C.4b (2026-05-08): aux prediction horizon ── + // Two slots [450, 451] driving the new ISV-driven aux label + // horizon. AUX_PRED_HORIZON_BARS_INDEX=450 is the H consumed by + // both `aux_sign_label_kernel.cu` and + // `aux_sign_label_per_step_kernel.cu`; sentinel 60.0 (≈1hr @ + // 1-min bars) is the cold-start default until the adaptive + // producer (`aux_horizon_update_kernel`) lands the first + // observation. AVG_WIN_HOLD_TIME_BARS_INDEX=451 is the EMA + // target produced by `avg_win_hold_time_update_kernel` from + // the existing per-sample `hold_at_exit_per_sample` + + // `trade_profitable_per_sample` buffers; sentinel 0.0 is the + // Pearl-A bootstrap value (no winning trades observed yet). + RegistryEntry { + name: "sp14_c4b_aux_pred_horizon_bars", + category: ResetCategory::FoldReset, + description: "ISV[AUX_PRED_HORIZON_BARS_INDEX=450] — SP14 Layer C Phase C.4b adaptive aux prediction horizon H. Read by `aux_sign_label_kernel.cu` + `aux_sign_label_per_step_kernel.cu` to compute label = (p[bar+H] > p[bar]). FoldReset sentinel SENTINEL_AUX_PRED_HORIZON_BARS=60.0 (≈1hr @ 1-min bars). Pearl-A first-observation bootstrap on first valid winning-trade observation; Wiener-α slow EMA blend (α=0.01 fixed) thereafter via `aux_horizon_update_kernel`. Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.4b.", + }, + RegistryEntry { + name: "sp14_c4b_avg_win_hold_time_bars", + category: ResetCategory::FoldReset, + description: "ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] — SP14 Layer C Phase C.4b avg winning hold time EMA (bars). Aggregated from `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers (populated by `unified_env_step_core` in `experience_kernels.cu`) by `avg_win_hold_time_update_kernel` at per-epoch boundary. FoldReset sentinel SENTINEL_AVG_WIN_HOLD_TIME_BARS=0.0 — Pearl-A first-observation bootstrap. Drives the horizon producer for ISV[450]. Plan: §C.4b.", + }, // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── // Two ISV slots [407, 408]: // - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 21a582e04..1c0dd2957 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -391,6 +391,26 @@ impl DQNTrainer { if let Err(e) = fused.trainer().launch_sp5_pearl_6_kelly(ps_dev_ptr, n_envs) { tracing::warn!(error = %e, "SP5 Pearl 6 producer launch failed"); } + // SP14 Layer C Phase C.4b (2026-05-08): adaptive aux + // prediction horizon — per-epoch boundary launch chain. + // Step 1 (`avg_win_hold_time_update`) sweeps the + // per-sample `hold_at_exit_per_sample` + + // `trade_profitable_per_sample` buffers (populated by + // `unified_env_step_core`) and writes the winning-trade + // hold-time EMA to ISV[451]. Step 2 (`aux_horizon_update`) + // drives ISV[450] from ISV[451] via Pearl-A bootstrap + + // slow Wiener-α EMA. Both signals are slow-moving; + // per-epoch cadence is appropriate. Plan: §C.4b. + let hold_at_exit_ptr = collector.hold_at_exit_per_sample_dev_ptr(); + let trade_profitable_ptr = collector.trade_profitable_per_sample_dev_ptr(); + let total_samples = (collector.alloc_episodes() * collector.alloc_timesteps()) as i32; + if let Err(e) = fused.trainer().launch_aux_horizon_chain( + hold_at_exit_ptr, + trade_profitable_ptr, + total_samples, + ) { + tracing::warn!(epoch, "launch_aux_horizon_chain failed (non-fatal): {e}"); + } } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. @@ -8050,6 +8070,33 @@ impl DQNTrainer { ); } } + // SP14 Layer C Phase C.4b (2026-05-08): aux prediction horizon + // sentinels. Both slots use Pearl-A first-observation bootstrap + // — the FoldReset rewrites them to sentinel so the kernels + // detect "first observation" cleanly on the new fold's first + // launch (avoids cross-fold EMA contamination). + "sp14_c4b_aux_pred_horizon_bars" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + AUX_PRED_HORIZON_BARS_INDEX, SENTINEL_AUX_PRED_HORIZON_BARS, + }; + fused.trainer().write_isv_signal_at( + AUX_PRED_HORIZON_BARS_INDEX, + SENTINEL_AUX_PRED_HORIZON_BARS, + ); + } + } + "sp14_c4b_avg_win_hold_time_bars" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + AVG_WIN_HOLD_TIME_BARS_INDEX, SENTINEL_AVG_WIN_HOLD_TIME_BARS, + }; + fused.trainer().write_isv_signal_at( + AVG_WIN_HOLD_TIME_BARS_INDEX, + SENTINEL_AVG_WIN_HOLD_TIME_BARS, + ); + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/tests/aux_trunk_oracle_tests.rs b/crates/ml/tests/aux_trunk_oracle_tests.rs index a72bc0946..db933beb3 100644 --- a/crates/ml/tests/aux_trunk_oracle_tests.rs +++ b/crates/ml/tests/aux_trunk_oracle_tests.rs @@ -735,4 +735,434 @@ mod gpu { kernel_path.display() ); } + + // ════════════════════════════════════════════════════════════════════ + // SP14 Layer C Phase C.4b oracle tests (2026-05-08). + // + // Added 4 tests exercising the new ISV-driven aux prediction horizon: + // + // - aux_sign_label_h_bar_horizon — label correctness for a known H + // - aux_sign_label_lookahead_mask — masking when t+H runs past total_bars + // - aux_horizon_pearl_a_bootstrap — sentinel REPLACED, not blended + // - aux_horizon_converges_to_steady_target — Wiener-α EMA convergence + // - aux_horizon_holds_sentinel_with_no_winning_trades — cold-start guard + // + // Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.4b + // ════════════════════════════════════════════════════════════════════ + + const AUX_SIGN_LABEL_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/aux_sign_label_kernel.cubin")); + const AUX_HORIZON_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/aux_horizon_update_kernel.cubin")); + + /// AUX_PRED_HORIZON_BARS_INDEX — must match `sp14_isv_slots.rs`. + const AUX_PRED_HORIZON_BARS_INDEX: usize = 450; + /// AVG_WIN_HOLD_TIME_BARS_INDEX — must match `sp14_isv_slots.rs`. + const AVG_WIN_HOLD_TIME_BARS_INDEX: usize = 451; + /// SENTINEL_AUX_PRED_HORIZON_BARS — must match `sp14_isv_slots.rs`. + const SENTINEL_H: f32 = 60.0; + /// AUX_PRED_HORIZON_BARS_MIN/MAX — fundamental bounds. + const H_MIN: f32 = 1.0; + const H_MAX: f32 = 240.0; + + /// ISV size used by the horizon producer test (must be > 451). + const ISV_TEST_SIZE: usize = 460; + + fn load_aux_sign_label(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(AUX_SIGN_LABEL_CUBIN.to_vec()) + .expect("load aux_sign_label cubin"); + module + .load_function("aux_sign_label_kernel") + .expect("load aux_sign_label_kernel function") + } + + fn load_aux_horizon_update(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(AUX_HORIZON_UPDATE_CUBIN.to_vec()) + .expect("load aux_horizon_update cubin"); + module + .load_function("aux_horizon_update") + .expect("load aux_horizon_update function") + } + + /// Helper: launch `aux_sign_label_kernel` on the given pinned-mapped + /// targets / bar_indices / isv buffers. Returns the device-side + /// labels copied back to host. + fn run_aux_sign_label( + stream: &Arc, + kernel: &CudaFunction, + targets_buf: &MappedF32Buffer, + bar_indices_buf: &MappedF32Buffer, + isv_buf: &MappedF32Buffer, + total: usize, + total_bars: i32, + ) -> Vec { + let labels_dev = stream + .alloc_zeros::(total) + .expect("alloc out_labels"); + let total_i32 = total as i32; + let isv_h_idx_i32 = AUX_PRED_HORIZON_BARS_INDEX as i32; + let blocks = ((total + 255) / 256) as u32; + let targets_dev = targets_buf.dev_ptr; + let bar_indices_dev = bar_indices_buf.dev_ptr; + let isv_dev = isv_buf.dev_ptr; + unsafe { + stream + .launch_builder(kernel) + .arg(&targets_dev) + .arg(&bar_indices_dev) + .arg(&isv_dev) + .arg(&isv_h_idx_i32) + .arg(&labels_dev) + .arg(&total_i32) + .arg(&total_bars) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .expect("aux_sign_label_kernel launch"); + } + let mut out = vec![0_i32; total]; + stream.memcpy_dtoh(&labels_dev, &mut out).expect("dtoh labels"); + stream.synchronize().expect("sync"); + out + } + + /// Helper: build a synthetic `targets [total_bars, 6]` mapped-pinned + /// buffer with `prices` written into column 2 (raw_close). All other + /// columns are zero-filled. + fn build_targets_buf(prices: &[f32]) -> MappedF32Buffer { + let total_bars = prices.len(); + let buf = unsafe { MappedF32Buffer::new(total_bars * 6) } + .expect("MappedF32Buffer::new(targets)"); + let host = unsafe { std::slice::from_raw_parts_mut(buf.host_ptr, total_bars * 6) }; + host.fill(0.0); + for (i, &p) in prices.iter().enumerate() { + host[i * 6 + 2] = p; + } + buf + } + + /// Helper: build a `bar_indices [total]` mapped-pinned buffer with + /// the i32 values reinterpreted as f32 (MappedF32Buffer is our only + /// generic mapped-pinned allocator). The kernel reads it as `int*` + /// — bit-pattern preservation matters more than f32 representability. + fn build_bar_indices_buf(indices: &[i32]) -> MappedF32Buffer { + let n = indices.len(); + let buf = unsafe { MappedF32Buffer::new(n) } + .expect("MappedF32Buffer::new(bar_indices)"); + // bytewise copy: same memory layout, no f32 interpretation + unsafe { + let dst = buf.host_ptr as *mut i32; + std::ptr::copy_nonoverlapping(indices.as_ptr(), dst, n); + } + buf + } + + /// Helper: build a `[ISV_TEST_SIZE]` mapped-pinned ISV buffer with + /// AUX_PRED_HORIZON_BARS_INDEX seeded to `h` and other slots zero. + fn build_isv_buf(h: f32) -> MappedF32Buffer { + let buf = unsafe { MappedF32Buffer::new(ISV_TEST_SIZE) } + .expect("MappedF32Buffer::new(isv)"); + let host = unsafe { std::slice::from_raw_parts_mut(buf.host_ptr, ISV_TEST_SIZE) }; + host.fill(0.0); + host[AUX_PRED_HORIZON_BARS_INDEX] = h; + buf + } + + /// Step 6 oracle: aux_sign_label produces correct labels for a + /// known H on three deterministic price series. + /// + /// 1. Linearly increasing prices, H=30, T=200 → labels 1 in + /// `[0..170)` and -1 in `[170..200)` (lookahead mask). + /// 2. Linearly decreasing prices, H=30, T=200 → labels 0 in + /// `[0..170)` and -1 in `[170..200)`. + /// 3. Sine wave, H=30 → analytic comparison of `p[t+30] > p[t]`. + #[test] + #[ignore = "requires GPU"] + fn aux_sign_label_h_bar_horizon() { + let stream = make_test_stream(); + let kernel = load_aux_sign_label(&stream); + + const T: usize = 200; + const H: i32 = 30; + let total_bars = T as i32; + let bar_indices: Vec = (0..T as i32).collect(); + let bar_indices_buf = build_bar_indices_buf(&bar_indices); + let isv_buf = build_isv_buf(H as f32); + + // Case 1: linearly increasing + let prices_up: Vec = (0..T).map(|i| i as f32).collect(); + let targets_up = build_targets_buf(&prices_up); + let labels_up = + run_aux_sign_label(&stream, &kernel, &targets_up, &bar_indices_buf, &isv_buf, T, total_bars); + let lookahead_cutoff = T as i32 - H; + for (i, &lbl) in labels_up.iter().enumerate() { + let i_i32 = i as i32; + if i_i32 + H < total_bars { + assert_eq!(lbl, 1, "increasing series: bar {i} expected 1 (p[t+H]>p[t]), got {lbl}"); + } else { + assert_eq!(lbl, -1, "increasing series: bar {i} expected mask -1 (t+H>=total_bars), got {lbl}"); + } + } + eprintln!("Case 1 (increasing): valid bars={lookahead_cutoff}, all 1 — OK"); + + // Case 2: linearly decreasing + let prices_dn: Vec = (0..T).map(|i| (T - i) as f32).collect(); + let targets_dn = build_targets_buf(&prices_dn); + let labels_dn = + run_aux_sign_label(&stream, &kernel, &targets_dn, &bar_indices_buf, &isv_buf, T, total_bars); + for (i, &lbl) in labels_dn.iter().enumerate() { + let i_i32 = i as i32; + if i_i32 + H < total_bars { + assert_eq!(lbl, 0, "decreasing series: bar {i} expected 0 (p[t+H] = + (0..T).map(|i| (i as f32 * 0.1).sin()).collect(); + let targets_sin = build_targets_buf(&prices_sin); + let labels_sin = + run_aux_sign_label(&stream, &kernel, &targets_sin, &bar_indices_buf, &isv_buf, T, total_bars); + for (i, &lbl) in labels_sin.iter().enumerate() { + let i_i32 = i as i32; + if i_i32 + H >= total_bars { + assert_eq!(lbl, -1, "sine: bar {i} mask expected -1, got {lbl}"); + continue; + } + let expected = if prices_sin[i + H as usize] > prices_sin[i] { 1 } else { 0 }; + assert_eq!(lbl, expected, + "sine: bar {i} expected {expected} (p[{}]={:.4} vs p[{i}]={:.4}), got {lbl}", + i + H as usize, prices_sin[i + H as usize], prices_sin[i]); + } + eprintln!("Case 3 (sine): valid bars={lookahead_cutoff}, all match analytic — OK"); + } + + /// Step 7 oracle: lookahead mask invariant — for H=60 with + /// total_bars=100, labels in `[40..100)` must be -1 and labels in + /// `[0..40)` must be ∈ {0, 1} (valid). + #[test] + #[ignore = "requires GPU"] + fn aux_sign_label_lookahead_mask() { + let stream = make_test_stream(); + let kernel = load_aux_sign_label(&stream); + + const T: usize = 100; + const H: i32 = 60; + let total_bars = T as i32; + let bar_indices: Vec = (0..T as i32).collect(); + let bar_indices_buf = build_bar_indices_buf(&bar_indices); + let isv_buf = build_isv_buf(H as f32); + let prices: Vec = (0..T).map(|i| (i as f32 * 0.5).sin()).collect(); + let targets = build_targets_buf(&prices); + let labels = + run_aux_sign_label(&stream, &kernel, &targets, &bar_indices_buf, &isv_buf, T, total_bars); + + for (i, &lbl) in labels.iter().enumerate() { + let i_i32 = i as i32; + if i_i32 + H >= total_bars { + assert_eq!(lbl, -1, + "lookahead mask: bar {i} (t+H={}) >= total_bars={total_bars} expected -1, got {lbl}", + i_i32 + H); + } else { + assert!(lbl == 0 || lbl == 1, + "valid window: bar {i} expected ∈ {{0, 1}}, got {lbl}"); + } + } + + let mask_count = labels.iter().filter(|&&l| l == -1).count(); + let valid_count = labels.iter().filter(|&&l| l == 0 || l == 1).count(); + assert_eq!(mask_count, H as usize, + "expected {} masked bars (last H rows), got {mask_count}", H); + assert_eq!(valid_count, T - H as usize, + "expected {} valid bars, got {valid_count}", T - H as usize); + eprintln!("lookahead_mask: {mask_count}/{T} masked, {valid_count}/{T} valid — OK"); + } + + /// Helper: launch `aux_horizon_update` once on the given mapped-pinned + /// ISV buffer. + fn run_aux_horizon_update( + stream: &Arc, + kernel: &CudaFunction, + isv_buf: &MappedF32Buffer, + isv_h_idx: i32, + isv_h_target_idx: i32, + isv_h_target_var_idx: i32, + h_min: f32, + h_max: f32, + sentinel_h: f32, + ) { + let isv_dev = isv_buf.dev_ptr; + unsafe { + stream + .launch_builder(kernel) + .arg(&isv_dev) + .arg(&isv_h_idx) + .arg(&isv_h_target_idx) + .arg(&isv_h_target_var_idx) + .arg(&h_min) + .arg(&h_max) + .arg(&sentinel_h) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("aux_horizon_update launch"); + } + stream.synchronize().expect("sync"); + } + + /// Step 7b oracle: Pearl-A first-observation bootstrap — sentinel + /// must be REPLACED, not blended. + #[test] + #[ignore = "requires GPU"] + fn aux_horizon_pearl_a_bootstrap() { + let stream = make_test_stream(); + let kernel = load_aux_horizon_update(&stream); + + // 1. Sentinel state: H=60.0, h_target=90.0 (first valid observation). + let isv_buf = build_isv_buf(SENTINEL_H); + let host = unsafe { std::slice::from_raw_parts_mut(isv_buf.host_ptr, ISV_TEST_SIZE) }; + host[AVG_WIN_HOLD_TIME_BARS_INDEX] = 90.0; + + run_aux_horizon_update( + &stream, + &kernel, + &isv_buf, + AUX_PRED_HORIZON_BARS_INDEX as i32, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + -1, // no variance EMA → fixed α + H_MIN, + H_MAX, + SENTINEL_H, + ); + + let h_after_bootstrap = host[AUX_PRED_HORIZON_BARS_INDEX]; + assert!( + (h_after_bootstrap - 90.0).abs() < 1e-3, + "Pearl-A bootstrap: expected H=90.0 (REPLACED, not blended), got {h_after_bootstrap}", + ); + eprintln!("Pearl-A bootstrap: H={SENTINEL_H} sentinel + target=90.0 → H={h_after_bootstrap} (REPLACED)"); + + // 2. Subsequent observation: target=100.0 → blended (NOT replaced). + // Current H is now 90.0, NOT sentinel — kernel takes the EMA path. + host[AVG_WIN_HOLD_TIME_BARS_INDEX] = 100.0; + run_aux_horizon_update( + &stream, + &kernel, + &isv_buf, + AUX_PRED_HORIZON_BARS_INDEX as i32, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + -1, + H_MIN, + H_MAX, + SENTINEL_H, + ); + + let h_after_blend = host[AUX_PRED_HORIZON_BARS_INDEX]; + // α=0.01 → blend = 0.99*90 + 0.01*100 = 90.1 ± float noise + let expected = 0.99_f32 * 90.0 + 0.01_f32 * 100.0; + assert!( + (h_after_blend - expected).abs() < 1e-3, + "second observation: expected blended H≈{expected}, got {h_after_blend} \ + (must be in (90, 100) — NOT replaced)", + ); + assert!( + h_after_blend > 90.0 && h_after_blend < 100.0, + "blended H={h_after_blend} not in (90, 100) — should be small step from 90 toward 100", + ); + eprintln!("Second obs: H={h_after_blend} (blended toward 100, expected ≈{expected})"); + } + + /// Step 7c oracle: Wiener-α EMA convergence — with α=0.01 fixed + /// (no variance EMA), H starting at 50.0 with steady target=100.0 + /// should converge per the geometric `(1-α)^k * 50 + (1-(1-α)^k) * 100`. + #[test] + #[ignore = "requires GPU"] + fn aux_horizon_converges_to_steady_target() { + let stream = make_test_stream(); + let kernel = load_aux_horizon_update(&stream); + + // Post-bootstrap state: H=50.0, target=100.0 (steady). + // 50.0 differs from sentinel 60.0 → bootstrap path NOT triggered. + let isv_buf = build_isv_buf(50.0); + let host = unsafe { std::slice::from_raw_parts_mut(isv_buf.host_ptr, ISV_TEST_SIZE) }; + host[AVG_WIN_HOLD_TIME_BARS_INDEX] = 100.0; + + const ITERS: usize = 100; + for _ in 0..ITERS { + run_aux_horizon_update( + &stream, + &kernel, + &isv_buf, + AUX_PRED_HORIZON_BARS_INDEX as i32, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + -1, + H_MIN, + H_MAX, + SENTINEL_H, + ); + } + + let h_final = host[AUX_PRED_HORIZON_BARS_INDEX]; + // Closed-form for fixed-α geometric blend: + // H_k = (1-α)^k * H_0 + (1 - (1-α)^k) * target + // α=0.01, k=100, H_0=50, target=100 + // (1-α)^100 ≈ 0.366032 + // H_100 ≈ 0.366 * 50 + 0.634 * 100 ≈ 81.7 + let alpha: f32 = 0.01; + let one_minus_alpha = 1.0 - alpha; + let decay = one_minus_alpha.powi(ITERS as i32); + let expected = decay * 50.0 + (1.0 - decay) * 100.0; + let tol = 0.05; // accumulated f32 noise across 100 iters + assert!( + (h_final - expected).abs() < tol, + "Wiener-α convergence: expected H≈{expected:.4} after {ITERS} iters, got {h_final:.4} \ + (decay=(1-{alpha})^{ITERS}≈{decay:.6})", + ); + eprintln!("Convergence: H_0=50, target=100, α={alpha}, after {ITERS} iters → H={h_final:.4} (expected {expected:.4})"); + } + + /// Step 7d oracle: cold-start "no winning trades" guard — when + /// `h_target == 0.0` (no winning trade observed yet), the kernel + /// MUST keep the sentinel and not blend toward zero. + #[test] + #[ignore = "requires GPU"] + fn aux_horizon_holds_sentinel_with_no_winning_trades() { + let stream = make_test_stream(); + let kernel = load_aux_horizon_update(&stream); + + // Sentinel state: H=60.0, target=0.0 (no winning trade yet). + let isv_buf = build_isv_buf(SENTINEL_H); + let host = unsafe { std::slice::from_raw_parts_mut(isv_buf.host_ptr, ISV_TEST_SIZE) }; + host[AVG_WIN_HOLD_TIME_BARS_INDEX] = 0.0; + + run_aux_horizon_update( + &stream, + &kernel, + &isv_buf, + AUX_PRED_HORIZON_BARS_INDEX as i32, + AVG_WIN_HOLD_TIME_BARS_INDEX as i32, + -1, + H_MIN, + H_MAX, + SENTINEL_H, + ); + + let h_after = host[AUX_PRED_HORIZON_BARS_INDEX]; + assert!( + (h_after - SENTINEL_H).abs() < 1e-6, + "no winning trades guard: expected H=sentinel {SENTINEL_H} unchanged, got {h_after}", + ); + eprintln!("No-winning-trades guard: H={h_after} (sentinel preserved)"); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 0280a76ef..ffabd27f4 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -7254,3 +7254,60 @@ The error surfaces as "load sp15_dd_state cubin: ILLEGAL_ADDRESS" via sticky-cas - Pre-init wiring (this commit): `crates/ml/src/trainers/dqn/trainer/training_loop.rs:1716`, immediately before `self.gpu_experience_collector = Some(collector);`. - Per-epoch re-wiring (preserved): `crates/ml/src/trainers/dqn/trainer/training_loop.rs:855-957`, inside the `for epoch in 0..self.hyperparams.epochs` loop. - Producer setters: `GpuExperienceCollector::set_isv_signals_ptr` / `set_sp15_alpha_warm_count_ptr` / `set_sp15_plasticity_target`; `GpuReplayBuffer::set_isv_signals_ptr` / `set_sp15_dd_trajectory_per_env_ptr` / `set_sp15_per_env_dims`. + + +## SP14 Layer C Phase C.4b — aux label horizon ISV-driven (multi-bar pivot, 2026-05-08) + +**Rationale.** Aux's original label was `(p_{t+1} > p_t)` — pure HFT-scale microstructure noise that's unlearnable at our HFT-MFT trading frequency (multi-bar holds). Without this fix, aux flatlines at ln(2) regardless of trunk separation: the *target itself* is unlearnable. Pivoted to `(p_{t+H} > p_t)` where H is read from `ISV[AUX_PRED_HORIZON_BARS_INDEX=450]` and adaptively driven from observed avg winning hold time. + +**Step 5b finding (Case B).** Searched `crates/ml/src/cuda_pipeline/*.cu` + `*.rs` for existing avg-winning-hold-time infrastructure: +- ✅ Per-sample buffers exist: `hold_at_exit_per_sample [N*L]` (`saved_hold_time` on exit/reverse, populated by `unified_env_step_core` in `experience_kernels.cu:2698`) + `trade_profitable_per_sample [N*L]` (1 iff trade_close && step_ret>0, `experience_kernels.cu:2727`). +- ❌ No aggregate ISV slot tracking the EMA. Existing host-side reduction in `trail_fire_and_hold_per_mag` is per-magnitude HEALTH_DIAG output, not a controller signal. +- ⇒ **Case B**: added `AVG_WIN_HOLD_TIME_BARS_INDEX=451` aggregate slot + new `avg_win_hold_time_update_kernel.cu` producer that block-tree-reduces winning-trade hold-time mean and EMA-blends into the slot. ISV_TOTAL_DIM bumped 450 → **452** (one extra over plan's 451 estimate). + +**Slot layout (added 2026-05-08):** +| Slot | Name | Sentinel | Producer | Consumer(s) | +|------|------|----------|----------|-------------| +| 450 | `AUX_PRED_HORIZON_BARS_INDEX` | 60.0 | `aux_horizon_update_kernel` | `aux_sign_label_kernel.cu`, `aux_sign_label_per_step_kernel.cu` | +| 451 | `AVG_WIN_HOLD_TIME_BARS_INDEX` | 0.0 | `avg_win_hold_time_update_kernel` | `aux_horizon_update_kernel` | + +Both slots use Pearl-A first-observation bootstrap. Slot 451 has α=0.05 EMA; slot 450 has fixed α=0.01 slow blend (no target-variance EMA available, so Wiener-α reduces to its fallback per `pearl_wiener_optimal_adaptive_alpha.md`). + +**Atomic migration (per `feedback_no_partial_refactor`).** Both `aux_sign_label_kernel.cu` (trajectory variant — collector-end) and `aux_sign_label_per_step_kernel.cu` (per-rollout-step variant — used by sp14-β EGF chain) migrated together to the new ISV-driven signature `(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)`. The `lookahead` host-passed scalar argument is removed from both kernels and both launch sites; H is read from ISV inside the kernel (broadcast value, single ISV read per thread). If only one had been migrated, training-time labels (trajectory) and eval-time per-step labels would diverge — fatal inconsistency. + +**Lookahead masking (correctness-critical).** Both kernels preserve the existing `-1` skip sentinel for `bar + H >= total_bars`. The downstream loss-reduce kernel `aux_next_bar_loss_reduce` already supports `-1` masking natively (`aux_heads_kernel.cu:281-336`: rows with `label == -1` contribute 0 to numerator AND 0 to valid_count; backward kernel reads the saved `B_valid` and zeros d_logits for masked rows at line 489-526). No new `valid_mask` parameter required — existing `-1` sentinel convention is sufficient. + +**Producer chain (per-epoch boundary).** New launcher method `GpuDqnTrainer::launch_aux_horizon_chain` orchestrates two sequential single-launch kernels at epoch boundary alongside `launch_kelly_cap_update`: +1. `avg_win_hold_time_update` — single-block, block-tree-reduce over `[total_samples]` (no atomicAdd per `feedback_no_atomicadd`); writes `ISV[451]`. Pearl-A bootstrap (sentinel=0); α=0.05 EMA blend thereafter. +2. `aux_horizon_update` — single-thread; reads `ISV[451]`, drives `ISV[450]` via Pearl-A bootstrap (sentinel=60.0) + fixed-α=0.01 slow blend. Bounds `[h_min=1, h_max=240]` are fundamental floors/ceilings (cannot predict past; rollout buffer guard) per `feedback_isv_for_adaptive_bounds`. + +**Files modified (this commit):** +- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (+2 const slot indices, +4 sentinels/bounds, +2 lock tests) +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (ISV_TOTAL_DIM 450→452; +2 cubin static refs; +2 ops fields; +2 ops constructors; +1 `launch_aux_horizon_chain` method) +- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (+`AuxHorizonUpdateOps` + `AvgWinHoldTimeUpdateOps` + their `new()` + `launch()`) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (2 launch sites: trajectory + per-step kernel signatures bump; +2 dev-ptr getters for `hold_at_exit_per_sample` + `trade_profitable_per_sample`) +- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu` (signature: drop `lookahead` arg, add `isv` + `isv_h_idx`; read H on-device with clamp `[1, 240]`) +- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu` (same migration) +- Create: `crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu` +- Create: `crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu` +- Modify: `crates/ml/build.rs` (+2 cubin entries) +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+2 FoldReset entries) +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+2 reset_named_state dispatch arms; +1 launch site for `launch_aux_horizon_chain` at epoch boundary) +- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (+5 oracle tests) +- Modify: `docs/dqn-wire-up-audit.md` (this section) + +**Tests (8 oracle tests in `aux_trunk_oracle_tests` — 3 preserved + 5 new):** +- ✅ `aux_trunk_forward_matches_numpy_reference` (C.3 — preserved) +- ✅ `aux_trunk_backward_gradient_check` (C.4 — preserved) +- ✅ `aux_trunk_backward_does_not_write_dx` (C.4 — preserved) +- ✅ `aux_sign_label_h_bar_horizon` (NEW — 3 deterministic series: increasing/decreasing/sine, H=30, T=200; verifies labels match analytic and tail H bars are masked -1) +- ✅ `aux_sign_label_lookahead_mask` (NEW — H=60, T=100: verifies last H rows are -1 and first T-H rows are valid {0,1}) +- ✅ `aux_horizon_pearl_a_bootstrap` (NEW — sentinel 60.0 + first observation 90.0 → REPLACE; second observation 100.0 → blend = 0.99*90 + 0.01*100 = 90.1, in (90,100)) +- ✅ `aux_horizon_converges_to_steady_target` (NEW — H_0=50, target=100, α=0.01: closed-form `(1-α)^100 * 50 + (1-(1-α)^100) * 100 = 81.6984` — bit-precise match within 0.05 tol) +- ✅ `aux_horizon_holds_sentinel_with_no_winning_trades` (NEW — sentinel 60.0 + target=0.0 → unchanged 60.0) + +8/8 pass on RTX 3050 Ti. + +**Open follow-ups:** +- Add a target-variance EMA for `AVG_WIN_HOLD_TIME_BARS_INDEX` (Welford-style, single new ISV slot) in a future phase if the fixed α=0.01 blend is too slow/fast on real-data validation. The producer kernel already has the variance-slot path wired; pass `isv_h_target_var_idx >= 0` to enable Wiener-α. +- Pearl note for `pearl_event_driven_reward_density_alignment.md` cross-reference: this phase aligns the AUX *label density* to event density (winning trade close), the same principle as the SP12 reward-density alignment. diff --git a/docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md b/docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md new file mode 100644 index 000000000..5755fd391 --- /dev/null +++ b/docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md @@ -0,0 +1,1243 @@ +# SP14 Layer C: Separate Aux Trunk Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Give the auxiliary next-bar prediction head its own dedicated trunk (parallel to Q's GRN trunk) so it can shape a representation suited to next-bar prediction, while Q's trunk remains shaped purely by Q-loss. Delete the obsolete EGF α-coupling machinery (Coupling B + scale_wire kernel + Schmitt gate). Keep the forward feature wire (Coupling A: `dir_concat_qaux`) so Q's head still sees aux's directional softmax-diff as a feature. + +**Architecture:** Encoder remains shared and Q-shaped. After encoder, the forward path branches: GRN trunk produces `h_s2_q` → Q-heads; new aux trunk (3-layer Linear→ELU MLP) produces `h_s2_aux` → aux head (K=2 softmax). Backward: Q-loss propagates Q-trunk → encoder; aux-loss propagates aux-trunk only (stop-gradient at encoder boundary). ISV is the control plane for aux trunk hyperparameters and adaptive bounds; ISV is NOT used as aux input features. + +**Tech Stack:** Rust 1.85, cudarc 0.16+, CUDA 12.4, custom kernels. RTX 3050 Ti (sm_86) for dev tests; L40S (sm_89) for full validation. + +**Branch:** Continue on `sp15-phase1-honest-numbers`, fork sub-worktree if scope warrants. Forks from `0ec791734` (4-fix HEAD validated by `train-thtxx`). + +--- + +## Design Decisions Locked + +| Question | Decision | Rationale | +|---|---|---| +| Aux trunk topology | 3-layer Linear→ELU→Linear→ELU→Linear (~150K params) | Capacity sufficient for K=2 prediction; GRN gating is Q-specific | +| Aux trunk input | Encoder output (post-attention, pre-GRN) | Reused; no need to re-encode raw features | +| Encoder backward from aux | **Stop-gradient** at encoder boundary | Encoder remains Q-shaped per today's stop-grad philosophy | +| Coupling A (forward feature wire `dir_concat_qaux`) | KEEP unchanged | Aux-as-feature-for-Q is the actual benefit | +| Coupling B (α-gated backward `sp14_scale_wire_col`) | DELETE atomically | Q co-training aux's head defeats "independent predictor" premise | +| α machinery (`alpha_grad_compute_kernel` + variance EMAs + Schmitt gate) | DELETE atomically with Coupling B | No remaining consumer after Coupling B deletion | +| Q-disagreement signal `ISV[383..390)` | KEEP | Useful diagnostic; producer remains active for HEALTH_DIAG | +| ISV role for aux trunk | Control plane only (hyperparams, adaptive bounds, diagnostic outputs) | NOT data plane — per-step broadcast violates per-env granularity | +| Aux loss weight (`aux_w`) | Keep ISV-driven (already exists) | No change | +| New aux trunk hyperparam slots | 6 ISV slots (LR, β1, β2, ε, grad_clip, h_s2_aux_rms_ema) | Mirrors SP5 Pearl 4 Adam pattern + GRN's h_s2_rms_ema | + +## File Structure + +**Create:** +- `crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu` — 3-layer MLP forward +- `crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu` — 3-layer MLP backward (stop-grad on encoder) +- `crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu` — adaptive scale producer for aux trunk output magnitude +- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — Rust wrapper (forward op + backward op + Adam state holders) +- `crates/ml/tests/aux_trunk_oracle_tests.rs` — gradient check + forward output check + stop-grad invariant + +**Modify:** +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — add 6 new slots (`AUX_TRUNK_LR_INDEX`, etc.) + mark α slots as deprecated +- `crates/ml/build.rs` — register new cubins +- `crates/ml/src/cuda_pipeline/mod.rs` — module declarations +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — aux trunk param tensors + Adam state allocation in trainer constructor; delete α launch sites +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — aux trunk forward + backward in collector chain; delete α launch sites; aux head now reads `h_s2_aux` instead of `h_s2` +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — register new aux trunk slots; remove α slots +- `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` — aux head input switches from `h_s2` to `h_s2_aux` (parameter rename + dimension verify) +- `docs/dqn-wire-up-audit.md` — close-out entry + +**Delete:** +- `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` +- `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu` +- `crates/ml/src/cuda_pipeline/q_disagreement_variance_kernel.cu` (if separate file; if inlined into alpha_grad_compute, deletion is in C.7) +- All α-related ISV slot constants in `sp14_isv_slots.rs` (`VAR_AUX_INDEX`, `VAR_Q_INDEX`, `VAR_ALPHA_INDEX`, `ALPHA_RAW_INDEX`, `ALPHA_SMOOTHED_INDEX`, `GATE1_OPEN_STATE_INDEX`, `GATE2_OPEN_STATE_INDEX`) + +--- + +## Phase C.0: Pre-flight Audit + +### Task C.0.1: Verify ISV slot vacancy + EGF wire-up snapshot + +**Files:** +- Read: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` +- Read: `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` +- Read: `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` +- Read: `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu` +- Modify: `docs/dqn-wire-up-audit.md` (audit-only entry, no code change) + +- [ ] **Step 1: Confirm ISV vacancy after slot 443 (SP15 took [397..443))** + +Run: `grep -nE "INDEX:.*[0-9]+ as usize" crates/ml/src/cuda_pipeline/sp15_isv_slots.rs | tail -5 && grep -nE "ISV_TOTAL_DIM" crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` + +Expected: max slot < 443; `ISV_TOTAL_DIM = 443`. + +- [ ] **Step 2: Identify exact α-machinery launch sites** + +Run: `grep -nE "alpha_grad_compute|scale_wire_col|VARIANCE_REF_AUX|SCHMITT_BAND|gate1_open|gate2_open" crates/ml/src/cuda_pipeline/*.rs crates/ml/src/cuda_pipeline/*.cu crates/ml/src/trainers/dqn/**/*.rs 2>/dev/null` + +Capture line numbers. Save list inline in the audit doc. + +- [ ] **Step 3: Add audit doc entry** + +Append to `docs/dqn-wire-up-audit.md` top: +- Date 2026-05-07 +- Title: "SP14 Layer C pre-flight — α machinery deletion targets" +- List of files + line numbers from Step 2 +- Note: deletion is atomic in Phase C.7; pre-flight establishes "before" state for review + +- [ ] **Step 4: Commit** + +```bash +git add docs/dqn-wire-up-audit.md +git commit -m "docs(sp14-c-preflight): catalogue α-machinery launch sites pending atomic deletion" +``` + +--- + +## Phase C.1: ATOMIC — Delete α Machinery + Add Aux Trunk ISV Slots + +> **CRITICAL: this is ONE atomic commit per `feedback_no_partial_refactor` and `feedback_no_legacy_aliases`** — no DEPRECATED stage. α machinery deletes simultaneously with new slot allocation. Build must be clean before AND after this commit. + +### Task C.1.1: Atomic α deletion + 6 new aux trunk ISV slots + +**Files** (per C.0 audit `a7a162d1f` — full launch site catalogue): +- Delete: `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` +- Delete: `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu` +- Delete: `crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu` (EGF circuit breaker — reads α-machinery products; with EGF gone, breaker has no purpose) +- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (add 6 new + delete 9+ α slots) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (9 launch sites) +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (21 reference sites — kernel fields + launcher fns + load sites) +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (4 sites) +- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs` (1 site) +- Modify: `crates/ml/src/trainers/dqn/batched_backward.rs` (1 site) +- Modify: `crates/ml/build.rs` (remove cubin entries) +- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (remove module decls) +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (remove α slot reset entries; add new aux trunk slot entries) +- Modify: `crates/ml/tests/sp14_oracle_tests.rs` (delete tests for deleted machinery: `alpha_grad_adaptive_beta`, `alpha_grad_schmitt_hysteresis`, `gradient_hack_circuit_breaker_fires`. Preserve `dir_concat_qaux_correct` and `q_disagreement_*`.) +- Modify: `docs/dqn-wire-up-audit.md` (atomic deletion entry) + +- [ ] **Step 1: Add 6 new slots after current SP15 max** + +C.0 audit (commit `a7a162d1f`) confirmed `ISV_TOTAL_DIM = 444` (SP15 1.3.b-followup occupied 443). New slots start at **444**, not 443. + +In `sp14_isv_slots.rs`, append after the SP15 block: + +```rust +// SP14 Layer C — separate aux trunk control plane (added 2026-05-08) +pub const AUX_TRUNK_LR_INDEX: usize = 444; // Adam LR for aux trunk params +pub const AUX_TRUNK_BETA1_INDEX: usize = 445; // Adam β1 +pub const AUX_TRUNK_BETA2_INDEX: usize = 446; // Adam β2 +pub const AUX_TRUNK_EPS_INDEX: usize = 447; // Adam ε +pub const AUX_TRUNK_GRAD_CLIP_INDEX: usize = 448; // grad-clip threshold +pub const H_S2_AUX_RMS_EMA_INDEX: usize = 449; // aux trunk output magnitude EMA (mirrors H_S2_RMS_EMA_INDEX) +``` + +- [ ] **Step 2: Bump ISV_TOTAL_DIM** to `450` (was 444). + +- [ ] **Step 3: DELETE α-related slot constants** (no deprecation step — atomic with kernel deletion in Step 7) + +C.0 audit identified **9** α-machinery slots. Real names: +- `K_AUX_ADAPTIVE_INDEX` (385) — written by alpha_grad_compute +- `K_Q_ADAPTIVE_INDEX` (386) — written by alpha_grad_compute +- `BETA_RATE_LIMITER_ADAPTIVE_INDEX` (387) — written by alpha_grad_compute +- `ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX` +- `GATE1_OPEN_STATE_INDEX` (single Schmitt gate; no GATE2) +- `ALPHA_GRAD_RAW_INDEX` +- `ALPHA_SMOOTHED_INDEX` +- Plus any var_aux/var_q-equivalent constants in the file (verify exact names via `grep -nE "VAR_AUX|VAR_Q|AUX_DIR_ACC_POST_OPEN_MIN" crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — use whatever the file actually has). +- `AUX_DIR_ACC_POST_OPEN_MIN_INDEX` if present (read by gradient_hack_detect_kernel). + +Delete the constant definitions outright. Leave the freed indices as RESERVED gap (do NOT compact — preserves checkpoint layout fingerprint compatibility). + +- [ ] **Step 3a: Delete α kernel files** + +```bash +git rm crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu +git rm crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu +git rm crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu +``` + +- [ ] **Step 3b: Remove from build.rs cubin manifest** + +Search build.rs for `alpha_grad_compute`, `sp14_scale_wire_col`, `gradient_hack_detect` — delete those compilation entries. + +- [ ] **Step 3c: Remove module decls** in `mod.rs`. + +- [ ] **Step 3d: Remove all launch / consumer sites** (use C.0 audit `a7a162d1f` as the work list — 31 reference lines across 7 files): + - `gpu_experience_collector.rs` (9 sites): remove `launch_alpha_grad_compute(...)`, `launch_sp14_scale_wire_col(...)`, and any `launch_gradient_hack_detect(...)` calls + - `gpu_dqn_trainer.rs` (21 sites): remove `alpha_grad_compute_kernel: CudaFunction`, `sp14_scale_wire_col_kernel: CudaFunction`, `gradient_hack_detect_kernel: CudaFunction` fields; remove their loads in `compile_training_kernels`; remove the launcher fns + - `training_loop.rs` (4 sites): remove any `gate1_open` / `α_smoothed` reads or HEALTH_DIAG entries + - `trainer/mod.rs` (1 site) + - `batched_backward.rs` (1 site) + - `sp14_isv_slots.rs` (3 sites — beyond the 9 slot constants, also delete any helper fns / re-exports) + +- [ ] **Step 3e: Delete obsolete oracle tests** + +In `crates/ml/tests/sp14_oracle_tests.rs`, delete: +- `gpu::alpha_grad_adaptive_beta` +- `gpu::alpha_grad_schmitt_hysteresis` +- `gpu::gradient_hack_circuit_breaker_fires` + +Preserve: +- `gpu::dir_concat_qaux_correct` (Coupling A — survives) +- `gpu::q_disagreement_*` (3 tests — diagnostic kernel survives) + +- [ ] **Step 3f: Verify q_disagreement signal preserved** + +`q_disagreement_short_ema` and `q_disagreement_long_ema` (`ISV[383..390)`) are NOT deleted — they remain useful diagnostics consumed by HEALTH_DIAG. Verify their producer kernel `q_disagreement_update_kernel.cu` is NOT in the deletion list. + +Run: `ls crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu` — file must still exist. + +- [ ] **Step 4: Add reset_registry entries for new slots** + +In `state_reset_registry.rs`, follow the existing per-slot pattern (find an existing entry like `H_S2_RMS_EMA_INDEX => StateReset::SentinelF32(0.0)` for reference). New entries: + +| Slot | Reset value | Sentinel meaning | +|---|---|---| +| `AUX_TRUNK_LR_INDEX` | `SentinelF32(1e-4)` | Default aux-LR until Pearl-A bootstrap | +| `AUX_TRUNK_BETA1_INDEX` | `SentinelF32(0.9)` | Adam default | +| `AUX_TRUNK_BETA2_INDEX` | `SentinelF32(0.999)` | Adam default | +| `AUX_TRUNK_EPS_INDEX` | `SentinelF32(1e-8)` | Adam default | +| `AUX_TRUNK_GRAD_CLIP_INDEX` | `SentinelF32(1.0)` | Default norm-clip | +| `H_S2_AUX_RMS_EMA_INDEX` | `SentinelF32(0.0)` | Pearl-A first-observation bootstrap (replace, not blend) | + +- [ ] **Step 4b: Add dispatch arm for `reset_named_state`** if the registry has a parallel by-name dispatch (mirrors SP5 Layer A bug-fix pattern). + +- [ ] **Step 4c: Remove reset_registry entries for the 9 deleted α slots** (atomic with Step 3 deletion). + +- [ ] **Step 4d: Update audit doc** — append entry to `docs/dqn-wire-up-audit.md`: + - Title: "SP14 Layer C — α machinery deleted atomically (Phase C.1)" + - List of deleted files (3 kernel files) + 9 deleted slot constants + - Reference to C.0 pre-flight audit + commit SHA `a7a162d1f` + - Note: `q_disagreement_*` slots preserved as diagnostic-only + +- [ ] **Step 5: Verify** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -15` + +Expected: **clean build, all targets**. Compile errors here mean a launch site or import was missed in Step 3d — fix before commit. Build must compile because: +- All α launches deleted simultaneously with the kernels they call +- All α slot reads deleted simultaneously with the slot constants +- Aux head's input still flows from h_s2 (Q's trunk) — that's fine; it just means aux head has slightly different gradient flow now (Q-loss can pull on aux head freely without α scaling). Aux head behavior will be replaced in Phase C.5 anyway. + +Run remaining oracle tests to confirm nothing else broke: +`SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15` + +Expected: 4 tests pass (dir_concat_qaux_correct + 3 q_disagreement_*). Three deleted tests no longer exist. + +- [ ] **Step 6: Commit** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +refactor(sp14-c): atomic α machinery deletion + aux trunk ISV slot allocation + +Phase C.1 of SP14 Layer C separate-aux-trunk refactor. Single atomic +commit per feedback_no_partial_refactor and feedback_no_legacy_aliases — +no DEPRECATED stage. + +Deleted (per C.0 audit a7a162d1f): +- alpha_grad_compute_kernel.cu +- sp14_scale_wire_col_kernel.cu +- gradient_hack_detect_kernel.cu (EGF circuit breaker — α-coupled) +- 9 ISV slot constants (K_AUX_ADAPTIVE, K_Q_ADAPTIVE, + BETA_RATE_LIMITER_ADAPTIVE, ALPHA_GRAD_RAW_VARIANCE_EMA, + GATE1_OPEN_STATE, ALPHA_GRAD_RAW, ALPHA_SMOOTHED, + AUX_DIR_ACC_POST_OPEN_MIN, plus var_aux/var_q if present) +- 31 reference sites across 7 files (collector, trainer, training_loop, + batched_backward, mod.rs, build.rs cubin manifest) +- 3 oracle tests (alpha_grad_adaptive_beta, alpha_grad_schmitt_hysteresis, + gradient_hack_circuit_breaker_fires) + +Added: +- 6 aux trunk control plane ISV slots [444..450): + AUX_TRUNK_LR (444), AUX_TRUNK_BETA1 (445), AUX_TRUNK_BETA2 (446), + AUX_TRUNK_EPS (447), AUX_TRUNK_GRAD_CLIP (448), H_S2_AUX_RMS_EMA (449) +- Reset registry entries (Pearl-A first-observation bootstrap) +- ISV_TOTAL_DIM bumped 444 → 450 + +Preserved: +- dir_concat_qaux_kernel.cu (Coupling A: forward feature wire) +- q_disagreement_update_kernel.cu (diagnostic-only) +- q_disagreement_* ISV slots [383..390) (HEALTH_DIAG consumer) + +Build clean; 4 surviving oracle tests pass. ISV layout: α slots left as +RESERVED gap for checkpoint fingerprint compatibility. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Phase C.2: Aux Trunk Parameter Tensors + Adam State + +### Task C.2.1: Allocate aux trunk params + Adam m/v in trainer constructor + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` + +**Topology** (sized to ~150K params total): + +| Layer | Shape | Init | +|---|---|---| +| `aux_trunk_w1` | `[encoder_out_dim, 256]` | Kaiming-He fan_in=encoder_out_dim | +| `aux_trunk_b1` | `[256]` | zeros | +| `aux_trunk_w2` | `[256, 128]` | Kaiming-He fan_in=256 | +| `aux_trunk_b2` | `[128]` | zeros | +| `aux_trunk_w3` | `[128, AUX_HIDDEN_DIM]` (= existing aux head input dim) | Kaiming-He fan_in=128 | +| `aux_trunk_b3` | `[AUX_HIDDEN_DIM]` | zeros | + +`AUX_HIDDEN_DIM` must match the aux head's expected input dimension. Read `aux_heads_kernel.cu` to find the constant (likely `SH2` or a separate aux-specific constant). If aux head currently consumes `SH2`, then `AUX_HIDDEN_DIM = SH2`. + +- [ ] **Step 1: Add struct fields to the trainer struct** (`GpuDqnTrainer` in `gpu_dqn_trainer.rs`) + +Find an existing param-group struct (e.g., GRN trunk params at the established location — grep for `grn_trunk_w1` or `aux_w1`). Mirror the pattern. Each tensor is a `cudarc::driver::CudaSlice`. Add 6 param tensors + 12 Adam state tensors (m, v for each). + +- [ ] **Step 2: Allocation in `compile_training_kernels` or wherever existing params are allocated** + +Look for the existing pattern (e.g., `let aux_w1 = stream.alloc_zeros::(...)?`). Allocate aux_trunk_w1/b1/w2/b2/w3/b3 + their Adam m/v. Store in trainer struct fields. + +- [ ] **Step 3: Kaiming-He init via cuRAND** + +Find existing init pattern (e.g., GRN init or VSN init). Kaiming-He fan_in formula: `std = sqrt(2.0 / fan_in)`. Apply to each weight tensor. Biases zero-init (already done by `alloc_zeros`). + +- [ ] **Step 4: Mirror in collector** if the collector independently holds copies (per β migration pattern). Check `gpu_experience_collector.rs` for a parallel `aux_trunk_w1` field — if present, the trainer-side allocation must propagate (use `clone_async` or share the underlying `CudaSlice` Arc). + +- [ ] **Step 5: Verify** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --lib --tests 2>&1 | tail -10` + +Expected: clean build. Param tensors allocated but not yet wired to forward/backward. + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +git commit -m "feat(sp14-c): allocate aux trunk params (~150K) + Adam m/v state + +3-layer MLP: encoder_out_dim → 256 → 128 → AUX_HIDDEN_DIM. Kaiming-He +weights, zero biases, separate Adam m/v buffers. Allocated in trainer +constructor + collector mirror per β-migration pattern. Not yet wired +to forward/backward — pure allocation." +``` + +--- + +## Phase C.3: Aux Trunk Forward Kernel + Oracle Test + +### Task C.3.1: Forward kernel `aux_trunk_forward_kernel.cu` + +**Files:** +- Create: `crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu` +- Modify: `crates/ml/build.rs` (register cubin) +- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (module decl) +- Create: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (Rust wrapper, forward only — backward in C.4) + +**Kernel signature** (mirrors `aux_next_bar_forward` in `aux_heads_kernel.cu`): + +```c +extern "C" __global__ void aux_trunk_forward( + const float* __restrict__ x_in, // [B, encoder_out_dim] post-encoder + const float* __restrict__ w1, // [encoder_out_dim, 256] + const float* __restrict__ b1, // [256] + const float* __restrict__ w2, // [256, 128] + const float* __restrict__ b2, // [128] + const float* __restrict__ w3, // [128, AUX_HIDDEN_DIM] + const float* __restrict__ b3, // [AUX_HIDDEN_DIM] + float* __restrict__ h_aux1_out, // [B, 256] saved for backward + float* __restrict__ h_aux2_out, // [B, 128] saved for backward + float* __restrict__ h_s2_aux_out, // [B, AUX_HIDDEN_DIM] final aux trunk output + int B, + int ENCODER_OUT_DIM, + int AUX_HIDDEN_DIM +); +``` + +Body: standard 3-layer MLP. Linear → ELU → Linear → ELU → Linear (no activation on final output — let aux head's own activation handle it). Use shared memory for intermediate hidden vectors per existing aux_heads kernel pattern. + +**Reference:** `aux_heads_kernel.cu:117-189` for the 2-layer pattern; extend to 3 layers. + +- [ ] **Step 1: Write the kernel file** + +Follow `aux_heads_kernel.cu`'s style: thread-block per batch row, threadIdx.x sweeps the output dimension. ELU device function (already exists at `aux_heads_kernel.cu:38-44` — use the same). + +- [ ] **Step 2: Register cubin in `build.rs`** + +Find the existing pattern (grep `aux_next_bar` in `build.rs`). Add `aux_trunk_forward` to the cubin compilation list. + +- [ ] **Step 3: Add module decl in `mod.rs`** (if not auto-discovered). + +- [ ] **Step 4: Write Rust wrapper `gpu_aux_trunk.rs`** + +Pattern: `pub struct AuxTrunkForwardOps { kernel: CudaFunction, ... }`. `pub fn launch(stream, x_in, w1, b1, w2, b2, w3, b3, scratch_h_aux1, scratch_h_aux2, h_s2_aux_out, B, encoder_out_dim, aux_hidden_dim) -> Result<(), MLError>`. Mirror `gpu_aux_heads.rs` if exists, else `gpu_grn.rs`. + +**Critical:** pre-load `CudaFunction` once at construction; do NOT `load_cubin` in the launch hot path (per `pearl_no_host_branches_in_captured_graph`). + +- [ ] **Step 5: Write oracle test (forward output check)** + +Create `crates/ml/tests/aux_trunk_oracle_tests.rs`: + +```rust +//! Oracle tests for aux trunk forward+backward. + +use cudarc::driver::*; +use ml::cuda_pipeline::gpu_aux_trunk::AuxTrunkForwardOps; + +#[test] +#[ignore] // GPU test +fn aux_trunk_forward_matches_numpy_reference() { + // Setup: deterministic seed, B=4, enc=64, h1=256, h2=128, aux=128 + // 1. Allocate input + weights with fixed values + // 2. Launch aux_trunk_forward + // 3. Compute reference on host: h_aux1 = elu(x @ w1 + b1); h_aux2 = elu(h_aux1 @ w2 + b2); out = h_aux2 @ w3 + b3 + // 4. Compare element-wise; tol 1e-4 (f32 SGEMM precision) +} +``` + +Full body using `cudarc::driver::CudaContext` plus `ndarray` for the host reference. + +- [ ] **Step 6: Run test, verify pass** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored aux_trunk_forward_matches_numpy_reference --nocapture 2>&1 | tail -10` + +Expected: PASS. + +- [ ] **Step 7: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/src/cuda_pipeline/mod.rs crates/ml/build.rs crates/ml/tests/aux_trunk_oracle_tests.rs +git commit -m "feat(sp14-c): aux trunk forward kernel + Rust wrapper + oracle test + +3-layer MLP forward (Linear→ELU→Linear→ELU→Linear). Pre-loaded +CudaFunction for graph-capture safety. Oracle test verifies bit-for-bit +match against numpy reference within 1e-4 tol." +``` + +--- + +## Phase C.4: Aux Trunk Backward Kernel + Oracle Test + +### Task C.4.1: Backward kernel + stop-grad on encoder boundary + +**Files:** +- Create: `crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu` +- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add backward op) +- Modify: `crates/ml/src/cuda_pipeline/mod.rs` +- Modify: `crates/ml/build.rs` +- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (add gradient-check test) + +**Kernel signature:** + +```c +extern "C" __global__ void aux_trunk_backward( + const float* __restrict__ dh_s2_aux_in, // [B, AUX_HIDDEN_DIM] gradient from aux head + const float* __restrict__ x_in, // [B, ENCODER_OUT_DIM] saved fwd input + const float* __restrict__ h_aux1, // [B, 256] saved fwd + const float* __restrict__ h_aux2, // [B, 128] saved fwd + const float* __restrict__ w1, // [ENCODER_OUT_DIM, 256] + const float* __restrict__ w2, // [256, 128] + const float* __restrict__ w3, // [128, AUX_HIDDEN_DIM] + float* __restrict__ dW1, // [ENCODER_OUT_DIM, 256] grad accum + float* __restrict__ db1, // [256] + float* __restrict__ dW2, // [256, 128] + float* __restrict__ db2, // [128] + float* __restrict__ dW3, // [128, AUX_HIDDEN_DIM] + float* __restrict__ db3, // [AUX_HIDDEN_DIM] + int B, + int ENCODER_OUT_DIM, + int AUX_HIDDEN_DIM +); +``` + +**Critical: NO `dx_in_out` parameter.** The kernel does NOT compute or write a gradient back to `x_in` (the encoder output). This is the stop-gradient at the encoder boundary — aux trunk's gradient terminates here. + +**ELU derivative reminder:** `d_elu(x) = 1 if x > 0 else exp(x)` (continuous version: `elu(x) + 1 if x ≤ 0 else 1`). + +**Reference:** `aux_heads_kernel.cu:481-614` for the 2-layer backward pattern; extend to 3 layers, omit dx propagation. + +- [ ] **Step 1: Write backward kernel** + +Three steps in reverse order: +1. d_logits = dh_s2_aux_in directly (output is linear, no activation) +2. dh_aux2 = (d_logits @ w3.T) ⊙ d_elu(h_aux2 pre-activation); accumulate dW3, db3 +3. dh_aux1 = (dh_aux2 @ w2.T) ⊙ d_elu(h_aux1 pre-activation); accumulate dW2, db2 +4. (NO Step 4 dx propagation — stop at w1 grad accumulation): accumulate dW1, db1 + +`atomicAdd` on dW gradients is forbidden per `feedback_no_atomicadd`. Use block-tree-reduce per `aux_heads_kernel.cu` precedent (lines ~520-570 in `aux_next_bar_backward`). + +- [ ] **Step 2: Add backward op to Rust wrapper** + +In `gpu_aux_trunk.rs`, add `pub struct AuxTrunkBackwardOps { kernel: CudaFunction, ... }` with `launch(...)`. Pre-load CudaFunction. + +- [ ] **Step 3: Register cubin** in `build.rs`. + +- [ ] **Step 4: Write gradient-check oracle test** + +```rust +#[test] +#[ignore] +fn aux_trunk_backward_gradient_check() { + // Setup deterministic input + weights + // 1. Forward to get h_s2_aux + // 2. Compute analytic gradient via aux_trunk_backward (using identity dh_s2_aux for clarity) + // 3. Compute numerical gradient via finite differences: + // grad[i,j] ≈ (loss(W + ε·e_ij) - loss(W - ε·e_ij)) / (2ε) + // where loss = sum(h_s2_aux²)/2 (so dloss/dh_s2_aux = h_s2_aux directly) + // 4. Compare analytic vs numerical for each of dW1/dW2/dW3, db1/db2/db3 + // 5. Tolerance: 1e-3 (numerical gradient noisy in f32) +} +``` + +- [ ] **Step 5: Write stop-grad invariant test** + +```rust +#[test] +#[ignore] +fn aux_trunk_backward_does_not_write_dx() { + // Allocate dx_in buffer with known sentinel pattern (e.g., 0xDEADBEEF / -1.0) + // Call aux_trunk_backward + // Verify dx_in unchanged — kernel must NOT touch encoder gradient buffer +} +``` + +This is critical: encoder gradient must remain Q-shaped only. + +- [ ] **Step 6: Run tests** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15` + +Expected: 3 tests pass (forward + gradient check + stop-grad invariant). + +- [ ] **Step 7: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/build.rs crates/ml/tests/aux_trunk_oracle_tests.rs +git commit -m "feat(sp14-c): aux trunk backward kernel + gradient check + stop-grad test + +Backward propagates dh_s2_aux through w3/w2/w1 with block-tree-reduce +(no atomicAdd per feedback_no_atomicadd). Critical: kernel does NOT +write dx_in — encoder gradient remains Q-shaped only. Stop-grad +invariant verified by sentinel-pattern test." +``` + +--- + +## Phase C.4b: ISV-Driven Aux Prediction Horizon (multi-bar pivot) + +> **Added 2026-05-08 per user design correction.** Aux's original label `(p_{t+1} > p_t)` from `aux_sign_label_kernel.cu` predicts next-bar at HFT scale — pure microstructure noise. User trades at HFT-MFT scale (multi-bar holds), so aux must predict over a horizon matching actual hold times. Without this, aux flatlines at ln(2) regardless of trunk separation: the *target itself* is unlearnable. + +> **Scope:** label production + 1 new ISV slot + adaptive producer. Trunk math (C.2/C.3/C.4) is label-agnostic and unchanged. Cold-start sentinel H=60 (≈1hr @ 1min); first winning-trade observation triggers Pearl-A bootstrap (replace, not blend); thereafter Wiener-α EMA tracks avg winning hold time. + +### Task C.4b.1: Migrate aux label kernels to ISV-driven H-bar horizon + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (add 1 new slot, bump ISV_TOTAL_DIM 450→451; possibly add 1-2 hold-time tracker slots if not already present) +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (Pearl-A bootstrap entry) +- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu` (use H from ISV, mask `t+H >= total_bars`) +- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu` (same) +- Create: `crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu` (adaptive producer reading avg winning hold time, writing to AUX_PRED_HORIZON_BARS_INDEX) +- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add launcher op for producer) +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (kernel field + load + launch site) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (per-epoch boundary launch) +- Modify: `crates/ml/build.rs` (cubin entry) +- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (module decl) +- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (4 tests: H-bar correctness, lookahead mask, Pearl-A bootstrap, Wiener-α EMA convergence) +- Modify: `docs/dqn-wire-up-audit.md` (Phase C.4b entry) + +- [ ] **Step 1: Add ISV slot** + +In `sp14_isv_slots.rs`: + +```rust +// SP14 Layer C — aux prediction horizon (added 2026-05-08, multi-bar pivot) +pub const AUX_PRED_HORIZON_BARS_INDEX: usize = 450; +``` + +Bump `ISV_TOTAL_DIM` from 450 to **451**. + +- [ ] **Step 2: Reset registry entry** + +In `state_reset_registry.rs`: + +```rust +AUX_PRED_HORIZON_BARS_INDEX => StateReset::SentinelF32(60.0), +``` + +Sentinel 60.0 = 60 bars ≈ 1 hour @ 1-min bars. This is the cold-start default. Pearl-A bootstrap will replace on first observation when adaptive producer lands (deferred follow-up). + +Add the dispatch arm to `reset_named_state` if applicable. + +- [ ] **Step 3: Modify `aux_sign_label_kernel.cu`** + +Find the current label computation (likely a single line: `out_labels[i] = (p_fut > p_now) ? 1 : 0`). Replace with: + +```c +extern "C" __global__ void aux_sign_label( + const float* __restrict__ prices, // [B*T, 1] price column from fxcache + const float* __restrict__ isv, // [ISV_TOTAL_DIM] + int* __restrict__ out_labels, // [B*T] -1 sentinel for masked + int total_bars, + int B, + int T +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= B * T) return; + + int b = idx / T; + int t = idx % T; + + int H = (int)__float2int_rn(isv[AUX_PRED_HORIZON_BARS_INDEX]); + if (H < 1) H = 1; // floor: can't predict past + if (H > 240) H = 240; // ceiling: rollout buffer guard + + int t_future = t + H; + if (t_future >= total_bars) { + // Lookahead mask: no label available for this position + out_labels[idx] = -1; // sentinel for "ignore in loss" + return; + } + + float p_now = prices[b * total_bars + t]; + float p_fut = prices[b * total_bars + t_future]; + out_labels[idx] = (p_fut > p_now) ? 1 : 0; +} +``` + +Adjust signature to match the EXISTING kernel's parameter list. The key changes: +- Read `H` from ISV (read once per thread, fine — broadcast value) +- Index `t_future = t + H` instead of `t + 1` +- Mask `out_labels[idx] = -1` when `t_future >= total_bars` + +**Critical:** the loss reduction kernel (`aux_next_bar_loss_reduce` in `aux_heads_kernel.cu`) must be checked — if it reads label as `int` and casts to softmax index, a `-1` sentinel will go out of bounds. Verify the consumer handles `-1` (skip the sample) or change the masking convention. The existing `out_labels` is `i32` per SP13 B1.1a (K=2 softmax CE). + +If the loss-reduce kernel does NOT support a -1 mask, add masking handling: +- Easier alternative: instead of -1 mask, FORWARD-FILL with a deterministic but safe label (e.g., 0). The loss term will be incorrect but bounded; lookahead loss is ~6% of training data so signal is preserved. +- OR: add a `valid_mask` parameter that the loss-reduce kernel multiplies by. + +Implementer choice: pick the simpler one given the existing kernel structure. Prefer the valid_mask approach (correctness > simplicity here since we're already touching label production). + +- [ ] **Step 4: Modify `aux_sign_label_per_step_kernel.cu`** + +Same migration. The per-step kernel runs at every rollout step (not at end-of-rollout); H is read from ISV at that step. Per-step kernel has access to forward bars in the rollout buffer; mask if `t + H` exceeds rollout length. + +- [ ] **Step 5: Update launch sites** + +Find launch sites in `gpu_dqn_trainer.rs` and `gpu_experience_collector.rs`. If they don't already pass the ISV pointer to the label kernel, add it. ISV ptr is already passed to dozens of other kernels — find an example pattern (`launch_aux_horizon_compute` or similar from neighboring launches) and mirror. + +- [ ] **Step 5b: Identify avg winning hold time source in existing infrastructure** + +Run: +```bash +grep -nE "winning_hold|win_hold|hold_time|HOLD_TIME|trade_close|trade_dur|HOLDING|holding_bars" \ + crates/ml/src/cuda_pipeline/*.rs crates/ml/src/cuda_pipeline/*.cu \ + crates/ml/src/cuda_pipeline/sp14_isv_slots.rs \ + crates/ml/src/cuda_pipeline/sp15_isv_slots.rs 2>/dev/null +``` + +SP12/SP15 introduced trade-close events (per-trade reward emission). Expected to find: per-env "ticks since position opened" tracker + on-close emission of duration to a hold-time EMA. + +Three possible findings: + +**Case A: Existing slot tracks avg winning hold time directly** (e.g., `AVG_WIN_HOLD_TIME_EMA_INDEX`). +→ Producer reads from that slot. Done. + +**Case B: Existing infrastructure tracks per-env open-tick + emits on close, but no aggregate slot.** +→ Add a small aggregate kernel: averaged-on-close winning hold time → new ISV slot `AVG_WIN_HOLD_TIME_BARS_INDEX = 451`. Producer reads from that. +→ Bumps ISV_TOTAL_DIM 451→452. Document. + +**Case C: No existing hold-time infrastructure.** +→ Out of scope for this phase; report BLOCKED. (We don't want to introduce trade-close tracking infrastructure here — that's SP15 territory.) + +Document which case you found in the audit doc. + +- [ ] **Step 5c: Write `aux_horizon_update_kernel.cu` (adaptive producer)** + +```c +extern "C" __global__ void aux_horizon_update( + float* __restrict__ isv, // [ISV_TOTAL_DIM] + int isv_h_idx, // AUX_PRED_HORIZON_BARS_INDEX = 450 + int isv_h_target_idx, // index of avg winning hold time slot (Case A or B above) + int isv_h_target_var_idx, // optional: variance EMA of target for Wiener-α (-1 if not available) + float h_min, // 1.0 + float h_max, // 240.0 + float sentinel_h // 60.0 (cold-start default) +) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; // single-thread compute + + float h_target = isv[isv_h_target_idx]; + float h_current = isv[isv_h_idx]; + + // No winning trade observed yet — keep sentinel + if (h_target <= 0.0f) { + isv[isv_h_idx] = sentinel_h; // keep sentinel + return; + } + + // Pearl-A first-observation bootstrap: replace sentinel directly + if (fabsf(h_current - sentinel_h) < 1e-6f) { + // First valid observation: replace sentinel directly + float h_clamped = fmaxf(h_min, fminf(h_max, h_target)); + isv[isv_h_idx] = h_clamped; + return; + } + + // Steady-state Wiener-α EMA blend + float alpha; + if (isv_h_target_var_idx >= 0) { + // Wiener-optimal: α = sample_var / (sample_var + diff_var + ε) + // For horizon (slow signal): use signal-derived α with floor for stability + float sample_var = isv[isv_h_target_var_idx]; + float diff = h_target - h_current; + float diff_var = diff * diff; // single-sample approximation + alpha = sample_var / (sample_var + diff_var + 1e-6f); + alpha = fmaxf(0.001f, fminf(0.1f, alpha)); // bounds: slow horizon + } else { + alpha = 0.01f; // fixed slow blend for stable horizon + } + + float h_blended = (1.0f - alpha) * h_current + alpha * h_target; + isv[isv_h_idx] = fmaxf(h_min, fminf(h_max, h_blended)); +} +``` + +Single-thread kernel (one block, one thread) — compute is trivial; just need to write a single ISV slot. No atomicity concerns. + +Per `pearl_first_observation_bootstrap`: sentinel = 60.0; first observation replaces directly (not blends). + +Per `pearl_wiener_optimal_adaptive_alpha`: if target variance is available, use Wiener-α; else fixed 0.01 (slow blend). + +Per `feedback_isv_for_adaptive_bounds`: bounds h_min=1, h_max=240 are constants (acceptable — fundamental floors/ceilings, not tuning parameters). + +- [ ] **Step 5d: Wire producer** + +In `gpu_aux_trunk.rs`, add `AuxHorizonUpdateOps` struct mirroring forward/backward op pattern. Pre-load CudaFunction. Launch sig matches kernel. + +In `gpu_experience_collector.rs`, add launch at **per-epoch boundary** (horizon is slow-moving — epoch-level cadence is appropriate, not per-step). Find the existing per-epoch ISV update block (mirrors SP4 producer launches at end of epoch). + +In `gpu_dqn_trainer.rs`, add `aux_horizon_update_ops: AuxHorizonUpdateOps` field; load in trainer constructor. + +- [ ] **Step 6: Add label correctness oracle test** + +In `aux_trunk_oracle_tests.rs`: + +```rust +#[test] +#[ignore] +fn aux_sign_label_h_bar_horizon() { + // 1. Synthetic price series: linearly increasing for B=2, T=200 + // 2. Set ISV[AUX_PRED_HORIZON_BARS_INDEX] = 30.0 + // 3. Launch aux_sign_label + // 4. Verify out_labels[0..170] == 1 (price 30 bars ahead is higher) + // 5. Verify out_labels[170..200] == -1 (lookahead masked) + + // Then with descending series: out_labels[0..170] == 0 + + // Then with sine wave: verify labels match analytic (p(t+30) > p(t)) +} +``` + +- [ ] **Step 7: Add lookahead mask invariant test** + +```rust +#[test] +#[ignore] +fn aux_sign_label_lookahead_mask() { + // Set H = 60. Set total_bars = 100. + // Launch. + // Verify out_labels[40..100] == -1 (cannot look 60 bars past total_bars - 60 = 40) + // Verify out_labels[0..40] are valid (∈ {0, 1}) +} +``` + +- [ ] **Step 7b: Pearl-A bootstrap test for horizon producer** + +```rust +#[test] +#[ignore] +fn aux_horizon_pearl_a_bootstrap() { + // 1. ISV[AUX_PRED_HORIZON_BARS_INDEX] = 60.0 (sentinel) + // 2. ISV[h_target_idx] = 90.0 (first observation) + // 3. Launch aux_horizon_update + // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] == 90.0 (REPLACE, not blend) + + // 5. ISV[h_target_idx] = 100.0 (subsequent observation) + // 6. Launch + // 7. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] in (90, 100) — blended, not replaced +} +``` + +- [ ] **Step 7c: Wiener-α convergence test** + +```rust +#[test] +#[ignore] +fn aux_horizon_converges_to_steady_target() { + // 1. Bootstrap ISV[AUX_PRED_HORIZON_BARS_INDEX] = 50.0 (post-bootstrap state) + // 2. ISV[h_target_idx] = 100.0 (steady target) + // 3. Launch aux_horizon_update 100 times + // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] converges to within 5% of 100.0 + // (with α=0.01 fixed blend, after 100 iters: (1-0.01)^100 ≈ 0.366 + // so distance halves every ~70 iters; 100 iters → ~63% closure → ~78.5) + // 5. Adjust assertion to (1 - (1-α)^100) * 100 + (1-α)^100 * 50 ≈ 81.7 +} +``` + +Test the actual convergence math, not a hand-wavy "approaches 100". + +- [ ] **Step 7d: Cold-start "no winning trades" guard test** + +```rust +#[test] +#[ignore] +fn aux_horizon_holds_sentinel_with_no_winning_trades() { + // 1. ISV[AUX_PRED_HORIZON_BARS_INDEX] = 60.0 (sentinel) + // 2. ISV[h_target_idx] = 0.0 (no winning trades observed yet) + // 3. Launch + // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] == 60.0 (unchanged — kernel guard) +} +``` + +- [ ] **Step 8: Verify** + +```bash +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -10 +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -20 +``` + +Expected: clean build; all aux_trunk tests pass + 2 new label tests pass. + +- [ ] **Step 9: Audit doc** + +Append to `docs/dqn-wire-up-audit.md`: +- Title: "SP14 Layer C Phase C.4b — aux label horizon ISV-driven" +- Date 2026-05-08 +- Rationale: HFT-scale next-bar prediction is unlearnable noise; user trades multi-bar +- Scope: 1 new ISV slot (450), 2 label kernels migrated, lookahead masking +- Cold-start: H=60 bars sentinel, adaptive producer deferred +- Open follow-up: ISV producer for `AUX_PRED_HORIZON_BARS_INDEX` from observed winning hold time (likely sourced from existing trade-close event tracking infrastructure introduced in SP15) + +- [ ] **Step 10: Commit** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +feat(sp14-c): aux prediction horizon ISV-driven (multi-bar pivot) + +Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure +noise that's unlearnable at our HFT-MFT trading frequency. Migrated to +(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX]. + +Adaptive producer drives H from observed avg winning hold time: +- Pearl-A first-observation bootstrap: replace sentinel H=60 directly + on first valid observation +- Steady-state Wiener-α EMA blend, slow (α≈0.01) for stable horizon +- "No winning trades yet" guard keeps sentinel until first valid observation + +Lookahead truncation: labels at t where t+H >= total_bars are masked +(sentinel -1, loss-reduce skips). ~H/T fraction of training data lost +per fold; for H=60, T=1000 → ~6% acceptable. + +Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label- +agnostic. Validation in C.10 will use H=60 cold-start; if aux learns, +adaptive H lands as follow-up. + +Phase C.4b of SP14 Layer C separate-aux-trunk refactor. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +DO NOT push. + +## Hard rules + +1. **`feedback_no_partial_refactor`**: both label kernels (regular + per_step) migrate together. If only one is migrated, validation will be inconsistent (training-time labels vs eval-time labels diverge). +2. **Lookahead masking is correctness-critical**: a label at `t + H >= total_bars` reads garbage memory. Mask MUST be applied or the kernel produces undefined labels. +3. **Loss-reduce consumer**: verify the existing aux loss-reduce kernel handles the chosen mask convention. If `-1` causes OOB, use `valid_mask` parameter approach instead. +4. **No host branches in graph capture**: all logic in the label kernel runs on-device. + +## Falsification gate (post-validation) + +If C.10 L40S validation shows aux loss STILL flatlines at ln(2) with H=60: +- Check label distribution (rolling p(up) — if extremely imbalanced, that's the bug) +- Check encoder output normalization (regime-spike instability) +- Escalate to "aux pulls from PRE-attention encoder layer" or "aux trunk needs LayerNorm" — out of scope for this plan + +--- + +## Phase C.5: Atomic Forward+Backward Wire-Up + +> **CRITICAL: this is ONE atomic commit per `feedback_no_partial_refactor`.** All consumers of aux head's input migrate from `h_s2` to `h_s2_aux` simultaneously with the new kernel launches in collector. If the build is broken between forward-only and backward-wired, the commit is incomplete. + +### Task C.5.1: Wire aux trunk forward + backward into collector chain + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (if backward chain lives here) +- Modify: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (input parameter rename `h_s2` → `h_s2_aux`) +- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs` (Rust wrapper input plumbing) + +- [ ] **Step 1: Locate aux head's current forward call site in collector** + +Run: `grep -nE "aux_next_bar_forward|aux_regime_forward|launch_aux_forward" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` + +Save line numbers. + +- [ ] **Step 2: Locate encoder output buffer** + +The encoder output is the input to GRN trunk. Find it via: `grep -nE "encoder_out|h_s1|grn_input|grn_trunk_forward" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs | head -20` + +Identify the buffer that GRN reads as input. That's also aux trunk's input. + +- [ ] **Step 3: Allocate `h_s2_aux` buffer** (rollout-sized, mirrors aux's existing forward buffers in collector) + +Find where `aux_next_bar_logits` or similar is allocated; allocate `h_s2_aux` same way + scratch buffers for `h_aux1`, `h_aux2` saved-for-backward. + +- [ ] **Step 4: Insert aux trunk forward launch** AFTER encoder forward, BEFORE aux head forward. + +Pseudocode order in the collector forward chain: +``` +1. encoder_forward → encoder_out +2. grn_trunk_forward(encoder_out) → h_s2 // Q's path, unchanged +3. aux_trunk_forward(encoder_out) → h_s2_aux // NEW +4. dir_concat_qaux(h_s2, aux_softmax_diff) → x_concat // Coupling A unchanged, but reads aux's softmax which now comes from h_s2_aux +5. q_head_forward(x_concat) → q_logits +6. aux_head_forward(h_s2_aux) → aux_logits // INPUT CHANGED from h_s2 to h_s2_aux +``` + +- [ ] **Step 5: Update aux head input plumbing** + +In `aux_heads_kernel.cu`, the `aux_next_bar_forward` and `aux_regime_forward` kernels currently take `const float* h_s2` as input. Rename parameter to `h_s2_aux` (semantic rename — interface contract unchanged at call site since both buffers have same shape `[B, AUX_HIDDEN_DIM]`). + +In `gpu_aux_heads.rs` Rust wrapper, update the launch call to pass `h_s2_aux` buffer ptr instead of `h_s2`. + +- [ ] **Step 6: Wire backward chain** + +After aux loss reduce produces `dh_s2_aux` (currently it produces `dh_s2`): +1. **OLD path (delete):** `dh_s2_aux` was zero-filled by today's stop-grad; no further propagation. +2. **NEW path:** `dh_s2_aux` → `aux_trunk_backward` → `dW1/db1/dW2/db2/dW3/db3` for aux trunk Adam update. NO propagation back to encoder (stop-grad). + +Find aux backward call site in trainer (likely `gpu_dqn_trainer.rs`, search for `aux_next_bar_backward`). Insert `aux_trunk_backward` launch immediately after. + +Also: the aux head's backward kernels (`aux_next_bar_backward`, `aux_regime_backward`) currently zero-fill `dh_s2_out`. With this commit, they should write the actual `dh_s2_aux` gradient (no zero-fill) — since the consumer is now aux_trunk_backward, not Q's trunk. **Revert the zero-fills landed today** (commits 872bd7392 and 411a30473) — they're no longer needed because aux's gradient now flows into aux_trunk, not h_s2. + +- [ ] **Step 7: Add Adam update launch for aux trunk params** + +Find existing Adam launch pattern (e.g., `apply_adam_update` for GRN params). Add aux trunk param-group with ISV-driven β1/β2/ε/LR (slots from C.1). + +- [ ] **Step 8: Verify** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -15` + +Expected: clean build. Build failure here means the contract migration is non-atomic — fix before proceeding. + +- [ ] **Step 9: Run all oracle tests** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -20` + +Expected: all pre-existing tests still pass + 3 new aux_trunk tests pass. + +- [ ] **Step 10: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/aux_heads_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_heads.rs +git commit -m "feat(sp14-c): atomic aux trunk wire-up — fwd+bwd+aux head input switch + +Aux trunk inserted between encoder and aux head. Aux head input migrates +h_s2 → h_s2_aux atomically. Aux backward chain: aux CE → dh_s2_aux → +aux_trunk_backward → param Adam updates. Encoder gradient untouched +(stop-grad at trunk boundary). Today's stop-grad zero-fills in +aux_next_bar_backward + aux_regime_backward reverted — gradient now has +a valid consumer (aux_trunk_backward) instead of being structurally +discarded. Coupling A (dir_concat_qaux) unchanged — still feeds aux's +softmax-diff to Q's x_concat as forward feature." +``` + +--- + +## Phase C.6: h_s2_aux_rms_ema Producer Kernel + +### Task C.6.1: Adaptive scale producer for aux trunk output magnitude + +**Files:** +- Create: `crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu` +- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add producer launch op) +- Modify: `crates/ml/src/cuda_pipeline/mod.rs` +- Modify: `crates/ml/build.rs` +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (launch site) + +**Goal:** mirror `h_s2_rms_ema` (existing slot 96 ISV producer for GRN trunk's output magnitude) but for aux trunk's output `h_s2_aux`. Used downstream for adaptive bound computation. + +**Reference:** find existing `h_s2_rms_ema` kernel via `grep -rln "H_S2_RMS_EMA_INDEX\|h_s2_rms_ema" crates/ml/src/cuda_pipeline/`. Copy structure verbatim, retarget input buffer. + +- [ ] **Step 1: Write kernel** (mirrors h_s2_rms_ema): + - Input: `h_s2_aux [B, AUX_HIDDEN_DIM]` + - Output: `ISV[H_S2_AUX_RMS_EMA_INDEX]` (single float) + - Math: per-batch RMS → reduce-mean across batch → blend with EMA via Pearl-A bootstrap (sentinel-0 → batch_mean replace; subsequent: Welford EMA blend) + +- [ ] **Step 2: Wrapper, cubin registration, mod decl** — standard. + +- [ ] **Step 3: Wire into collector** AFTER `aux_trunk_forward`, BEFORE consumers (Adam launches that read this slot for adaptive bounds). + +- [ ] **Step 4: Unit test** + +```rust +#[test] +#[ignore] +fn h_s2_aux_rms_ema_pearl_a_bootstrap() { + // 1. Set ISV[H_S2_AUX_RMS_EMA_INDEX] = 0.0 (sentinel) + // 2. Launch with known h_s2_aux (controlled RMS) + // 3. Assert ISV slot == known RMS (replace, not blend) + // 4. Launch again with different RMS + // 5. Assert ISV slot is Welford-EMA blend of the two +} +``` + +- [ ] **Step 5: Verify + commit** + +```bash +git add crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/build.rs crates/ml/src/cuda_pipeline/mod.rs crates/ml/tests/aux_trunk_oracle_tests.rs +git commit -m "feat(sp14-c): h_s2_aux_rms_ema producer — adaptive scale for aux trunk + +Mirrors GRN trunk's h_s2_rms_ema (slot 96) for aux trunk output. +Pearl-A first-observation bootstrap, Welford EMA blend thereafter. +Wired into collector chain after aux_trunk_forward." +``` + +--- + +## Phase C.7: ~~Atomic Deletion of α Machinery~~ — MERGED INTO C.1 + +α machinery deletion folded into Phase C.1 (no DEPRECATED stage; atomic with new ISV slot allocation per `feedback_no_legacy_aliases`). Phase C.7 task #380 should be marked completed once C.1 lands. + +--- + +## Phase C.8: ISV-Driven Aux Trunk Hyperparameters + +### Task C.8.1: Wire Adam β1/β2/ε/LR/grad-clip from ISV slots + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (Adam launch site for aux trunk) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (if Adam launches live there) + +- [ ] **Step 1: Find existing per-group Adam pattern** + +SP5 Pearl 4 wired per-group Adam β1/β2/ε from ISV. Reference: search for `adaptive_adam_beta1` or `ISV_BETA1_BASE` in production launches. Find the existing pattern. + +- [ ] **Step 2: Wire aux trunk param-group** + +Aux trunk's Adam launch reads: +- `β1` from `ISV[AUX_TRUNK_BETA1_INDEX]` +- `β2` from `ISV[AUX_TRUNK_BETA2_INDEX]` +- `ε` from `ISV[AUX_TRUNK_EPS_INDEX]` +- `lr` from `ISV[AUX_TRUNK_LR_INDEX]` +- `grad_clip_norm` from `ISV[AUX_TRUNK_GRAD_CLIP_INDEX]` applied before Adam step + +Each is read from device-mapped ISV (no DtoH copy). + +- [ ] **Step 3: Cold-start defaults** + +Sentinel values from C.1 reset registry serve as bootstrapping defaults. Verify Pearl-A first-observation behavior: the producer that adapts these slots (if any — could be a controller TBD in a future phase) replaces sentinel on first valid observation. + +**Note:** this plan does NOT include adaptive controllers driving these slots — only the static sentinel→Adam path. Adaptive driving is a follow-up (mirrors how SP5 Pearl 4 was extended over multiple sub-tasks). + +- [ ] **Step 4: Verify** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --lib 2>&1 | tail -10` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +git commit -m "feat(sp14-c): aux trunk Adam reads β1/β2/ε/LR/grad-clip from ISV + +Per-group Adam pattern (mirrors SP5 Pearl 4). Sentinel defaults from +C.1 reset registry serve as bootstrap. Adaptive controllers driving +these slots deferred to follow-up phase." +``` + +--- + +## Phase C.9: Local Validation + Audit Doc + Memory Pearl + +### Task C.9.1: Synthetic-data smoke test (RTX 3050 Ti) + +**Files:** +- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (add smoke test) + +- [ ] **Step 1: Synthetic up-trend smoke** + +```rust +#[test] +#[ignore] +fn aux_trunk_learns_synthetic_uptrend() { + // 1. Generate synthetic batch: prices monotone increasing + // 2. Labels: all UP (label=1, K=2) + // 3. Run 100 forward+backward+Adam steps with deterministic seed + // 4. Assert: aux_loss < 0.1 by step 100 (ln(2)≈0.693 = random; <0.1 = clear signal of learning) + // 5. Assert: aux_dir_acc > 0.95 on the same synthetic data +} +``` + +- [ ] **Step 2: Run + verify** + +Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored aux_trunk_learns_synthetic_uptrend --nocapture 2>&1 | tail -20` + +Expected: aux_loss < 0.1 by step 100, dir_acc > 0.95. + +If the smoke test FAILS (aux can't learn synthetic up-trend), the architectural premise is wrong. Stop and investigate before L40S dispatch. + +### Task C.9.2: Audit doc close-out + memory pearl + +- [ ] **Step 1: Audit doc final entry** + +Append to `docs/dqn-wire-up-audit.md`: +- Title: "SP14 Layer C — separate aux trunk landed" +- Phase summary (C.1 through C.8 commit SHAs) +- Decision rationale: stop-grad at encoder boundary, ISV as control-plane-only, α machinery deleted +- Open follow-ups: ISV-driven adaptive controllers for aux LR/β/grad-clip (deferred to next phase) + +- [ ] **Step 2: Memory pearl candidate** + +Create `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_separate_aux_trunk_when_shared_starves.md`: + +``` +--- +name: separate aux trunk when shared trunk starves +description: When auxiliary head can't learn from a shared trunk despite gradient stop-grad, the cause is shared-trunk task incompatibility (Q-loss-shaped representation lacks aux-relevant features). Solution is separate aux trunk with stop-grad at encoder boundary, NOT more stop-grad / regularization tweaks +type: pearl +--- + +When an auxiliary head's loss flatlines at random baseline (ln(2) for K=2 CE) +even after stop-gradient prevents the aux loss from corrupting the shared trunk: +the bug is structural — Q-loss alone shapes the trunk to support value +estimation, which doesn't surface aux's task-relevant features. The aux head +reads the trunk forward but cannot reshape it backward (stop-grad). Result: +aux is starving for input. + +Fix is **separate aux trunk** parallel to Q's trunk, both reading shared +encoder output. Aux trunk receives gradient only from aux loss; encoder +remains Q-shaped via additional stop-grad at encoder boundary. + +Identification signal: aux_loss converges to and STAYS at the random-prediction +baseline (ln(2) ≈ 0.693 for binary CE; ln(K) generally) for several epochs +despite hyperparameter tuning. Pure stop-grad, regularization changes, or +loss weight adjustments will not fix this — they're treating the symptom of a +structural decoupling. + +Anti-pattern: invent ever-cleverer α coupling formulas (Schmitt gates, +variance-driven scaling) to bridge the gap. The gap is structural; bridging +it via gradient-flow tricks fails because the representation gap can't be +crossed by gradient scaling. + +Canonical: SP14 Layer C (2026-05-07), separate aux trunk for next-bar K=2 +prediction in DQN. Replaced EGF α machinery (Coupling B + Schmitt gate + +variance EMAs) with simple "aux gets its own representation pipeline." +Net delete ~300 LOC, net add ~750 LOC for trunk + tests. +``` + +Add index entry to `MEMORY.md` under "Architectural constraints": + +```markdown +- [pearl_separate_aux_trunk_when_shared_starves.md](pearl_separate_aux_trunk_when_shared_starves.md) — aux head flatlining at random baseline despite stop-grad means structural representation mismatch; fix is separate trunk, not gradient-flow cleverness; canonical SP14 Layer C (2026-05-07) +``` + +- [ ] **Step 3: Commit** + +```bash +git add docs/dqn-wire-up-audit.md +git commit -m "docs(sp14-c): close-out audit + memory pearl candidate + +Audit doc entry summarizing C.1-C.8 phases, decision rationale, and +deferred follow-ups (adaptive controllers for aux trunk hyperparams). +Memory pearl candidate: 'separate aux trunk when shared trunk starves' +— diagnostic for aux-loss-flatlines-at-ln(2) class of bugs." +``` + +(Memory file lives outside repo; commit only the audit doc.) + +--- + +## Phase C.10: L40S 30-Epoch Validation + +### Task C.10.1: Push + dispatch + monitor + +**Acceptance criteria:** +- aux `next_bar_mse` (CE) drops below **0.65** by epoch 5 (vs. ln(2)=0.693 random baseline) +- aux `next_bar_mse` continues dropping (not flat) through epoch 10 +- val_Sharpe trajectory does NOT collapse to <30 by epoch 15 (current bqtnd/thtxx baseline shows Sharpe→40 by ep10 with Hold-collapse) +- val_dir_dist `hold` stays below 50% through epoch 15 +- No NaN/Inf in any metric + +- [ ] **Step 1: Push** + +```bash +git push origin sp15-phase1-honest-numbers +``` + +- [ ] **Step 2: Verify HEAD on remote** + +```bash +git ls-remote origin sp15-phase1-honest-numbers +``` + +Expected: SHA matches local HEAD. + +- [ ] **Step 3: Dispatch** + +```bash +./scripts/argo-train.sh dqn --branch sp15-phase1-honest-numbers --gpu-pool ci-training-l40s --epochs 30 --baseline --tag sp14-c-aux-trunk +``` + +- [ ] **Step 4: Monitor with tightened filter** + +Arm Monitor on log stream filtering for: aux_loss trajectory, val_Sharpe, val_dir_dist hold%, train_active_frac, NaN/Inf, FATAL/panic, workflow Succeeded/Failed. + +Kill-fast triggers (per `feedback_kill_runs_on_anomaly_quickly`): +- aux next_bar_mse > 0.69 at epoch 5 (no learning — separate trunk hypothesis falsified, escalate) +- val_Sharpe collapse to <20 by epoch 10 (model worse than baseline) +- NaN/Inf anywhere +- CUDA_ERROR_ILLEGAL_ADDRESS or panic + +- [ ] **Step 5: On success — close out** + +If aux_loss < 0.65 by ep5 and trajectory is healthy through ep30: +- Mark memory pearl as canonical (move from "candidate" status) +- Promote SP14 Layer C task in MEMORY.md project state +- Begin scoping ISV-driven adaptive controllers for aux trunk hyperparams (deferred phase) + +- [ ] **Step 6: On failure — escalate** + +If aux_loss STILL flatlines at ln(2) with separate trunk: +- The deeper hypothesis (next-bar prediction is intrinsically too hard at HFT scale) is confirmed +- Options: (a) drop aux entirely, replace EGF intent with q_disagreement-only signal; (b) keep aux as feature-only via Coupling A but accept it's untrainable; (c) change aux's task (next-bar → multi-bar return, regime → trend phase) +- Brainstorm separately; not in this plan + +--- + +## Self-Review Checklist (run after writing complete) + +- [ ] **Spec coverage:** every locked decision has a phase that implements it +- [ ] **No placeholders:** zero `TODO`, `tbd`, `// fill in later`, `(implementation per existing pattern)` without a specific file:line reference +- [ ] **Type consistency:** `AUX_HIDDEN_DIM` referenced uniformly; `h_s2_aux` buffer name consistent across phases; `CudaFunction` pre-load discipline consistent +- [ ] **Atomic refactor rule:** Phase C.5 wire-up and Phase C.7 deletion are each ONE commit (not split) +- [ ] **Hard rules met:** no atomicAdd in any new kernel; ISV slot consumers never read pre-bootstrap sentinels without Pearl-A first-observation handling; no host branches in graph-captured paths +- [ ] **Test coverage:** forward oracle, backward gradient check, stop-grad invariant, Pearl-A bootstrap test, synthetic learning smoke +- [ ] **Validation gate:** acceptance criteria in C.10 are measurable and falsifiable + +--- + +## Execution Handoff + +**Recommended execution path:** subagent-driven-development. + +Reasoning: 10 phases, each phase is one commit (mostly mechanical implementation given the locked design), independent enough that reviewer subagents can verify per-phase. Subagent-driven beats inline because: +- Each phase requires reading several files to find existing patterns — subagents can do that exploration without polluting controller context +- Spec-compliance review catches scope creep (e.g., adding adaptive controllers in C.8 when they're explicitly deferred) +- Code-quality review catches missing CudaFunction pre-loads, atomicAdd violations, missing reset_registry entries + +**REQUIRED SUB-SKILL:** Use superpowers:subagent-driven-development to execute.