From 0ca45ef61d67fe2ed3ebbdb9d7bd105b9f63f0fe Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 22:24:16 +0200 Subject: [PATCH] docs(sp13): v2 spec + plan after critical review MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Spec v2 supersedes v1 with five major fixes from the critical review: - Drop alpha-vs-benchmark reward (mathematically tautological — Long trades had alpha = -costs always against always-long benchmark) - Split Phase 0 into 0a (Hold-only, clean test of user hypothesis) + 0b (aux amplification probe, only if 0a partial), avoiding the v1 confounded experiment - Bound direction-skill bonus relative to |alpha| (cap_ratio × |alpha|) to prevent reward gaming on small-alpha correct-direction losers; spec adds 8-quadrant worked-example matrix verifying the no-negative-EV invariant - Replace 5-epoch absolute-threshold gate with 10-epoch trajectory criterion to avoid false-negatives from aux head underconvergence - Aux_w controller adds dual-EMA stagnation detector (decays toward base when no improvement) — prevents permanent destabilization of Q-head in data-limited case Plan v2 mirrors spec changes: - Phase 0a (Hold elimination + dir_acc instrumentation, ~300 LOC, 1 atomic commit) - Phase 0b (aux_w controller replacement, conditional, ~80 LOC) - Layer B (aux head regression -> binary classification, ~120 LOC) - Layer C (skill bonus + luck discount with calibrated bounds, ~150 LOC) - Layer D (30-epoch validation + 3 new pearls) ISV slot allocation [372..380): drops slot 371 (BENCHMARK_PNL_CUMULATIVE), adds 374 (AUX_DIR_ACC_LONG_EMA — stagnation), 379 (SKILL_BONUS_CAP_RATIO). Plan integrates Explore-agent touch-list (50-70 sites) with corrected enum ordering: Short=0, Long=1, Flat=2 (preserves codebase Short-first convention, avoids 30+ stale-comment churn). Adds direction-bias signal handling at experience_kernels.cu:5159 (Hold's 0.5 softening gate vanishes -> [1, 1, 0]). Three new pearls planned for Layer D close-out: - pearl_redefine_success_for_predictive_skill - pearl_skill_bonus_must_be_alpha_bounded (calibration lesson) - pearl_reward_quadrant_audit_required (meta-pearl) Co-Authored-By: Claude Opus 4.7 (1M context) --- ...3-redefine-success-for-predictive-skill.md | 1228 +++++++++++++++++ ...3-redefine-success-for-predictive-skill.md | 386 ++++++ 2 files changed, 1614 insertions(+) create mode 100644 docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md create mode 100644 docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md diff --git a/docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md b/docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md new file mode 100644 index 000000000..ab2fd5833 --- /dev/null +++ b/docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md @@ -0,0 +1,1228 @@ +# SP13 Redefine Success for Predictive Skill — Implementation Plan (v2) + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Test the user's hypothesis that the persistent ~46% win rate is caused by Hold action hiding directional signal (not by absence of signal in data) by removing Hold cleanly in Phase 0a. If confirmed, build out direction-skill bonus + correct-call asymmetry (calibrated to prevent reward gaming) and aux head binary classification. v1 design dropped alpha-vs-benchmark reward (mathematically broken — see spec v1→v2 changes). + +**Architecture:** 8 new ISV slots `[372..380)`, 1 new GPU producer kernel (`aux_dir_acc_reduce_kernel.cu`), corrected aux_w controller with stagnation detector (Phase 0b only), full action-space refactor (4 → 3, eliminating Hold), aux head refactor (regression → binary classification), reward composition extension (skill bonus + luck discount on top of SP12). Phases: P0a (Hold-only, decisive gate, 1 commit), P0b (aux amplification, only if P0a partial/red, 1 commit), B (aux classification, 1 commit), C (reward composition, 1 atomic commit), D (30-epoch validation + close-out). + +**Tech Stack:** Rust 1.85+, CUDA 12.4, cudarc, mapped-pinned memory, block tree-reduce (no atomicAdd), GPU-only Pearls A+D via `apply_pearls_ad_kernel`, L40S smoke validation via Argo Workflows, RTX 3050 Ti for local oracle tests. + +**Spec source:** `docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md` (v2) + +**Branch:** `sp11-reward-as-controlled-subsystem` (greenfield refactor — replay buffer compatibility intentionally broken in P0a) + +**Operating principles** (from MEMORY.md): +- `feedback_no_cpu_compute_strict` — every producer GPU; Pearls A+D via `apply_pearls_ad_kernel` +- `feedback_no_atomicadd` — block tree-reduce only +- `feedback_no_partial_refactor` — P0a / P0b / B / C each single atomic commit +- `feedback_isv_for_adaptive_bounds` — all 8 knobs ISV-driven +- `feedback_no_htod_htoh_only_mapped_pinned` — mapped-pinned only +- `feedback_no_cpu_test_fallbacks` — GPU oracle testing only +- `feedback_no_stubs` / `feedback_no_todo_fixme` / `feedback_no_feature_flags` +- End-to-end training: NO pre-training of aux head +- `pearl_audit_unboundedness_for_implicit_asymmetry` — SP12 asymmetric cap stays +- `pearl_event_driven_reward_density_alignment` — reward fires only at trade events +- **NEW reward-quadrant audit rule** — every reward change must come with a 4-quadrant + worked-example table (correct/wrong × winner/loser); enforced in Layer C tests + +**Hypothesis under test** (from `project_hold_hides_directional_signal_hypothesis.md`): +- P0a green (10-epoch smoke shows dir_acc and WR climbing) → CONFIRMED, proceed +- P0a partial → P0b adds aux amplification probe +- Both red → DISCONFIRMED, halt SP13, plan SP14 data-augmentation work separately + +--- + +## File structure + +### New files (Phase 0a) + +| File | Phase | Responsibility | +|---|---|---| +| `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` | P0a | 8 SP13 slot index constants `[372..380)` + defaults | +| `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` | P0a | Single-block tree-reduce: 4 shared arrays (correct, pos_pred, pos_label, valid) → `(dir_acc, pos_pred_frac, pos_label_frac)` | +| `crates/ml/tests/sp13_phase0_oracle_tests.rs` | P0a | 6+ GPU oracle tests for dir_acc kernel + 4 action-selection tests for 3-direction argmax | + +### Modified files (Phase 0a) + +| File | Change | +|---|---| +| `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs` | Bump `SP5_SLOT_END = 380`, `ISV_TOTAL_DIM = 380`, layout fingerprint seed | +| `crates/ml/src/cuda_pipeline/state_layout.cuh` | `NUM_DIRECTIONS: 4 → 3`; `DIRECTION_SHORT=0, DIRECTION_LONG=1, DIRECTION_FLAT=2` (DIRECTION_HOLD deleted); `ISV_TOTAL_DIM` mirror; SP13 slot defines | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | 2 new entries (slots 373, 374) per-fold sentinel 0.5; line :974 `TRAIL_DIST_PER_DIR_BASE=270..274` comment update | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Cubin include + launcher for `aux_dir_acc_reduce`; `aux_dir_acc_buf: MappedF32Buffer` (3 floats); static-init slot 372 = 0.55 | +| `crates/ml/build.rs` | Cubin registration; line :490 stale `Short=0, Hold=1, Long=2, Flat=3` doc-comment update | +| **`crates/ml-core/src/common/action.rs:124-132`** | Canonical `DirectionAction` Rust enum — eliminate `Hold` variant; `Short=0, Long=1, Flat=2` | +| **`crates/ml-core/src/state_layout.rs:48-49`** | `pub const NUM_DIRECTIONS: usize = 3` (was 4); doc-comment update | +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Multi-site: line :1013 `experience_action_select` direction Thompson 4 → 3; line :4059 direction accumulation loop; line :5159 **direction-bias signal `[1.0, 0.5, 1.0, 0.0]` → `[1.0, 1.0, 0.0]`** (Hold's 0.5 softening gate removed; see P0a.T3 step 4 below); line :5306 conviction EMA loop; lines :6371-6397 per-atom direction loops (4 occurrences) | +| `crates/ml/src/cuda_pipeline/c51_loss_kernel.cu` | Lines :394 (doc), :460 (`for k<4` argmax), :1088 (argmax across direction bins) | +| `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` | Implementer AUDIT for `for.*<4` direction loops (Explore agent flagged for verification) | +| `crates/ml/src/cuda_pipeline/iqn_loss_kernel.cu` / `iqn_dual_head_kernel.cu` | Lines :245, :1199, :200 — `actions[sample * 4 + 0]` direction index UNCHANGED dimensionally (action layout still [direction, magnitude, order, urgency]); only direction VALUE range shrinks to 0..2 | +| `crates/ml/src/cuda_pipeline/dqn_inference.cu` | Direction Q-head output dim 4 → 3 | +| `crates/ml/src/cuda_pipeline/dqn_param_layout.rs` | Direction Q-head weight/bias shapes; layout fingerprint seed bump | +| `crates/ml/src/trainers/dqn/financials.rs:27,197` | `exp_idx = dir*3 + mag` packing — **range shrinks 12 → 9** (was 4 dirs × 3 mags = 12; now 3 dirs × 3 mags = 9); doc-comment :27 update | +| `crates/ml/src/trainers/dqn/monitoring.rs:13,341` | `action_counts[9..12]` Flat-mag indexing recomputes; doc-comment :13 stale `S=Short(0), H=Hold(1), L=Long(2), F=Flat(3)` update | +| `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs:272` | `TRAIL_DIST_PER_DIR_BASE` doc-comment `[4 dirs: Short, Hold, Long, Flat]` update — slot range KEPT at 270..274 (4 entries) per spec; slot 273 becomes dead-space (consumers read by direction index 0..2; never touched) | +| `crates/ml/src/cuda_pipeline/pearl_8_trail_kernel.cu` | Producer kernel writes per-direction trail dist at NEW indices `{Short→0, Long→1, Flat→2, dead→3}`; consumer reads unchanged | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs:4046-4048` | Pearl 8 `apply_pearls_ad_kernel` calls — 4 calls become 3 (dead slot 273 not smoothed); also dir_acc launch + ISV write + HEALTH_DIAG `aux_dir_acc` + val_dir_dist 3-bucket emit | +| `crates/ml/src/trainers/dqn/trainer/metrics.rs:949` | `val_dir_dist [short=… hold=… long=… flat=…]` → `[short=… long=… flat=…]` | +| `crates/ml/src/trainers/dqn/smoke_tests/magnitude_distribution.rs:88,217` | `EVAL_DIR_DIST` 4-tuple → 3-tuple emit | +| `crates/ml/src/trainers/dqn/distributional_q_tests.rs:790,888` | `σ_C51 per direction` test output 4 → 3; `PROD_B0=4` test constant → 3 | +| `crates/ml/src/cuda_pipeline/kl_divergence_kernel.cu` (if exists) | 3-bucket update — Explore agent did NOT find this file; verify nonexistence or update if found | +| `crates/ml/src/trainers/dqn/trainer/tests.rs:542` | `assert_eq!(config.num_actions, 4, …)` → `3` | +| `crates/ml/src/trainers/dqn/trainer/mod.rs:1966` | Doc-comment `Layout: [Short, Hold, Long, Flat]` → `[Short, Long, Flat]` | +| `crates/ml/src/cuda_pipeline/thompson_test_kernel.cu` | Test kernel direction count 4 → 3 | +| 20+ stale doc lines across kernel headers (`.cu` files) | Implementer audit: `grep -rn "Short=0\|Hold=1\|Long=2\|Flat=3" crates/ml/src/` and update narrative comments to `Short=0, Long=1, Flat=2` | + +**Note on slot 375 (`AUX_DIR_PREDICTION_INDEX`) in P0a:** the existing aux head still emits a regression scalar. Phase 0a writes `tanh(scalar_pred)` into slot 375 to fit the [-1, +1] interface so Layer B can refactor without changing wire format. This is NOT a stub — it produces a real prediction signal, just with the legacy regression head. Layer B replaces with softmax-derived logit_diff. + +### Modified files (Phase 0b — only if P0a partial/red) + +| File | Change | +|---|---| +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs:4103-4104` | Replace inverted SP11 formula with deficit + stagnation controller | +| `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` | Add `AUX_W_BASE_PHASE0B = 0.5`, `AUX_W_STAGNATION_DECAY_K = 0.7` constants | + +### Modified files (Layer B) + +| File | Change | +|---|---| +| `crates/ml/src/cuda_pipeline/aux_head_kernel.cu` | Output dim 1 → 2; activation linear → softmax; loss MSE → CrossEntropy | +| `crates/ml/src/cuda_pipeline/aux_label_gen_kernel.cu` | Label = `sign(price[t+30] - price[t])` as 0/1 class index | +| `crates/ml/src/cuda_pipeline/aux_dir_acc_reduce_kernel.cu` | Read 2-class softmax instead of regression scalar; pred_pos = (softmax[1] > softmax[0]) | +| `crates/ml/src/cuda_pipeline/dqn_inference.cu` | Direction Q-head input concat: `[trunk_features, ISV[AUX_DIR_PREDICTION_INDEX]]`; in_dim grows by 1 | +| `crates/ml/src/cuda_pipeline/dqn_param_layout.rs` | Direction Q-head input dim +1 | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Replace `tanh(scalar)` ISV[375] write with `softmax[1] - softmax[0]` | +| `crates/ml/tests/sp13_phase0_oracle_tests.rs` | +3 cross-entropy classification tests | + +### New files (Layer C) + +| File | Phase | Responsibility | +|---|---|---| +| `crates/ml/src/cuda_pipeline/sp13_reward_math_test_kernel.cu` | C | GPU oracle wrapper for 2 device-inline reward funcs (skill bonus, luck discount); mirrors SP12 sp12_reward_math_test_kernel.cu | +| `crates/ml/tests/sp13_reward_math_tests.rs` | C | 16 GPU oracle tests covering full 8-quadrant matrix from spec | + +### Modified files (Layer C) + +| File | Change | +|---|---| +| `crates/ml/src/cuda_pipeline/trade_physics.cuh` | Add 2 device-inline funcs: `compute_direction_skill_bonus_bounded`, `compute_luck_win_discount` | +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` (~line 2788, exit-event reward) | After SP12 asymmetric cap, ADD skill bonus then ADD luck discount; min_hold_factor unchanged | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | 4 new entries (slots 376, 377, 378, 379) per-fold | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | At fold-end HEALTH_DIAG: emit `per_fold_alpha_vs_buy_and_hold = total_pnl - (price[fold_end] - price[fold_start])` as METRIC ONLY | + +### Documentation files (Layer D) + +| File | Phase | Content | +|---|---|---| +| `docs/dqn-wire-up-audit.md` | D | SP13 entry: hypothesis, P0a/P0b results, Layer B/C deltas, 30-epoch metrics | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_redefine_success_for_predictive_skill.md` | D | NEW pearl | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_skill_bonus_must_be_alpha_bounded.md` | D | NEW pearl — calibration lesson from v1 review | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_reward_quadrant_audit_required.md` | D | NEW meta-pearl from v1 review | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` | D | Index entries | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_hold_hides_directional_signal_hypothesis.md` | D | Status → CONFIRMED / DISCONFIRMED with smoke results | + +--- + +## Phase 0a — Hold elimination + dir_acc instrumentation (1 atomic commit) + +**Per `feedback_no_partial_refactor`**: every consumer of the 4-way direction action space migrates to 3-way in a single commit. Replay buffer compatibility intentionally broken (greenfield). + +### Task P0a.T1: ISV slot constants + state-reset registry + +**Files:** Create `sp13_isv_slots.rs`; modify `sp5_isv_slots.rs`, `state_layout.cuh`, `state_reset_registry.rs`, `gpu_dqn_trainer.rs` + +- [ ] **Step 1:** Create `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs`: + +```rust +//! SP13 ISV slot constants — see spec docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md + +pub const TARGET_DIR_ACC_INDEX: usize = 372; +pub const AUX_DIR_ACC_SHORT_EMA_INDEX: usize = 373; +pub const AUX_DIR_ACC_LONG_EMA_INDEX: usize = 374; +pub const AUX_DIR_PREDICTION_INDEX: usize = 375; +pub const DIR_SKILL_BONUS_ALPHA_INDEX: usize = 376; +pub const DIR_SKILL_BONUS_BETA_INDEX: usize = 377; +pub const LUCK_WIN_DISCOUNT_INDEX: usize = 378; +pub const SKILL_BONUS_CAP_RATIO_INDEX: usize = 379; + +pub const TARGET_DIR_ACC_DEFAULT: f32 = 0.55; +pub const DIR_SKILL_BONUS_ALPHA_DEFAULT: f32 = 1.0; +pub const DIR_SKILL_BONUS_BETA_DEFAULT: f32 = 1.0; +pub const LUCK_WIN_DISCOUNT_DEFAULT: f32 = 0.3; +pub const SKILL_BONUS_CAP_RATIO_DEFAULT: f32 = 0.3; + +/// Sentinel for both dir_acc EMAs at fold-start: 0.5 (random baseline) +/// — first observation replaces directly per pearl_first_observation_bootstrap +pub const DIR_ACC_EMA_SENTINEL: f32 = 0.5; +``` + +- [ ] **Step 2:** Bump SP5_SLOT_END in `sp5_isv_slots.rs`: + +```rust +pub const SP5_SLOT_END: usize = 380; // was 371 post-SP12; +8 for SP13 +pub const ISV_TOTAL_DIM: usize = SP5_SLOT_END; +``` + +Add `pub use crate::cuda_pipeline::sp13_isv_slots::*;` re-export. Bump layout fingerprint seed. + +- [ ] **Step 3:** Mirror in `state_layout.cuh` — bump `ISV_TOTAL_DIM` to 380; add `// === SP13 SLOT INDICES ===` block with all 8 defines. + +- [ ] **Step 4:** Add reset registry entries: + +```rust +FoldResetEntry { name: "aux_dir_acc_short_ema", slot_index: AUX_DIR_ACC_SHORT_EMA_INDEX, reset_value: DIR_ACC_EMA_SENTINEL, reset_kind: ResetKind::PerFold }, +FoldResetEntry { name: "aux_dir_acc_long_ema", slot_index: AUX_DIR_ACC_LONG_EMA_INDEX, reset_value: DIR_ACC_EMA_SENTINEL, reset_kind: ResetKind::PerFold }, +``` + +(Slots 372, 376-379 are static-initialised in trainer constructor — no reset entry needed; slot 375 is per-bar overwritten — no reset needed.) + +- [ ] **Step 5:** In `GpuDqnTrainer::new` (or equivalent constructor), add static-init for the 6 static slots: + +```rust +self.write_isv_slot(TARGET_DIR_ACC_INDEX, TARGET_DIR_ACC_DEFAULT)?; +self.write_isv_slot(DIR_SKILL_BONUS_ALPHA_INDEX, DIR_SKILL_BONUS_ALPHA_DEFAULT)?; +self.write_isv_slot(DIR_SKILL_BONUS_BETA_INDEX, DIR_SKILL_BONUS_BETA_DEFAULT)?; +self.write_isv_slot(LUCK_WIN_DISCOUNT_INDEX, LUCK_WIN_DISCOUNT_DEFAULT)?; +self.write_isv_slot(SKILL_BONUS_CAP_RATIO_INDEX, SKILL_BONUS_CAP_RATIO_DEFAULT)?; +``` + +(The α/β/luck/cap-ratio slots aren't *consumed* until Layer C, but writing them in P0a is harmless and avoids a follow-up registry update in C.) + +- [ ] **Step 6:** Verify compile: `SQLX_OFFLINE=true cargo check -p ml --lib` + +- [ ] **Step 7:** Stage (no commit yet — atomic at end of P0a). + +--- + +### Task P0a.T2: aux_dir_acc_reduce kernel + GPU oracle tests + +**Goal:** Single-block tree-reduce kernel with 4 shared-mem arrays (correct, pos_pred, pos_label, valid). Reads regression-mode aux scalar in P0a (Layer B refactors to softmax). 6+ GPU oracle tests on RTX 3050 Ti. + +- [ ] **Step 1:** Write the kernel (`aux_dir_acc_reduce_kernel.cu`): + +```cuda +#include "state_layout.cuh" + +extern "C" __global__ void aux_dir_acc_reduce_kernel( + const float* __restrict__ aux_pred, // [B] regression scalar (Phase 0a); Layer B will swap to softmax read + const float* __restrict__ next_bar_label, // [B] sign-encoded: +1 / 0 / -1 + int batch_size, + float* __restrict__ out_3 // [3] = {dir_acc, pos_pred_frac, pos_label_frac} +) { + const int tid = threadIdx.x; + const int bdim = blockDim.x; + + extern __shared__ int shared[]; + int* sh_correct = &shared[0 * bdim]; + int* sh_pos_pred = &shared[1 * bdim]; + int* sh_pos_label = &shared[2 * bdim]; + int* sh_valid = &shared[3 * bdim]; + + int local_correct = 0, local_pos_pred = 0, local_pos_label = 0, local_valid = 0; + for (int i = tid; i < batch_size; i += bdim) { + const float p = aux_pred[i]; + const float l = next_bar_label[i]; + if (l == 0.0f) continue; + local_valid += 1; + const int pred_pos = (p > 0.0f) ? 1 : 0; + const int label_pos = (l > 0.0f) ? 1 : 0; + if (pred_pos == label_pos) local_correct += 1; + local_pos_pred += pred_pos; + local_pos_label += label_pos; + } + sh_correct[tid] = local_correct; + sh_pos_pred[tid] = local_pos_pred; + sh_pos_label[tid] = local_pos_label; + sh_valid[tid] = local_valid; + __syncthreads(); + + for (int s = bdim / 2; s > 0; s >>= 1) { + if (tid < s) { + sh_correct[tid] += sh_correct[tid + s]; + sh_pos_pred[tid] += sh_pos_pred[tid + s]; + sh_pos_label[tid] += sh_pos_label[tid + s]; + sh_valid[tid] += sh_valid[tid + s]; + } + __syncthreads(); + } + + if (tid == 0) { + const float denom = (float)sh_valid[0]; + const float sentinel = 0.5f; + out_3[0] = (denom > 0.0f) ? ((float)sh_correct[0] / denom) : sentinel; + out_3[1] = (denom > 0.0f) ? ((float)sh_pos_pred[0] / denom) : sentinel; + out_3[2] = (denom > 0.0f) ? ((float)sh_pos_label[0] / denom) : sentinel; + } +} +``` + +Block dim: 256, shared mem: `4 × 256 × sizeof(int) = 4 KB`. + +- [ ] **Step 2:** Register cubin in `build.rs`. + +- [ ] **Step 3:** Add launcher in `gpu_dqn_trainer.rs` (mirrors existing reduction-kernel launcher patterns, e.g. `launch_q_stats_reduce`): + +```rust +pub fn launch_aux_dir_acc_reduce( + &self, + aux_pred: &CudaSlice, + next_bar_label: &CudaSlice, + batch_size: i32, + out_3: &mut CudaSlice, +) -> Result<()> { + let bdim: u32 = 256; + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (bdim, 1, 1), + shared_mem_bytes: 4 * bdim * std::mem::size_of::() as u32, + }; + unsafe { + self.aux_dir_acc_reduce + .clone() + .launch(cfg, (aux_pred, next_bar_label, batch_size, out_3))?; + } + Ok(()) +} +``` + +Add `aux_dir_acc_reduce: CudaFunction` field + `aux_dir_acc_buf: MappedF32Buffer` (3 floats). + +- [ ] **Step 4:** Write `crates/ml/tests/sp13_phase0_oracle_tests.rs`: + +```rust +#![cfg(feature = "gpu-tests")] + +use ml::cuda_pipeline::gpu_dqn_trainer::GpuDqnTrainer; + +fn test_trainer() -> GpuDqnTrainer { /* mirror sp12_reward_math_tests scaffold */ } +fn run_dir_acc(trainer: &GpuDqnTrainer, pred: &[f32], label: &[f32]) -> [f32; 3] { /* ... */ } + +#[test] +fn dir_acc_all_correct_returns_one() { + let pred = vec![1.0_f32, 1.0, -1.0, -1.0]; + let label = vec![1.0_f32, 1.0, -1.0, -1.0]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!((out[0] - 1.0).abs() < 1e-6); +} + +#[test] +fn dir_acc_all_wrong_returns_zero() { + let pred = vec![1.0_f32, 1.0, -1.0, -1.0]; + let label = vec![-1.0_f32, -1.0, 1.0, 1.0]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!(out[0].abs() < 1e-6); +} + +#[test] +fn dir_acc_half_correct() { + let pred = vec![1.0_f32, 1.0, 1.0, 1.0]; + let label = vec![1.0_f32, 1.0, -1.0, -1.0]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!((out[0] - 0.5).abs() < 1e-6); + assert!((out[1] - 1.0).abs() < 1e-6); + assert!((out[2] - 0.5).abs() < 1e-6); +} + +#[test] +fn dir_acc_skips_zero_labels() { + let pred = vec![1.0_f32, 1.0, -1.0]; + let label = vec![1.0_f32, 0.0, -1.0]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!((out[0] - 1.0).abs() < 1e-6); +} + +#[test] +fn dir_acc_zero_pred_is_negative() { + // (p > 0.0f) ? 1 : 0 — zero predicts class 0 (negative) + let pred = vec![0.0_f32, -0.0]; + let label = vec![-1.0_f32, -1.0]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!((out[0] - 1.0).abs() < 1e-6); +} + +#[test] +fn dir_acc_empty_batch_returns_sentinel() { + let pred: Vec = vec![]; + let label: Vec = vec![]; + let out = run_dir_acc(&test_trainer(), &pred, &label); + assert!((out[0] - 0.5).abs() < 1e-6); +} +``` + +- [ ] **Step 5:** Run on RTX 3050 Ti: `SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp13_phase0_oracle_tests`. Expected: 6/6 passing. + +- [ ] **Step 6:** Stage. + +--- + +### Task P0a.T3: Direction action space refactor — atomic Hold removal + +**Goal:** `NUM_DIRECTIONS: 4 → 3`. Eliminate Hold from `DirectionAction` enum. Update all 7+ consumers in lockstep. Per `feedback_no_partial_refactor` — single commit, no partial migration. + +**Action semantics in 3-way world** (replaces `project_hold_action_design`): + +| Prior position | Action picked | Result | +|---|---|---| +| Flat | Long | Open Long | +| Long | Long | Keep Long (was old Hold-Long) | +| Short | Long | Flip to Long | +| Long | Short | Flip to Short | +| Long | Flat | Close position | +| Flat | Flat | No-op | + +`Long` is "have Long position", not "open Long". Same for Short. Flat = "no position". + +**Encoding (CRITICAL — preserves codebase Short-first canonical anchor):** +``` +OLD: Short=0, Hold=1, Long=2, Flat=3 (4-way) +NEW: Short=0, Long=1, Flat=2 (3-way) — Hold removed; Long & Flat shift down +``` + +This is NOT `Long=0, Short=1, Flat=2` — keeping Short=0 avoids 30+ stale-comment churn across kernel headers and tests. + +- [ ] **Step 1:** Update `state_layout.cuh`: + +```c +// SP13: direction action space reduced 4 → 3. Hold removed; Long shifts 2→1, Flat 3→2. +// Short-first ordering preserved per codebase canonical anchor. +#define NUM_DIRECTIONS 3 +#define DIRECTION_SHORT 0 +#define DIRECTION_LONG 1 +#define DIRECTION_FLAT 2 +// (DIRECTION_HOLD definition deleted — was =1 in 4-way scheme) +``` + +- [ ] **Step 2:** Update canonical Rust enum at **`crates/ml-core/src/common/action.rs:124-132`** (NOT `crates/ml/src/agents/policy.rs` — that path was speculative; Explore agent confirmed the canonical site): + +```rust +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum DirectionAction { + Short = 0, + Long = 1, + Flat = 2, +} +``` + +And `crates/ml-core/src/state_layout.rs:48-49`: + +```rust +// (3 actions: Short / Long / Flat; Hold removed in SP13) +pub const NUM_DIRECTIONS: usize = 3; // was 4 +``` + +Compiler will surface every `Hold` match arm. Fix each — do NOT add `_ => unreachable!()` cover-ups. + +- [ ] **Step 3:** Update `experience_action_select` direction branch at `experience_kernels.cu:1013`: argmax over 3 instead of 4. Thompson sampling on 3 atoms (`pearl_thompson_for_distributional_action_selection`). + +- [ ] **Step 4:** **NEW — Direction-bias signal at `experience_kernels.cu:5159`:** + +The 4-tuple per-direction bias weights `[Short=1.0, Hold=0.5, Long=1.0, Flat=0.0]` encoded a softening gate where Hold attenuated the Q-signal. With Hold removed: + +```cuda +// SP13: 4-tuple direction-bias weights → 3-tuple (Hold's 0.5 softening gate removed +// since Hold itself is removed; Long and Short stay full-weight, Flat stays soft). +// OLD: const float dir_bias[4] = {1.0f, 0.5f, 1.0f, 0.0f}; // [Short, Hold, Long, Flat] +const float dir_bias[3] = {1.0f, 1.0f, 0.0f}; // [Short, Long, Flat] +``` + +Verify the loop bounds at `:5159` use `NUM_DIRECTIONS` (now 3). If hardcoded `<4`, change. + +- [ ] **Step 5:** Update Q-head output dim in `dqn_inference.cu` and `dqn_param_layout.rs`: 4 → 3. Bump layout fingerprint seed (invalidates cached checkpoints — greenfield acceptable). + +- [ ] **Step 6:** Update C51/IQN per-direction-bucket loops at the specific sites Explore agent located: + - `c51_loss_kernel.cu:394` — doc comment update + - `c51_loss_kernel.cu:460` — `for (int k=1; k<4; k++)` → `< NUM_DIRECTIONS` + - `c51_loss_kernel.cu:1088` — argmax across direction bins → 3 + - `c51_grad_kernel.cu` — **AUDIT**: grep for `for.*<.*4` direction-axis loops, update each + - `experience_kernels.cu:4059` — direction accumulation loop + - `experience_kernels.cu:5306` — conviction EMA loop + - `experience_kernels.cu:6371-6397` — 4× per-atom direction loops, all need bound change + + Per-direction ISV slot ranges (SP4/SP5/SP6 atoms_pos/tau/...) KEEP 4-entry length; slot index 3 dead-space: + + ```rust + // SP13: NUM_DIRECTIONS reduced 4→3 (Hold removed). Slot index 3 dead-space until + // SP13 close-out recompaction; consumers read by direction index ∈ 0..2 so dead + // slot is never touched. + ``` + +- [ ] **Step 7:** SP5 Pearl 8 trail-distance migration (per-direction ISV slot range stays 4 entries; producer kernel rewrites for new direction indexing): + - `sp5_isv_slots.rs:272` — `TRAIL_DIST_PER_DIR_BASE=270..274` doc-comment update from `[4 dirs: Short, Hold, Long, Flat]` to `[3 dirs at Short=0, Long=1, Flat=2; slot 273 dead-space]` + - `pearl_8_trail_kernel.cu` — producer writes at new indices: Short→slot 270, Long→slot 271 (was Hold), Flat→slot 272 (was Long); slot 273 stays zero + - `state_reset_registry.rs:974` — comment update; reset values stay 4-entry (slot 273 stays sentinel) + - `training_loop.rs:4046-4048` — Pearl 8 `apply_pearls_ad_kernel` calls: 4 → 3 (skip slot 273 — no smoothing of dead slot) + +- [ ] **Step 8:** **Action packing recompacts 12 → 9 slots** (real shape change, not slot-3 dead-space): + - `crates/ml/src/trainers/dqn/financials.rs:27,197` — `exp_idx = dir*3 + mag` semantics: was 4 dirs × 3 mags = 12 indices `[Short[0..3], Hold[3..6], Long[6..9], Flat[9..12]]`; now 3 dirs × 3 mags = 9 indices `[Short[0..3], Long[3..6], Flat[6..9]]` + - `crates/ml/src/trainers/dqn/monitoring.rs:341` — `action_counts[9..12]` Flat-mag indexing: now `action_counts[6..9]` + - `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu:200,245,1199` — action layout `[direction, magnitude, order, urgency]` is 4 SUB-ACTIONS per sample (NOT 4-way direction); only the direction VALUE range shrinks 0..3 → 0..2; `actions[sample * 4 + 0]` indexing stays. Verify consumers expect dir < 3. + - Action-counts buffer allocation: shrink from `12 * batch_size` → `9 * batch_size` (search for allocation site) + +- [ ] **Step 9:** Update HEALTH_DIAG 4-bucket → 3-bucket emit at all sites: + - `metrics.rs:949` — `"val_dir_dist [short={:.4} hold={:.4} long={:.4} flat={:.4}]"` → `"val_dir_dist [short={:.4} long={:.4} flat={:.4}]"` + - `magnitude_distribution.rs:88,217` — `EVAL_DIR_DIST` 4-tuple → 3-tuple + - `distributional_q_tests.rs:790` — `σ_C51 per direction` test output + +- [ ] **Step 10:** Update KL divergence calc if present. Explore agent did NOT find a `kl_divergence_kernel.cu` — verify with `find crates -name "*kl*" -type f`. If absent, this step is a no-op (record in commit message). + +- [ ] **Step 11:** Update tests + build comments: + - `crates/ml/src/trainers/dqn/trainer/tests.rs:542` — `assert_eq!(config.num_actions, 4, …)` → `3` + - `crates/ml/src/trainers/dqn/distributional_q_tests.rs:888` — `const PROD_B0: i32 = 4` → `3` + - `crates/ml/src/cuda_pipeline/thompson_test_kernel.cu` — direction count constant 4 → 3 + - `crates/ml/build.rs:490,496` — comment `"4-thread single-block kernel (one per direction: Short=0, Hold=1, Long=2, Flat=3)"` → `"3-thread … Short=0, Long=1, Flat=2"` + +- [ ] **Step 12:** Add 5 GPU oracle tests for 3-direction action selection (NEW Short-first indexing) in `sp13_phase0_oracle_tests.rs`: + +```rust +// New encoding: Short=0, Long=1, Flat=2 + +#[test] fn action_select_3way_argmax_short() { /* q=[2.0, 1.0, 0.5] → 0 (Short wins) */ } +#[test] fn action_select_3way_argmax_long() { /* q=[1.0, 2.0, 0.5] → 1 (Long wins) */ } +#[test] fn action_select_3way_argmax_flat() { /* q=[0.5, 0.3, 2.0] → 2 (Flat wins) */ } +#[test] fn action_select_no_hold_leaks() { + // 100 random q-vectors → 0 ≤ action < 3 always; never 3 +} +#[test] fn action_packing_9_slot_invariant() { + // Pack (dir, mag) for all 9 (dir, mag) combinations; verify exp_idx ∈ [0, 9) +} +``` + +- [ ] **Step 13:** Stale doc cleanup — implementer audit: + +```bash +grep -rn "Short=0.*Hold=1.*Long=2.*Flat=3\|S=Short.*H=Hold.*L=Long.*F=Flat\|\[Short, Hold, Long, Flat\]" \ + crates/ --include="*.rs" --include="*.cu" --include="*.cuh" \ + | grep -v "// SP13:" +``` + +Update narrative comments to `Short=0, Long=1, Flat=2` ordering. Specific known sites: +- `crates/ml/src/trainers/dqn/monitoring.rs:13` +- `crates/ml/src/trainers/dqn/trainer/mod.rs:1966` +- 20+ kernel header comments + +- [ ] **Step 14:** Sanity-check no Hold leaks: + +```bash +grep -rn "DIRECTION_HOLD\|DirectionAction::Hold\|::Hold\b\|hold=\|action == 3\|< 4.*direction\|dir.*== 1.*Hold" \ + crates/ --include="*.rs" --include="*.cu" --include="*.cuh" \ + | grep -v test | grep -v "// SP13" +``` + +Expected: zero hits. + +- [ ] **Step 15:** Stage all P0a.T3 changes (atomic with T1, T2, T4 in P0a.T5 commit). + +--- + +### Task P0a.T4: Wire dir_acc launch + ISV write + HEALTH_DIAG + +**Files:** Modify `training_loop.rs` (validation/eval step where aux head is forward'd) + +- [ ] **Step 1:** Locate existing aux HEALTH_DIAG site: + +```bash +grep -n "aux_nb_mse\|aux_rg_ce" crates/ml/src/trainers/dqn/trainer/training_loop.rs +``` + +- [ ] **Step 2:** After existing aux forward on validation batch, add: + +```rust +// SP13 P0a: directional accuracy probe (regression-mode aux scalar in P0a; Layer B → softmax) +self.launch_aux_dir_acc_reduce( + &fused.aux_pred, + &fused.next_bar_label, + batch_size as i32, + self.aux_dir_acc_buf.device_slice_mut(), +)?; + +// Pearls A+D smoothing on slot 0 (dir_acc) — short EMA (α=0.3) → slot 373 +self.launch_apply_pearls_ad( + self.aux_dir_acc_buf.device_slice_at(0, 1), + 1, + AUX_DIR_ACC_SHORT_EMA_INDEX, + /* alpha */ 0.3, +)?; +// Slow EMA (α=0.05) → slot 374 (stagnation detector) +self.launch_apply_pearls_ad( + self.aux_dir_acc_buf.device_slice_at(0, 1), + 1, + AUX_DIR_ACC_LONG_EMA_INDEX, + /* alpha */ 0.05, +)?; + +// Per-bar logit signal write for direction Q-head input (slot 375) +// In P0a: regression scalar squashed to [-1, +1] via tanh; Layer B replaces with softmax logit_diff +self.launch_aux_pred_to_isv_tanh( + &fused.aux_pred, + AUX_DIR_PREDICTION_INDEX, +)?; + +let dir_acc_arr: [f32; 3] = self.aux_dir_acc_buf.read_pinned()?; +tracing::info!( + "HEALTH_DIAG aux_dir_acc accuracy={:.4} pos_pred={:.4} pos_label={:.4}", + dir_acc_arr[0], dir_acc_arr[1], dir_acc_arr[2], +); +``` + +(Adjust `apply_pearls_ad` signature to match project conventions; if it takes a single α with two slots needing different rates, write a small helper or two calls.) + +The `launch_aux_pred_to_isv_tanh` is a 1-line per-bar kernel — write it inline if no +existing helper fits, ~20 LOC. + +- [ ] **Step 3:** Verify compile + clippy clean. + +- [ ] **Step 4:** Stage. + +--- + +### Task P0a.T5: Atomic Phase 0a commit + smoke + +- [ ] **Step 1:** Final compile + tests: + +```bash +SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail +SQLX_OFFLINE=true cargo build -p ml --release 2>&1 | grep -E "error|warning" | head +SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp13_phase0_oracle_tests +SQLX_OFFLINE=true cargo test -p ml --features gpu-tests --test sp12_reward_math_tests # SP12 still green +``` + +- [ ] **Step 2:** Sanity grep for Hold leaks (Step 10 of T3) — must be zero. + +- [ ] **Step 3:** Commit: + +```bash +git commit -m "$(cat <<'EOF' +feat(sp13): P0a — Hold elimination + dir_acc instrumentation (atomic) + +Tests user's hypothesis (Hold hides directional signal) by removing Hold from the +direction action space. 4 → 3-way: NUM_DIRECTIONS, DirectionAction enum {Long, +Short, Flat}, Q-head output dim, C51/IQN per-direction-bucket state, action +selection argmax, replay buffer action encoding, HEALTH_DIAG val_dir_dist 3-bucket. + +Adds aux_dir_acc_reduce GPU kernel + 2 EMA slots (short α=0.3 / long α=0.05) for +dir_acc tracking, plus the aux logit-signal slot wired to ISV[375] via tanh of +the existing regression head (Layer B will replace with softmax logit_diff). + +8 SP13 ISV slots [372..380); SP5_SLOT_END = 380. Per-direction slot 3 dead-space +documented; recompaction deferred to close-out. + +Greenfield: replay buffer / checkpoints invalidated by design. + +Tests: 10 GPU oracle tests in sp13_phase0_oracle_tests (6 dir_acc + 4 action select). + +Spec: docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md (v2) + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +- [ ] **Step 4:** Push + 10-epoch L40S smoke: + +```bash +git push --set-upstream origin sp11-reward-as-controlled-subsystem +./scripts/argo-train.sh --model dqn --commit $(git rev-parse HEAD) --epochs 10 --pool ci-training-l40s +``` + +Capture workflow name. Monitor early per `feedback_kill_runs_on_anomaly_quickly`: +- ep1 HEALTH_DIAG `aux_dir_acc` line emitting (slot 373 wired) +- val_dir_dist 3-bucket (no `hold=` field) +- aux_w stays at SP11 inverted-formula values (~0.09) — controller unchanged in P0a +- No NaN/inf + +Kill on anomaly. + +- [ ] **Step 5:** Phase 0a gate decision (post-completion): + +```bash +argo logs -n foxhunt --tail=2000 | grep -E "HEALTH_DIAG.*epoch=9|aux_dir_acc|val_dir_dist|win_rate" | tail -30 +``` + +Apply gate: + +| `dir_acc` trajectory | `dir_acc[ep9]` | `win_rate[ep9]` | Decision | +|---|---|---|---| +| ↑ rising | > 0.52 | > 0.51 | ✅ **CONFIRMED** — proceed Layer B (skip P0b) | +| flat ~0.50 | < 0.51 | < 0.50 | proceed P0b (test aux amplification) | +| mixed (one ↑, one flat) | — | — | proceed P0b | + +Update `project_hold_hides_directional_signal_hypothesis.md` with smoke ID, commit SHA, final dir_acc and WR. + +If CONFIRMED → skip to **Layer B**. If proceeding to P0b → next section. + +--- + +## Phase 0b — Aux amplification probe (1 commit, only on partial/red P0a) + +### Task P0b.T1: Replace SP11 inverted aux_w with deficit + stagnation controller + +**Files:** Modify `training_loop.rs:4103-4104`; modify `sp13_isv_slots.rs` to add base/decay constants + +- [ ] **Step 1:** Add constants to `sp13_isv_slots.rs`: + +```rust +pub const AUX_W_BASE: f32 = 0.5; +pub const AUX_W_DEFICIT_GAIN: f32 = 5.0; // aux_w = base × (1 + 5 × deficit) +pub const AUX_W_STAGNATION_DECAY: f32 = 0.7; // (1 - 0.7 × stagnation) +pub const AUX_W_HARD_FLOOR_RATIO: f32 = 0.3; // never below 0.3 × base +pub const AUX_W_HARD_CEIL_RATIO: f32 = 3.0; // never above 3.0 × base +``` + +- [ ] **Step 2:** Replace the SP11-era formula at `training_loop.rs:4103-4104`: + +```rust +// SP13 P0b: deficit + stagnation aux_w controller. Replaces SP11 inverted +// `0.1 × health × (1 - sharpe)` formula which lowered aux_w when sharpe was +// high — exactly wrong for the failure mode (high sharpe + low WR). +let target = self.read_isv_slot(TARGET_DIR_ACC_INDEX); // 0.55 +let short_ema = self.read_isv_slot(AUX_DIR_ACC_SHORT_EMA_INDEX); // fast (α=0.3) +let long_ema = self.read_isv_slot(AUX_DIR_ACC_LONG_EMA_INDEX); // slow (α=0.05) + +let deficit = (target - short_ema).max(0.0); +let improvement = (short_ema - long_ema).max(0.0); +// stagnation = 1 means no improvement; 0 means improving fast +let stagnation = if deficit > 0.005 { + (1.0 - improvement / deficit).clamp(0.0, 1.0) +} else { + 0.0 // already at target — not stagnant, just satisfied +}; + +let aux_w_raw = AUX_W_BASE + * (1.0 + AUX_W_DEFICIT_GAIN * deficit) + * (1.0 - AUX_W_STAGNATION_DECAY * stagnation); +let aux_w = aux_w_raw.clamp( + AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO, + AUX_W_BASE * AUX_W_HARD_CEIL_RATIO, +); +fused.set_aux_weight(aux_w); + +tracing::debug!( + "aux_w controller: target={:.3} short={:.3} long={:.3} deficit={:.4} stag={:.3} aux_w={:.3}", + target, short_ema, long_ema, deficit, stagnation, aux_w, +); +``` + +Imports needed: +```rust +use crate::cuda_pipeline::sp13_isv_slots::{ + TARGET_DIR_ACC_INDEX, AUX_DIR_ACC_SHORT_EMA_INDEX, AUX_DIR_ACC_LONG_EMA_INDEX, + AUX_W_BASE, AUX_W_DEFICIT_GAIN, AUX_W_STAGNATION_DECAY, + AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO, +}; +``` + +- [ ] **Step 3:** Verify compile + 10-bar simulation in head: + - At fold start: short=long=0.5, deficit=0.05, improvement=0, stagnation=1.0, aux_w = 0.5 × 1.25 × 0.3 = 0.1875 → clamped to 0.15 (floor) — gentle start, controller waits for data + - After improvement (short=0.55, long=0.50): deficit=0, improvement=0.05, aux_w = 0.5 × 1.0 × 1.0 = 0.5 — back to base, satisfied + - After stagnation (short=0.50, long=0.50, target=0.55): deficit=0.05, improvement=0, stagnation=1.0, aux_w = floor 0.15 — backed off, no destabilisation + +- [ ] **Step 4:** Commit: + +```bash +git commit -m "$(cat <<'EOF' +feat(sp13): P0b — aux_w deficit + stagnation controller (replaces SP11 inverted) + +Replaces SP11-era `aux_w = 0.1 × health × (1 - sharpe_tanh).clamp(0.05, 0.3)` at +training_loop.rs:4103-4104 with deficit-driven amplification + stagnation decay. + +Mechanism: + deficit = max(0, target_dir_acc - short_ema) + improvement = max(0, short_ema - long_ema) + stagnation = (deficit > 0) ? clamp(1 - improvement/deficit, 0, 1) : 0 + aux_w = base × (1 + 5 × deficit) × (1 - 0.7 × stagnation), bounded [0.3×base, 3×base] + +Prevents the v1 failure mode where aux_w would push to hard cap forever in the +data-limited case (no signal → no improvement → controller stuck at 1.5× base +destabilising Q-head). Stagnation decay backs off when controller can't make progress. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +git push +``` + +- [ ] **Step 5:** 10-epoch L40S smoke. Same gate as P0a. If still red → halt SP13, update hypothesis memory to DISCONFIRMED, plan SP14 separately. + +--- + +## Layer B — Aux head regression → binary classification (1 atomic commit, only on green Phase 0) + +### Task B.T1: Aux head softmax + cross-entropy + +**Files:** `aux_head_kernel.cu`, `dqn_param_layout.rs` + +- [ ] **Step 1:** Update aux head forward: + +```cuda +// Output dim 1 (regression scalar) → 2 (softmax logits over [neg, pos]) +extern "C" __global__ void aux_head_forward_kernel( + /* ... existing inputs ... */ + float* __restrict__ aux_softmax_out // [B, 2] +) { + /* ... existing trunk → linear → 2 logits ... */ + float logit_neg = /* ... */; + float logit_pos = /* ... */; + float m = fmaxf(logit_neg, logit_pos); + float exp_neg = expf(logit_neg - m); + float exp_pos = expf(logit_pos - m); + float Z = exp_neg + exp_pos; + aux_softmax_out[b * 2 + 0] = exp_neg / Z; + aux_softmax_out[b * 2 + 1] = exp_pos / Z; +} +``` + +- [ ] **Step 2:** Update aux head backward — CrossEntropy gradient is `softmax - one_hot(label)`: + +```cuda +// d_logit[c] = softmax[c] - (label == c ? 1 : 0) +``` + +- [ ] **Step 3:** Update param shapes: aux head output features 1 → 2 in `dqn_param_layout.rs`. Bump layout fingerprint seed. + +- [ ] **Step 4:** Stage. + +--- + +### Task B.T2: Label generation kernel (sign of next-30-bar return) + +**Files:** `aux_label_gen_kernel.cu` (modify or create) + +- [ ] **Step 1:** Update / write kernel: + +```cuda +extern "C" __global__ void aux_label_gen_kernel( + const float* __restrict__ price, // [B, T] + int t_index, + int horizon, // 30 bars from spec + int t_max, // T (sequence length) + int* __restrict__ label_class_out, // [B] {0=negative, 1=positive}; bars too close to end set to -1 (skip) + int* __restrict__ label_valid_out // [B] 0/1 mask +) { + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= /* B */) return; + int t_target = t_index + horizon; + if (t_target >= t_max) { + label_class_out[b] = -1; + label_valid_out[b] = 0; + return; + } + float ret = price[b * t_max + t_target] - price[b * t_max + t_index]; + if (ret > 0.0f) { + label_class_out[b] = 1; + label_valid_out[b] = 1; + } else if (ret < 0.0f) { + label_class_out[b] = 0; + label_valid_out[b] = 1; + } else { + label_class_out[b] = -1; // skip flat bars + label_valid_out[b] = 0; + } +} +``` + +CrossEntropy loss masks invalid bars (multiply by `label_valid_out`). + +**No information leakage**: aux head sees state s_t at inference; only training-time +labels use price[t+30]. Same hindsight pattern as existing next_bar_mse. + +- [ ] **Step 2:** Wire label gen launch in training loop (replaces existing next-bar regression label gen). + +- [ ] **Step 3:** Adapt `aux_dir_acc_reduce_kernel` to read 2-class softmax: + +```cuda +// Layer B: softmax read instead of regression scalar +const float p_pos = aux_softmax[i * 2 + 1]; +const float p_neg = aux_softmax[i * 2 + 0]; +const int pred_pos = (p_pos > p_neg) ? 1 : 0; +// label comparison stays the same (sign-encoded label still works) +``` + +- [ ] **Step 4:** Stage. + +--- + +### Task B.T3: Wire aux logit_diff → ISV[375] → direction Q-head input + +**Files:** `training_loop.rs`, `dqn_inference.cu`, `dqn_param_layout.rs` + +- [ ] **Step 1:** Replace P0a `tanh(scalar)` write with softmax logit_diff: + +```rust +self.launch_aux_softmax_to_logit_diff_isv( + &fused.aux_softmax, // [B, 2] + AUX_DIR_PREDICTION_INDEX, +)?; +``` + +Tiny kernel: +```cuda +extern "C" __global__ void aux_softmax_to_logit_diff_kernel( + const float* __restrict__ aux_softmax, // [B, 2] + int batch_size, + float* __restrict__ isv, + int isv_stride +) { + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= batch_size) return; + float diff = aux_softmax[b * 2 + 1] - aux_softmax[b * 2 + 0]; // ∈ [-1, +1] + isv[b * isv_stride + AUX_DIR_PREDICTION_INDEX] = diff; +} +``` + +- [ ] **Step 2:** Concat `ISV[AUX_DIR_PREDICTION_INDEX]` into direction Q-head input: + +In `dqn_inference.cu`, locate the direction Q-head forward — input layer was `[trunk_features]`, becomes `[trunk_features, aux_logit_diff]`. In_dim grows by 1. + +In `dqn_param_layout.rs`, update direction Q-head weight shape: `(in_dim, 3)` → `(in_dim + 1, 3)`. Re-init. + +- [ ] **Step 3:** Add 3 GPU oracle classification tests: + +```rust +#[test] fn aux_softmax_logit_diff_positive_class() { /* logits=[-1, 1] → diff>0 */ } +#[test] fn aux_softmax_logit_diff_negative_class() { /* logits=[1, -1] → diff<0 */ } +#[test] fn aux_cross_entropy_grad_sign_correct() { /* softmax-onehot sign per class */ } +``` + +- [ ] **Step 4:** Atomic Layer B commit: + +```bash +git commit -m "$(cat <<'EOF' +feat(sp13): Layer B — aux head regression → 30-bar binary classification + +Replaces aux MSE on next-bar return with CrossEntropy on next-30-bar return sign. +Aux output dim 1 → 2 + softmax. logit_diff = softmax[1] - softmax[0] ∈ [-1, +1] +written to ISV[AUX_DIR_PREDICTION_INDEX] each bar, concatenated into direction +Q-head input layer (in_dim grows by 1). aux_dir_acc_reduce_kernel adapted to read +softmax instead of regression scalar. + +No information leakage: aux uses state s_t at inference; price[t+30] only at +training time (hindsight labels). Same pattern as the existing aux MSE head. + +Tests: 13 in sp13_phase0_oracle_tests (10 P0a + 3 classification). + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +git push +``` + +- [ ] **Step 5:** 5-epoch L40S smoke verifying: + - Aux CrossEntropy loss decreasing + - val_dir_acc continues at or above Phase 0 level + - WR not regressing + - aux signal feeding Q-head: perturbation test (zero out ISV[375] for 1 bar in eval, observe direction action distribution shift) + +If clean → proceed Layer C. + +--- + +## Layer C — Direction-skill bonus + luck discount (1 atomic commit) + +### Task C.T1: Device-inline reward functions in trade_physics.cuh + +**Files:** `trade_physics.cuh` + +- [ ] **Step 1:** Add 2 device-inline functions: + +```cuda +__device__ __forceinline__ float compute_direction_skill_bonus_bounded( + int action_dir, // 0=Long, 1=Short, 2=Flat + float realized_30_bar_return, // sign-encoded + float alpha_capped, // SP12 asymmetric-capped trade pnl, used for cap calc + float bonus_alpha_param, // ISV[376], default 1.0 + float bonus_beta_param, // ISV[377], default 1.0 + float cap_ratio // ISV[379], default 0.3 +) { + if (action_dir == DIRECTION_FLAT) return 0.0f; + int trade_sign = (action_dir == DIRECTION_LONG) ? +1 : -1; + int label_sign = (realized_30_bar_return > 0.0f) ? +1 + : ((realized_30_bar_return < 0.0f) ? -1 : 0); + if (label_sign == 0) return 0.0f; // no signal to score + float skill_cap = cap_ratio * fabsf(alpha_capped); + if (trade_sign == label_sign) { + return +fminf(bonus_alpha_param, skill_cap); + } else { + return -fminf(bonus_beta_param, skill_cap); + } +} + +__device__ __forceinline__ float compute_luck_win_discount( + float reward_pre_discount, + float realized_trade_pnl, + bool direction_was_correct, + float discount_factor // ISV[378], default 0.3 +) { + bool trade_won = realized_trade_pnl > 0.0f; + if (trade_won && !direction_was_correct) { + return reward_pre_discount * discount_factor; + } + return reward_pre_discount; +} +``` + +- [ ] **Step 2:** Stage. + +--- + +### Task C.T2: Replace exit-event reward composition in experience_kernels.cu + +**Files:** `experience_kernels.cu` (~line 2788, SP12 exit-event site) + +- [ ] **Step 1:** Locate the SP12 reward composition site: + +```bash +grep -n "fmaxf(-10\|asymmetric_capped\|min_hold_factor\|compute_min_hold_penalty" \ + crates/ml/src/cuda_pipeline/experience_kernels.cu +``` + +- [ ] **Step 2:** Layer in skill bonus + luck discount AFTER SP12's asymmetric cap, BEFORE min_hold_factor multiplication: + +```cuda +// SP12 asymmetric cap (preserved per pearl_audit_unboundedness_for_implicit_asymmetry) +float alpha_capped = compute_asymmetric_capped_pnl(realized_trade_pnl, -10.0f, +5.0f); + +// SP13 Layer C: direction-skill bonus (bounded relative to alpha) +int trade_sign = (trade_direction == DIRECTION_LONG) ? +1 + : ((trade_direction == DIRECTION_SHORT) ? -1 : 0); +int label_sign_for_30bar = (realized_30_bar_return > 0.0f) ? +1 + : ((realized_30_bar_return < 0.0f) ? -1 : 0); +bool dir_correct = (trade_sign != 0) && (trade_sign == label_sign_for_30bar); + +float skill_bonus = compute_direction_skill_bonus_bounded( + trade_direction, + realized_30_bar_return, + alpha_capped, + isv[DIR_SKILL_BONUS_ALPHA_INDEX], + isv[DIR_SKILL_BONUS_BETA_INDEX], + isv[SKILL_BONUS_CAP_RATIO_INDEX] +); + +float r_pre_discount = alpha_capped + skill_bonus; + +// SP13 Layer C: luck-win discount +float r_post_discount = compute_luck_win_discount( + r_pre_discount, + realized_trade_pnl, + dir_correct, + isv[LUCK_WIN_DISCOUNT_INDEX] +); + +// SP12 min-hold penalty (preserved — orthogonal to skill incentive) +float r = r_post_discount * min_hold_factor; +``` + +(Implementer adapts variable names to actuals at the call site.) + +- [ ] **Step 3:** Stage. + +--- + +### Task C.T3: GPU oracle tests — full quadrant matrix + +**Files:** Create `sp13_reward_math_test_kernel.cu` (mirrors SP12 pattern), Create `sp13_reward_math_tests.rs`, Modify `build.rs` + +- [ ] **Step 1:** Test wrapper kernel exposing the 2 device-inline funcs to host (mirror `sp12_reward_math_test_kernel.cu`). + +- [ ] **Step 2:** 16 oracle tests covering the 8-quadrant matrix from spec PLUS edge cases: + +```rust +// All defaults: α=β=1.0, cap_ratio=0.3, luck=0.3, SP12 cap [-10,+5] +// Direction encoding (post-P0a): Short=0, Long=1, Flat=2 + +#[test] +fn correct_long_winner_big() { + // alpha=+5, dir=Long(1), label=+1 → bonus=+min(1.0, 1.5)=+1.0 + let r = run_skill(/*alpha=*/5.0, /*dir=*/1, /*label=*/1.0); + assert!((r - 6.0).abs() < 1e-6); +} + +#[test] +fn correct_long_winner_small() { + // alpha=+0.5, dir=Long(1), bonus_cap = 0.3 × 0.5 = 0.15 + let r = run_skill(0.5, 1, 1.0); + assert!((r - 0.65).abs() < 1e-6); +} + +#[test] +fn correct_long_loser_small_still_negative() { + // alpha=-0.5, dir=Long(1), bonus=+0.15 → r=-0.35 (NOT positive — verify NO gaming) + let r = run_skill(-0.5, 1, 1.0); + assert!(r < 0.0, "correct-long-loser must stay negative; got {}", r); +} + +#[test] +fn correct_long_loser_big_dominated_by_loss() { + // alpha=-10, dir=Long(1), bonus_cap=3.0 → bonus=+1.0 → r=-9.0 + let r = run_skill(-10.0, 1, 1.0); + assert!((r - (-9.0)).abs() < 1e-6); +} + +#[test] +fn wrong_long_winner_lucky_discounted() { + // alpha=+5, dir=Long(1), label=-1 → bonus=-1.0; lucky win → × 0.3 + let r = run_full(/*pnl_realized=*/2.0, 5.0, 1, -1.0); // realized > 0, dir wrong + let pre = 5.0 - 1.0; // 4.0 + let expected = pre * 0.3; // 1.2 + assert!((r - expected).abs() < 1e-6); +} + +#[test] +fn wrong_long_loser() { + // alpha=-10, dir=Long(1), label=-1 → bonus=-1.0 → r=-11.0 (no luck discount on loss) + let r = run_full(-5.0, -10.0, 1, -1.0); + assert!((r - (-11.0)).abs() < 1e-6); +} + +#[test] +fn flat_zero_skill_bonus() { + let r = run_skill(2.0, 2 /* Flat */, 1.0); + assert!((r - 2.0).abs() < 1e-6); // alpha + 0 bonus +} + +#[test] +fn zero_label_zero_skill_bonus() { + let r = run_skill(2.0, 1 /* Long */, 0.0); // label=0 → no signal to score + assert!((r - 2.0).abs() < 1e-6); +} + +#[test] +fn correct_short_winner() { + let r = run_skill(5.0, 0 /* Short */, -1.0); // short + falling = correct + assert!((r - 6.0).abs() < 1e-6); +} + +#[test] +fn wrong_short_winner_lucky_discounted() { + // dir=Short(0), label=+1 → wrong; pnl_realized > 0 → lucky → × 0.3 + let r = run_full(2.0, 5.0, 0 /* Short */, 1.0); + let pre = 5.0 - 1.0; + let expected = pre * 0.3; + assert!((r - expected).abs() < 1e-6); +} + +#[test] +fn no_negative_ev_gaming_invariant() { + // Across many small-alpha correct-direction losers, reward must stay < 0 + for alpha_loss in (-10..=-1).map(|i| i as f32 * 0.01) { // -0.1, -0.09, ... + let r = run_skill(alpha_loss, 1 /* Long */, 1.0); + assert!(r < 0.0, + "negative-EV gaming invariant violated: alpha={} reward={}", + alpha_loss, r); + } +} + +// + 6 more covering: bonus magnitude calibration, cap ratio extremes, +// luck discount boundary (pnl exactly 0), Flat bypass invariant, etc. +``` + +- [ ] **Step 3:** Run on RTX 3050 Ti, expect 16/16 passing. + +- [ ] **Step 4:** Add 4 ISV reset entries (slots 376, 377, 378, 379) to `state_reset_registry.rs` (per-fold). + +- [ ] **Step 5:** Add per-fold HEALTH_DIAG metric: + +```rust +// At fold-end, after all per-fold metrics: +let bh_pnl = price_at_fold_end - price_at_fold_start; +let alpha_vs_bh = total_pnl - bh_pnl; +tracing::info!( + "HEALTH_DIAG fold_alpha_vs_buy_and_hold fold={} alpha={:.4} bh_pnl={:.4} our_pnl={:.4}", + fold_idx, alpha_vs_bh, bh_pnl, total_pnl, +); +``` + +(Metric only — never feeds into reward. Per Major-1 fix.) + +- [ ] **Step 6:** Atomic Layer C commit: + +```bash +git commit -m "$(cat <<'EOF' +feat(sp13): Layer C — direction-skill bonus + luck-win discount (atomic) + +Adds bounded skill bonus and luck-win discount on top of SP12 asymmetric-capped +P&L. Per spec v2 Major-3 fix, skill bonus is bounded relative to |alpha| via +cap_ratio (default 0.3) to prevent reward gaming on small-alpha correct-direction +losers. + + alpha_capped = SP12 cap [-10, +5] + skill_bonus = sign × min(α_param, cap_ratio × |alpha|) if action ∈ {L,S} & label ≠ 0 else 0 + luck_disc = (won AND wrong_dir) ? × 0.3 : 1.0 + reward = (alpha_capped + skill_bonus) × luck_disc × min_hold_factor + +4 new ISV slots [376, 377, 378, 379] populated in state-reset registry. Tests: +16-test quadrant matrix in sp13_reward_math_tests verifying no negative-EV gaming. + +Per-fold buy-and-hold delta added to HEALTH_DIAG as METRIC ONLY (per spec v2 +Major-1: alpha-vs-benchmark dropped as reward signal due to mathematical +tautology against always-long benchmark). + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +git push +``` + +--- + +## Layer D — 30-epoch validation + close-out + +### Task D.T1: 30-epoch L40S validation + +- [ ] **Step 1:** Push + deploy: + +```bash +./scripts/argo-train.sh --model dqn --commit $(git rev-parse HEAD) --epochs 30 --pool ci-training-l40s +``` + +- [ ] **Step 2:** Optional 3-action SP12 control run (~€1) for fair before/after baseline: + +Check out a parallel branch from pre-Layer-B SP12 + 3-action commit, deploy 30-epoch, compare. + +- [ ] **Step 3:** Monitor early; kill on anomaly. Capture final-epoch HEALTH_DIAG metrics. + +--- + +### Task D.T2: Close-out commit (audit doc + 3 new pearls + memory updates) + +**Files:** `docs/dqn-wire-up-audit.md`, 3 new pearl files, `MEMORY.md`, `project_hold_hides_directional_signal_hypothesis.md` + +- [ ] **Step 1:** Audit doc — SP13 entry: hypothesis, P0a result, P0b result (if run), Layer B/C deltas, 30-epoch metrics, control comparison (if run). + +- [ ] **Step 2:** Write 3 pearls: + +- `pearl_redefine_success_for_predictive_skill.md` — when policy plateaus on a structural-extraction objective across many tunes (12 across SP1-SP11 didn't move 46% WR), the fix is redefining the success criterion, not tuning. Canonical SP13. + +- `pearl_skill_bonus_must_be_alpha_bounded.md` — direction-skill bonuses must be bounded relative to alpha magnitude (e.g. `min(α_param, ratio × |alpha|)`) to prevent reward gaming on small-alpha correct-direction losers. Worked example from SP13 v1 review (where unbounded α=1.0 + 0.5 alpha = -0.5 trade gave +0.5 reward, model rationally entered expected losers). + +- `pearl_reward_quadrant_audit_required.md` — meta-pearl: any reward-composition change must include a 4-quadrant worked-example table (correct/wrong × winner/loser) showing all signs and the no-negative-EV-gaming invariant. Caught the SP13 v1 alpha-vs-benchmark math tautology and the unbounded skill bonus gaming. + +- [ ] **Step 3:** Update `MEMORY.md` index: 3 new pearls under "Architectural constraints". + +- [ ] **Step 4:** Update `project_hold_hides_directional_signal_hypothesis.md`: +- Status: UNVERIFIED → CONFIRMED / DISCONFIRMED / PARTIAL +- Phase 0a smoke ID, commit SHA, dir_acc, WR, val_dir_dist +- Phase 0b smoke if run + +- [ ] **Step 5:** Commit + push. + +--- + +## Estimated cost + +| Phase | LOC | Time | Cost | +|---|---|---|---| +| P0a | ~300 | 4-5h | local | +| P0a smoke | — | 30min | ~€0.60 (10-epoch) | +| P0b (conditional) | ~80 | 1h | local | +| P0b smoke (conditional) | — | 30min | ~€0.60 | +| Layer B | ~120 | 2-3h | + ~€0.30 (5-epoch) | +| Layer C | ~150 | 2-3h | local | +| Layer D | — | 1h | ~€1.00 (30-epoch) + ~€1.00 control (recommended) | +| **Best case** (P0a green, no P0b) | **~570** | **~10h** | **~€2.90** | +| **Worst confirm** (P0a partial → P0b green) | **~650** | **~12h** | **~€3.50** | +| **Disconfirm** (both red) | **~380** | **~6h** | **~€1.20** | + +Cheap-fail at ~€1.20 if hypothesis disconfirmed. Confirm cost ~€3 paid against the +chance of finding 55%+ WR (the prize across the entire SP1-SP12 lineage). diff --git a/docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md b/docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md new file mode 100644 index 000000000..e375f4a07 --- /dev/null +++ b/docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md @@ -0,0 +1,386 @@ +# SP13 — Redefine Success for Predictive Skill (v2) + +**Date**: 2026-05-04 (v2 — supersedes v1 in-place) +**Branch**: continues `sp11-reward-as-controlled-subsystem` (single coherent architectural lineage) +**Builds on**: SP11 (cap stability) + SP12 (event-driven density) + asymmetric loss aversion +**Anchor pearls**: +- `pearl_audit_unboundedness_for_implicit_asymmetry` +- `pearl_event_driven_reward_density_alignment` +- (NEW) `pearl_redefine_success_for_predictive_skill` — to be written post-validation +- (NEW) `pearl_skill_bonus_must_be_alpha_bounded` — to be written post-validation +- (NEW) `pearl_reward_quadrant_audit_required` — meta-pearl from v1 review + +## v1 → v2 changes (reasons here for traceability) + +| Issue | v1 | v2 | +|---|---|---| +| Alpha-vs-benchmark math broken — Long trades have alpha = −costs always (tautology against always-long benchmark) | reward = realized_pnl − benchmark_pnl_same_period | **Dropped.** Predictive-skill incentive carried by direction-skill bonus + correct-call asymmetry. Per-fold buy-and-hold delta tracked as HEALTH_DIAG metric only. | +| Phase 0 confounded Hold removal AND aux_w boost — couldn't tell which helped | Single P0 commit | **Split:** P0a Hold-only (cheap test of user hypothesis), P0b aux probe (only if 0a partial) | +| Skill bonus + luck-discount could produce positive reward for correct-direction expected-loser → reward gaming | α=β=1.0 raw constants | **Bounded:** `skill_bonus = sign × min(α_param, ratio × |alpha|)` with ratio=0.3 default. Worked examples below verify no negative-EV gaming. | +| 5-epoch gate too short for aux to converge → false-negative risk | 5-epoch absolute threshold | **10-epoch trajectory criterion:** dir_acc trending up AND final > 0.52 | +| Aux_w controller at hard cap forever in data-limited case → destabilises Q | `clamp(0.05, 1.5)` only | **Stagnation detector:** dual-EMA (fast/slow); if slow ≈ fast for ≥ N epochs, decay aux_w toward base | +| No worked examples of full reward composition | Inline formulas | **Worked-example matrix** for 4 trade-outcome quadrants (correct/wrong × winner/loser) | + +## TL;DR + +The reward function defines what "success" means to the model. SP1-SP12 lineage rewarded +**any positive bounded P&L** — including P&L from structural extraction (trail-stop +arithmetic on directionally-random bets at 46% win rate). SP13 rewards **predictive +skill specifically** via two surgical changes: + +1. **Eliminate Hold action** — model must commit. Tests user's hypothesis that Hold's + availability has been hiding directional signal across the entire SP1-SP12 lineage. +2. **Direction-skill bonus + correct-call asymmetry** — bounded explicit reward for + predicting direction correctly, bounded penalty for being wrong, and discount on + wrong-direction lucky wins. Mathematically calibrated so the model cannot game by + entering expected-loser trades. + +Plus three structural supports: + +3. **Aux head regression → binary classification** on next-30-bar return sign, wired + into direction Q-head input. +4. **ISV-driven aux_w controller** (replaces SP11 inverted formula) with deficit-driven + amplification AND stagnation detector to prevent destabilization in data-limited case. +5. **Per-fold buy-and-hold delta** tracked as HEALTH_DIAG metric (NOT reward signal). + +~600-750 LOC across 4 commits (P0a, P0b, B, C). Greenfield refactor — replay buffer +compatibility intentionally broken. + +**Hypothesis under test** (user, 2026-05-04): directional signal IS in the data; ~46% WR +persists because Hold hides it. P0a smoke decides cleanly: +- WR climbs > 0.52 with Hold removed → ✅ confirmed, hypothesis was right +- WR stays flat at 0.46-0.50 → P0b adds aux amplification as second probe +- Both P0a + P0b flat → ❌ data limit, halt SP13 + +## Background — why current system fails + +50-epoch validation `train-multi-seed-pmbwn` and 22-epoch SP12 v3 validation +`train-multi-seed-vd2v5`: + +- val_dir_dist: hold=0.45 (lazy default), Long+Short=0.36 (active bets) +- win_rate=0.4651 (slightly below random for direction) +- profit_factor=1.24 (winners 24% bigger than losers) +- aux_w stuck at 0.09 (suppressed by SP11-era `0.1 × health × (1 − sharpe)` controller) +- aux next_bar_mse=3.83 (regression error suggests poor signal extraction OR aux never got real gradient) + +**The mechanism**: model picks Hold ~half the time. Of the ~37% directional bets, only +46% are correct. Profitability comes from trail-stop arithmetic, not predictive skill. + +## The architectural shift + +| Current "success" | Proposed "success" | +|---|---| +| Maximize positive bounded P&L | Maximize bounded P&L + bounded direction-skill credit | +| Hold is safe default | Must commit (no Hold option) | +| Wrong direction + lucky exit = full reward | Wrong direction + win = 0.3× reward (luck discount) | +| Aux head suppressed when sharpe high | Aux amplified when WR below target, decayed when stagnant | +| 1-bar reward horizon (next-bar MSE) | 30-bar prediction horizon (sign classification) | + +## The five changes + +### Change 1: Eliminate Hold action + +**Current**: 4-way direction `{Short=0, Hold=1, Long=2, Flat=3}` × magnitude × ord × urg. +**New**: 3-way `{Short=0, Long=1, Flat=2}` — Hold removed, Short stays at 0 (preserves +canonical anchor used across 30+ kernel comments and tests), Long shifts down 2 → 1, +Flat shifts down 3 → 2. + +**Action semantics in 3-way world** (replaces `project_hold_action_design`): +- `Long` at t when prior position = Flat → open Long +- `Long` at t when prior position = Long → keep Long (same as old "Hold-Long") +- `Long` at t when prior position = Short → flip to Long +- `Flat` at t when any position → close all positions +- `Short` symmetric + +`Long` is no longer "open Long" specifically; it is "have Long position". `Long` is a +state, not a verb. + +**Affected** (atomic per `feedback_no_partial_refactor`, ~50–70 touch-sites mapped by +Explore agent 2026-05-04): +- Canonical Rust enum at `crates/ml-core/src/common/action.rs:124-132` and + `NUM_DIRECTIONS: 4 → 3` at `crates/ml-core/src/state_layout.rs:48-49` +- `state_layout.cuh` C/CUDA mirror: `#define NUM_DIRECTIONS 3` and `DIRECTION_SHORT=0, + DIRECTION_LONG=1, DIRECTION_FLAT=2` (DIRECTION_HOLD removed) +- Q-head output dim 4 → 3, weight/bias re-init, layout fingerprint bump +- Action selection kernel argmax over 3 (`experience_action_select`) +- C51/IQN per-direction-bucket loops (~6 `for d<4` sites become `< NUM_DIRECTIONS`) +- **Action packing recompacts 12 → 9 slots**: `action_counts` buffer and `exp_idx = + dir*3 + mag` packing in `financials.rs:197` shrink from 12 → 9 (3 dirs × 3 mags). + Real shape change, not slot-3 dead-space. Consumers in `monitoring.rs:341` reindex. +- **Direction-bias signal at `experience_kernels.cu:5159`**: 4-tuple per-direction + bias weights `[Short=1.0, Hold=0.5, Long=1.0, Flat=0.0]` re-derives to 3-tuple + `[Short=1.0, Long=1.0, Flat=0.0]`. Hold's 0.5 was a softening gate; with Hold gone + the gate is no longer needed (Long/Short are full-weight, Flat is the soft branch). +- SP5 Pearl 8 per-direction trail-distance ISV range `TRAIL_DIST_PER_DIR_BASE=270..274`: + producer rewrites at new indices `{0=Short, 1=Long, 2=Flat}`; slot 273 dead-space + (per-direction-indexed ISV ranges keep length 4 to avoid cascade-renumber of all + downstream slots — consumers read by direction index ∈ 0..2 so dead slot is never + touched). Recompaction deferred to SP13 close-out. +- HEALTH_DIAG val_dir_dist 3-bucket emit: `metrics.rs:949`, `magnitude_distribution.rs:88,217`, + `distributional_q_tests.rs:790` +- KL divergence calc (if present) +- 30+ stale doc/comment lines across kernel headers — implementer audits in same commit + +### Change 2: Direction-skill bonus (bounded) + +``` +if action ∈ {Long, Short}: + correct = (sign(action) == sign(realized_30_bar_return)) + skill_cap = SKILL_BONUS_CAP_RATIO_INDEX × |alpha_capped| # default 0.3 × |alpha| + if correct: + skill_bonus = +min(DIR_SKILL_BONUS_ALPHA_INDEX, skill_cap) + else: + skill_bonus = -min(DIR_SKILL_BONUS_BETA_INDEX, skill_cap) +else: # action == Flat + skill_bonus = 0 +``` + +**Calibration check** (worked examples — α=β=1.0, ratio=0.3, SP12 cap [-10,+5]): + +| Trade outcome | alpha (capped) | skill_bonus | luck-disc | reward | sane? | +|---|---|---|---|---|---| +| correct-Long-winner (big win) | +5 | +min(1.0, 1.5) = +1.0 | n/a | +6.0 | ✓ | +| correct-Long-winner (small win) | +0.5 | +min(1.0, 0.15) = +0.15 | n/a | +0.65 | ✓ | +| correct-Long-loser (small loss) | −0.5 | +min(1.0, 0.15) = +0.15 | n/a | **−0.35** | ✓ still negative | +| correct-Long-loser (big loss) | −10 | +min(1.0, 3.0) = +1.0 | n/a | **−9.0** | ✓ still very negative | +| wrong-Long-winner LUCKY (big) | +5 | −min(1.0, 1.5) = −1.0 | × 0.3 | (5−1)×0.3 = **+1.2** | ✓ reduced from naive +6 | +| wrong-Long-winner LUCKY (small) | +0.5 | −min(1.0, 0.15) = −0.15 | × 0.3 | (0.5−0.15)×0.3 = +0.105 | ✓ | +| wrong-Long-loser | −10 | −min(1.0, 3.0) = −1.0 | n/a | **−11.0** | ✓ full penalty + skill penalty | +| Flat in any direction | 0 | 0 | n/a | 0 | ✓ | + +**Invariant verified**: every correct-direction loser has reward < 0; every +wrong-direction expected-loser has reward < 0. The model cannot game by entering +trades it expects to lose. ✓ + +### Change 3: Conditional correct-call asymmetry (luck-win discount) + +``` +if trade_pnl > 0 AND direction_was_wrong: + reward *= LUCK_WIN_DISCOUNT_INDEX /* default 0.3 */ +# else: reward unchanged +``` + +Penalises specifically lucky wins (positive outcome despite wrong direction prediction). +Forces the model to optimize for correct prediction, not just positive outcome. + +ISV-driven discount factor; default 0.3 starts strong, can be relaxed by ISV controller +if luck-win frequency drops below threshold. + +### Change 4: Aux head — regression → binary classification + +**Current**: `next_bar_return_mse` (regression). MSE 3.83 — worse than predicting zero. + +**New**: `next_30_bar_return_sign` cross-entropy. 30 bars matches SP12 min_hold_target. + +**Output**: 2-class softmax `[P(positive), P(negative)]`. + +**Wired into direction Q-head**: aux signal `logit_diff = softmax[1] − softmax[0]` ∈ +[−1, +1] is concatenated into direction Q-head input layer via ISV[`AUX_DIR_PREDICTION_INDEX`]. +This explicitly couples predictive signal to action selection. + +**No information leakage**: aux head trained with hindsight labels (sign of +price[t+30] − price[t], known at training time only). At inference, aux predicts using +state s_t alone — no future info accessible. Same as existing next_bar_mse pattern. + +### Change 5: ISV-driven aux_w controller (with stagnation detector) + +**Current** (SP11-era, broken): +```rust +let aux_w = (0.1 * learning_health * (1 - sharpe_tanh)).clamp(0.05, 0.3); +``` +Lowers aux when sharpe is high — exactly wrong for the failure mode (high sharpe + low WR). + +**New** (deficit-driven + stagnation): +``` +target = ISV[TARGET_DIR_ACC_INDEX] # default 0.55 +short = ISV[AUX_DIR_ACC_SHORT_EMA_INDEX] # fast EMA (α=0.3, ~3-epoch lookback) +long = ISV[AUX_DIR_ACC_LONG_EMA_INDEX] # slow EMA (α=0.05, ~20-epoch lookback) + +deficit = max(0, target - short) +improvement = max(0, short - long) # positive when short pulling ahead of long +stagnation = max(0, 1 - improvement / max(0.005, deficit)) + # 1 = fully stagnant (no improvement); 0 = improving fast + +aux_w = base * (1 + 5 * deficit) * (1 - 0.7 * stagnation) +aux_w = clamp(aux_w, base * 0.3, base * 3.0) # never zero, never beyond 3× base +``` + +Rationale: +- When `short < target` and improving (`short > long`): aux_w rises strongly to push harder +- When `short < target` and stagnant (`short ≈ long`): aux_w decays back toward base — "we tried, it didn't help, stop destabilising Q" +- When `short ≥ target`: deficit = 0, aux_w = base. Steady-state. +- Hard floor at 0.3× base prevents complete suppression (different from SP11 inversion bug). + +base = 0.5 in Phase 0b and beyond (SP11 base 0.05 was the suppression bug; lifting is +the fix, not a confounder, because the SP11 formula was structurally wrong). + +**ISV slots used**: 372 (target), 373 (short EMA), 374 (long EMA). Both EMAs reset per fold. + +## ISV slot allocation + +8 new slots in `[372..380)`: + +| Slot | Name | Purpose | Reset | +|---|---|---|---| +| 372 | `TARGET_DIR_ACC_INDEX` | Directional accuracy target (default 0.55) | static | +| 373 | `AUX_DIR_ACC_SHORT_EMA_INDEX` | Fast EMA of dir_acc (α=0.3) | per-fold | +| 374 | `AUX_DIR_ACC_LONG_EMA_INDEX` | Slow EMA of dir_acc (α=0.05), stagnation detector | per-fold | +| 375 | `AUX_DIR_PREDICTION_INDEX` | Aux head logit_diff per bar (concat into Q-head input) | per-bar | +| 376 | `DIR_SKILL_BONUS_ALPHA_INDEX` | Skill bonus magnitude (correct direction), default 1.0 | per-fold | +| 377 | `DIR_SKILL_BONUS_BETA_INDEX` | Skill bonus penalty (wrong direction), default 1.0 | per-fold | +| 378 | `LUCK_WIN_DISCOUNT_INDEX` | Lucky-win discount factor, default 0.3 | per-fold | +| 379 | `SKILL_BONUS_CAP_RATIO_INDEX` | Cap ratio: bonus ≤ ratio × \|alpha\|, default 0.3 | per-fold | + +Bump `SP5_SLOT_END = 380`, `ISV_TOTAL_DIM = 380`. + +**Slot 371 (was BENCHMARK_PNL_CUMULATIVE_INDEX in v1) — DROPPED.** Per-fold buy-and-hold +delta is computed at fold-end in HEALTH_DIAG emit (trivial: `price[fold_end] − +price[fold_start]`). No kernel, no slot, no reward dependency. + +## Implementation phases + +### Phase 0a — Hold elimination only (1 atomic commit) + +**Purpose**: clean test of user's hypothesis. Hold removed; aux_w controller, dir_acc +instrumentation, skill bonus, aux classification all UNCHANGED from current main. + +- ISV slot 372 (target, static), 373 (short EMA), 374 (long EMA), 375 (aux prediction — + but produced by the existing regression aux head, just `tanh(scalar_pred)` for now to + fit the [−1, +1] interface) +- `aux_dir_acc_reduce_kernel` — single-block tree-reduce of (correct_count, pos_pred, + pos_label, valid_count) → 3 output floats per epoch +- HEALTH_DIAG `aux_dir_acc accuracy=… pos_pred=… pos_label=…` per epoch +- Hold elimination atomic refactor (state_layout, Q-head, action select, C51/IQN bucket + state, replay encoding, val_dir_dist, KL, Rust enum) +- ~300 LOC + +**Phase 0a smoke** (10-epoch L40S): +- ✅ **Pass**: `val_dir_acc[ep9] > val_dir_acc[ep1]` AND `val_dir_acc[ep9] > 0.52` AND + `val_win_rate[ep9] > 0.51` → user's hypothesis confirmed; proceed Layer B. +- ⚠️ **Partial** (one trending up, the other flat): proceed P0b. +- ❌ **Both flat** (`dir_acc < 0.51` AND `win_rate < 0.50` after 10 epochs): proceed P0b + to add aux amplification. If P0b also fails, halt SP13. + +10 epochs (not 5) because aux head regression has been suppressed for entire SP1-SP12 +lineage — needs time even at unchanged aux_w to show whether signal is recoverable. +Cost: ~€0.60. + +### Phase 0b — Add corrected aux_w controller (only if 0a partial/red, atomic commit) + +- Replace SP11 inverted formula at `training_loop.rs:4103-4104` with deficit-driven + + stagnation formula from Change 5 +- Bump aux_w base 0.05 → 0.5 +- ~80 LOC + +**Phase 0b smoke** (10-epoch): +- Same gate criterion as P0a +- If still red: halt SP13. Update `project_hold_hides_directional_signal_hypothesis.md` + status to DISCONFIRMED. Pivot to data-augmentation work (separate spec, SP14). + +### Layer B — Aux head regression → classification (atomic, only on green Phase 0) + +- Aux head output dim 1 → 2 + softmax + cross-entropy loss +- Label generation: `sign(price[t+30] − price[t])` as 0/1 class +- `aux_dir_acc_reduce_kernel` adapted: read 2-class softmax instead of regression scalar +- Wire `logit_diff = softmax[1] − softmax[0]` → ISV[375] → direction Q-head input concat + (in_dim grows by 1) +- ~120 LOC + +### Layer C — Reward composition (atomic) + +- ISV slots 376, 377, 378, 379 + state-reset registry +- Direction-skill bonus device-inline function (`compute_direction_skill_bonus_bounded`) +- Luck-win discount device-inline function (`compute_luck_win_discount`) +- Replace SP12 reward composition at exit-event sites: keep SP12 asymmetric cap, ADD + skill bonus, ADD luck discount +- HEALTH_DIAG: add `per_fold_alpha_vs_buy_and_hold = total_pnl − (price[fold_end] − + price[fold_start])` — METRIC ONLY, not reward signal +- 12+ GPU oracle tests for new reward math (extends SP12 sp12_reward_math_tests scaffold) +- ~150 LOC + +### Layer D — 30-epoch validation + close-out (no commit until green) + +- 30-epoch L40S full validation +- 3-action SP12 control run for fair baseline (~€1, optional but recommended for clean + before/after comparison) +- If WR ≥ 55% AND profit_factor improvement ≥ 15% over 3-action SP12 control → + close-out commit: audit doc + 2 new pearls (`pearl_redefine_success_for_predictive_skill`, + `pearl_skill_bonus_must_be_alpha_bounded`) + meta-pearl + `pearl_reward_quadrant_audit_required` +- If WR plateaus 50-54%: tune horizon (next_60 instead of next_30) or skill bonus α/β, + re-validate + +## Validation targets + +**Phase 0a smoke (10-epoch)** — primary gate: +- val_dir_acc trajectory: monotone non-decreasing OR final > start +- val_dir_acc[final] > 0.52 +- val_win_rate[final] > 0.51 +- Hold action eliminated: val_dir_dist sums to 1.0 over {Long, Short, Flat} +- No NaN/inf +- mean(weights) controller still ≈ 1.0 +- aux next_bar_mse trajectory: not blowing up + +**Phase 0b smoke (10-epoch)** — secondary gate (only if 0a partial): +- aux_w controller responding: aux_w[ep5] > base when deficit > 0.05 +- Stagnation detector engaging: aux_w decays toward base when short EMA ≈ long EMA +- Same dir_acc + win_rate criteria as 0a + +**Layer B 5-epoch smoke**: +- Aux CrossEntropy loss decreasing +- val_dir_acc continues at or above Phase 0 level +- aux signal feeding Q-head (verify by perturbing aux output → direction action distribution shifts) + +**Full validation (30-epoch)**: +- val_dir_dist: roughly {Long: 0.4, Short: 0.4, Flat: 0.2} — model committing +- **win_rate ≥ 0.55** PRIMARY SUCCESS CRITERION +- profit_factor ≥ 1.5 (SP12 baseline 1.24 × 21% improvement floor for "skill-driven" claim) +- per_fold_alpha_vs_buy_and_hold > 0 sustained (the dropped reward signal, now metric) +- max_drawdown < 0.005 + +## Operating principles (NON-NEGOTIABLE) + +- **`feedback_no_partial_refactor`**: P0a, P0b, B, C each single atomic commit +- **`feedback_isv_for_adaptive_bounds`**: target_dir_acc, aux_w base/cap, bonus α/β, + luck-discount, skill-bonus-cap-ratio — ALL ISV-driven +- **`feedback_no_quickfixes`**: this IS the proper fix per pearls +- **`feedback_no_feature_flags`**: behavior change unconditional +- **End-to-end training**: NO pre-training of aux head +- **`pearl_audit_unboundedness_for_implicit_asymmetry`**: SP12 asymmetric cap [-10,+5] + preserved +- **`pearl_event_driven_reward_density_alignment`**: skill bonus + luck discount fire + only at trade-exit events, not per-bar +- **(NEW) Reward-quadrant audit**: any reward composition change must include the 4- + quadrant table (correct/wrong × winner/loser) showing all signs and inequalities. This + rule will become `pearl_reward_quadrant_audit_required` post-validation. + +## Estimated cost + +**Implementation**: ~10-12h (~700 LOC across P0a + P0b + B + C + D close-out) +**P0a smoke**: ~€0.60 (10-epoch L40S) +**P0b smoke** (if needed): ~€0.60 +**Layer B smoke**: ~€0.30 (5-epoch) +**30-epoch validation**: ~€1.00 +**3-action SP12 control** (recommended): ~€1.00 + +**Total best case** (P0a green, no P0b): ~€2.90 + 10-12h +**Total worst case** (P0a red, P0b green): ~€3.50 +**Total disconfirm case** (both P0a + P0b red): ~€1.20 — cheapest learn + +## Open questions + +1. **Aux prediction horizon**: 30 bars matches SP12 min_hold_target. If 30-epoch + validation shows partial WR climb (50-54%), Layer D iteration tries 60-bar horizon. +2. **Skill-bonus cap ratio**: 0.3 chosen so worst-case bonus contribution to reward at + |alpha|=10 is ±1.0 — comparable to one tick of P&L. ISV-driven, can be tightened to + 0.2 if Layer B/C reveals reward gaming. +3. **Stagnation detection threshold**: 0.7 decay coefficient is heuristic. Worth probing + in Layer C smoke (read aux_w trajectory). +4. **Whether to keep SP12 min-hold penalty**: SP12's `min_hold_factor` curriculum still + applies. Spec recommends KEEP — it's orthogonal to skill incentive and prevents + compulsive churn. Re-evaluate at Layer D close-out. + +## Decision needed + +Approve SP13 v2 spec → update implementation plan to match → dispatch P0a implementer +chain → 10-epoch smoke → branch on outcome.