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=450].
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 target-variance EMA available, fallback per
pearl_wiener_optimal_adaptive_alpha)
- "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). The existing aux_next_bar_loss_reduce
in aux_heads_kernel.cu already supports the -1 mask convention via the
B_valid count — no new valid_mask parameter needed.
Step 5b finding: Case B — existing per-sample buffers
(hold_at_exit_per_sample, trade_profitable_per_sample) populated by
unified_env_step_core, but no aggregate ISV slot. Added new aggregator
slot AVG_WIN_HOLD_TIME_BARS_INDEX=451 + new producer kernel
avg_win_hold_time_update_kernel.cu (block-tree-reduce, no atomicAdd).
ISV_TOTAL_DIM bumped 450 → 452.
ATOMIC migration per feedback_no_partial_refactor: both label kernels
(aux_sign_label_kernel.cu trajectory + aux_sign_label_per_step_kernel.cu
per-rollout-step) migrated together to the new
(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)
signature. The lookahead host-passed scalar argument is removed; H is
read from ISV inside the kernel (broadcast value, single read per
thread, on-device clamp [1, 240]).
Producer chain (per-epoch boundary): new
GpuDqnTrainer::launch_aux_horizon_chain orchestrates
avg_win_hold_time_update → aux_horizon_update sequentially alongside
launch_kelly_cap_update at the existing epoch-boundary slot in
training_loop.rs.
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; the adaptive
producer drives H from real winning-trade observations.
Tests (8 oracle, 5 new + 3 preserved):
- aux_trunk_forward_matches_numpy_reference (C.3) ✓
- aux_trunk_backward_gradient_check (C.4) ✓
- aux_trunk_backward_does_not_write_dx (C.4) ✓
- aux_sign_label_h_bar_horizon (NEW) ✓
- aux_sign_label_lookahead_mask (NEW) ✓
- aux_horizon_pearl_a_bootstrap (NEW) ✓
- aux_horizon_converges_to_steady_target (NEW) ✓
- aux_horizon_holds_sentinel_with_no_winning_trades (NEW) ✓
8/8 pass on RTX 3050 Ti.
Phase C.4b of SP14 Layer C separate-aux-trunk refactor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
118
crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu
Normal file
118
crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu
Normal file
@@ -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));
|
||||
}
|
||||
@@ -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::<i32>(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"
|
||||
|
||||
@@ -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
|
||||
|
||||
137
crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu
Normal file
137
crates/ml/src/cuda_pipeline/avg_win_hold_time_update_kernel.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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<CudaStream>) -> Result<Self, MLError> {
|
||||
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<CudaStream>,
|
||||
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<CudaStream>) -> Result<Self, MLError> {
|
||||
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<CudaStream>,
|
||||
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::<f32>() 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(())
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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<f32>,
|
||||
@@ -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::<f32>(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,
|
||||
|
||||
@@ -4090,19 +4090,39 @@ impl GpuExperienceCollector {
|
||||
.alloc_zeros::<i32>(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 }
|
||||
|
||||
|
||||
@@ -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,
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream>) -> 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<CudaStream>,
|
||||
kernel: &CudaFunction,
|
||||
targets_buf: &MappedF32Buffer,
|
||||
bar_indices_buf: &MappedF32Buffer,
|
||||
isv_buf: &MappedF32Buffer,
|
||||
total: usize,
|
||||
total_bars: i32,
|
||||
) -> Vec<i32> {
|
||||
let labels_dev = stream
|
||||
.alloc_zeros::<i32>(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<i32> = (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<f32> = (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<f32> = (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]<p[t]), got {lbl}");
|
||||
} else {
|
||||
assert_eq!(lbl, -1, "decreasing series: bar {i} expected mask -1, got {lbl}");
|
||||
}
|
||||
}
|
||||
eprintln!("Case 2 (decreasing): valid bars={lookahead_cutoff}, all 0 — OK");
|
||||
|
||||
// Case 3: sine wave — analytic comparison
|
||||
let prices_sin: Vec<f32> =
|
||||
(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<i32> = (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<f32> = (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<CudaStream>,
|
||||
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)");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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.
|
||||
|
||||
1243
docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md
Normal file
1243
docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md
Normal file
File diff suppressed because it is too large
Load Diff
Reference in New Issue
Block a user