From 8da8e2e584957543b4aaea7d0eaa933e31747c9a Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 9 May 2026 10:16:49 +0200 Subject: [PATCH] feat(sp18 v2 P2.T1-T5): structural Hold opportunity-cost device fn + 3-site migration (INTERIM STATE LIFTED) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Atomic introduction of `compute_sp18_hold_opportunity_cost` device function in `trade_physics.cuh` + migration of the 3 placeholder-commented per-bar Hold sites in `experience_kernels.cu` + CPU/GPU oracle tests + test-wrapper cubin manifest entry — single atomic commit per `feedback_no_partial_refactor`. Closes the SP18 INTERIM STATE introduced by Phase 1 (3c318953a). The reward function is now structurally complete: no placeholder sites, no orphaned consumers. Wire-up audit fingerprint stable; production-call count for the new device fn = 3 (segment_complete + per-bar positioned-Hold + per-bar flat-Hold), test-wrapper count = 1, definition count = 1. Sites: - trade_physics.cuh: new device fn after `compute_lump_sum_opp_cost`. Returns 0 for `dir_idx != DIR_HOLD`. Sign convention DD5(b) MIRRORED: positive next_log_return → negative Hold reward (opp cost paid for sitting out a winner); negative next_log_return → positive Hold reward (correctly avoided a loser). Asymmetric-capped via the existing `compute_asymmetric_capped_pnl` primitive against ISV[483]/[484]; output multiplied by `shaping_scale` (validation = 0 → returns 0). No new infrastructure beyond existing primitives. - experience_kernels.cu Site 1 (segment_complete ~3143): unconditional add into `base_reward` BEFORE the SP12 asymmetric cap; preserves SP12 cap interaction. - experience_kernels.cu Site 2 (per-bar positioned-Hold ~3641): unconditional add into `r_micro` (rc[3]) — same component the deleted SP13/SP16 subtraction wrote to, preserving SP11 reward-subsystem controller signal attribution. `vol_proxy` recomputed locally from `features[bar_idx * market_dim + 9]` (per-bar branches are outside the segment_complete `vol_proxy` scope). - experience_kernels.cu Site 3 (per-bar flat-Hold ~3760): unconditional add into `r_opp_cost` (rc[4]) — same component the deleted subtraction wrote to. `vol_proxy` recomputed locally with the same pattern. Each site reads ISV[HOLD_REWARD_POS_CAP_INDEX=483] / ISV[HOLD_REWARD_NEG_CAP_INDEX=484] with sentinel-or-out-of-bounds guard falling through to SENTINEL_HOLD_REWARD_POS/NEG_CAP=+5/-10 macros. Bit- identical to a fixed +5/-10 cap during cold-start and Pre-Phase-3 (Phase 3 lands the producer kernel `hold_reward_cap_update_kernel`). Tests: - T1 CPU oracle (sp18_hold_reward_oracle_tests.rs): 6 cases covering positive/negative log-return, dir-gate (early-exit for non-HOLD), positive saturation (clamps at +pos_cap), negative saturation (clamps at -neg_cap), shaping_scale=0 validation mode (returns 0). - T2 GPU oracle (sp18_hold_reward_oracle_tests.rs::gpu): same 6 cases via `sp18_hold_opp_test_kernel.cu` wrapper; bit-equality to CPU at ε=1e-5. - sp18_hold_opp_test_kernel.cu: single-thread single-block test wrapper mirroring `sp12_reward_math_test_kernel.cu` structure. Mapped-pinned f32 input/output per `feedback_no_htod_htoh_only_mapped_pinned`; __threadfence_system() for PCIe visibility before host read. - build.rs: cubin manifest entry for sp18_hold_opp_test_kernel.cu. Validation: - SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean. - cargo test -p ml --lib sp18: 4 SP18 lib tests pass (slot layouts, fold-reset entries-present). - cargo test -p ml --lib state_reset_registry: 5 tests pass (including every_fold_and_soft_reset_entry_has_dispatch_arm). - cargo test -p ml --test sp18_hold_reward_oracle_tests --features cuda -- --include-ignored: 7 tests pass (2 CPU + 5 GPU on local RTX 3050 Ti), including the new compute_sp18_hold_opportunity_cost_cpu_oracle and compute_sp18_hold_opportunity_cost_gpu_oracle. - bash scripts/audit_sp18_consumers.sh --check: exit 0 (post-Phase-2 fingerprint snapshotted in docs/sp18-wireup-audit.md). - Wire-up audit: `grep -rn "compute_sp18_hold_opportunity_cost" crates/ml/` shows 1 device-fn definition + 3 production call sites + 1 test-wrapper kernel + CPU/GPU oracle tests. No orphans. Audit fingerprint diff (Phase 1 → Phase 2): - `Slot 380 (HOLD_COST_INDEX) consumers`: experience_kernels.cu hit count 1 → 0 (placeholder-comment reference to HOLD_COST_INDEX removed; new Phase 2 calls reference HOLD_REWARD_POS/NEG_CAP_INDEX=483/484 instead); doc references in dqn-wire-up-audit.md grew by 2 lines (Phase 2 close-out narrative). Total section count 42 → 43. - All other sections unchanged — Phase 2 introduces forward-tracking references on the SP18-bound side (slots 483/484, new device fn), not the deleted-chain side. Discipline: - feedback_no_partial_refactor: device fn + 3 call sites + tests + cubin manifest in same commit; INTERIM STATE LIFTED in same commit. - feedback_no_atomicadd: device fn is per-thread sample-local register arithmetic; test wrapper is single-thread single-block, no shared mem. - feedback_no_htod_htoh_only_mapped_pinned: GPU oracle test uses MappedF32Buffer. - feedback_isv_for_adaptive_bounds: pos/neg caps read from ISV[483]/[484] with sentinel fallback; Phase 3 producer `hold_reward_cap_update_kernel` drives them adaptively. - feedback_wire_everything_up: device fn + 3 call sites + test wrapper + CPU/GPU oracle tests + cubin manifest all in same commit. - pearl_one_unbounded_signal_per_reward: only `vol_normalized_return` is unbounded; everything else in [neg_cap, pos_cap] × shaping_scale. - pearl_first_observation_bootstrap: sentinels (5.0/-10.0) match SP14 P0-A REWARD_POS/NEG_CAP_ADAPTIVE for bit-identical cold-start. Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md Phase 2 Tasks 2.1-2.5. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 14 + .../src/cuda_pipeline/experience_kernels.cu | 205 ++++++++++--- .../sp18_hold_opp_test_kernel.cu | 62 ++++ crates/ml/src/cuda_pipeline/trade_physics.cuh | 82 +++++ .../ml/tests/sp18_hold_reward_oracle_tests.rs | 280 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 103 +++++++ docs/sp18-wireup-audit.md | 81 ++++- 7 files changed, 783 insertions(+), 44 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 1ef7bbb82..4696b5b56 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -896,6 +896,20 @@ fn main() { // Pure observability: pre-fix baseline for the B-DD9 ratio // gate. Plan: SP18 Phase 0 Task 0.2. "td_error_mag_ema_kernel.cu", + // SP18 v2 Phase 2 (2026-05-09): D-leg structural Hold + // opportunity-cost device-fn GPU oracle test wrapper. Single + // thread × single block; wraps `compute_sp18_hold_opportunity_cost` + // from `trade_physics.cuh`. Mapped-pinned f32 input/output (no + // htod/dtoh per `feedback_no_htod_htoh_only_mapped_pinned`); host + // reads via `read_volatile` after stream sync (the + // `__threadfence_system()` at thread exit guarantees PCIe + // visibility before the sync returns). Same structural pattern + // as `sp12_reward_math_test_kernel.cu`. Test-only; no production + // callers. Drives `compute_sp18_hold_opportunity_cost`'s closed- + // form math contract from the Rust GPU oracle in + // `tests/sp18_hold_reward_oracle_tests.rs` against the CPU + // oracle in the same file (ε=1e-5). Plan: SP18 Phase 2 Task 2.2. + "sp18_hold_opp_test_kernel.cu", // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. // Single-block 256-thread kernel computing RMS = sqrt(mean(x²)) // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 9c273b4da..bafacaf50 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -3140,19 +3140,54 @@ extern "C" __global__ void experience_env_step( * `exiting_trade=1` upstream (line ~2371), so this branch fires * for both — the segregation matters only for component attribution. */ float base_reward = 2.0f * vol_normalized_return; - /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction at the - * voluntary segment_complete branch. Goes BEFORE the SP12 asymmetric - * cap so the cost participates in the bounded clamp — keeps total - * reward in the same [-10, +5] range and avoids a long Hold run - * silently saturating the lower cap. Hold action rarely coincides - * with segment_complete (Hold keeps current position, doesn't close - * trades) but trail-fire can force-exit while the policy picks Hold; - * this branch covers that edge. Per-bar (non-segment_complete) - * Hold-cost subtraction is in the per-bar branches below. + /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost + * counterfactual replaces the SP13 P0a / SP16 P2 / SP16 T3 reactive + * per-bar Hold-cost-scale chain (deleted Phase 1 c66e15f6d / + * 3c318953a). Goes BEFORE the SP12 asymmetric cap so the + * counterfactual participates in the bounded clamp — keeps total + * reward in the [-10, +5] range and avoids a long Hold run silently + * saturating the lower cap. Hold action rarely coincides with + * segment_complete (Hold keeps current position, doesn't close + * trades) but trail-fire can force-exit while the policy picks + * Hold; this branch covers that edge. Per-bar (non-segment_complete) + * Hold sites are in the two branches below. * - * SP18 D-leg replacement: per-bar Hold cost subtraction removed. - * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. - * (No reward write here in interim; Phase 2 lands the new device fn.) */ + * The device function returns 0 for non-Hold so unconditional + * addition is safe. Caps come from ISV[HOLD_REWARD_NEG_CAP_INDEX=484] + * / ISV[HOLD_REWARD_POS_CAP_INDEX=483] with sentinel/window-guard + * fallback to the macro defaults — bit-identical to a fixed +5/-10 + * cap during cold-start (Phase 3 lands the adaptive producer). + * `shaping_scale` matches the SP12 contract (validation = 0 → no + * Hold opp cost in pure backtest). Per + * `pearl_one_unbounded_signal_per_reward.md` the only unbounded + * multiplicand is `next_log_return / sqrt(vol_proxy)`; everything + * else is in [floor, ceil]. */ + { + float sp18_pos_cap = SENTINEL_HOLD_REWARD_POS_CAP; /* +5.0 */ + float sp18_neg_cap = SENTINEL_HOLD_REWARD_NEG_CAP; /* -10.0 */ + if (isv_signals_ptr != NULL) { + const float isv_pos = + isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX]; + const float isv_neg = + isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX]; + if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f + && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND + && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) { + sp18_pos_cap = isv_pos; + sp18_neg_cap = isv_neg; + } + } + const float sp18_hold_opp = compute_sp18_hold_opportunity_cost( + dir_idx, + /*dir_hold=*/DIR_HOLD, + /*next_log_return=*/tgt[1], + /*vol_proxy=*/vol_proxy, + /*isv_neg_cap=*/sp18_neg_cap, + /*isv_pos_cap=*/sp18_pos_cap, + /*shaping_scale=*/shaping_scale + ); + base_reward += sp18_hold_opp; + } /* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion). * SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63 * PopArt EMA inflation. That fix preserved stability but erased the @@ -3603,21 +3638,73 @@ extern "C" __global__ void experience_env_step( * (which is a separate concern; spec carve-out: "ONLY micro_reward * and opp_cost are the per-bar shaping anti-patterns"). */ r_micro = 0.0f; - /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction. - * Hold action keeps the current position with zero PnL signal — - * pre-SP13 this was a free no-op the policy could pick indefinitely - * (CQL bias anchors toward minimum-variance Hold ~45% of bars across - * SP1-SP12). Hold-pricing makes Hold uneconomic when overused: the - * controller in `training_loop.rs` raises ISV[HOLD_COST_INDEX] when - * observed Hold-rate > target (default 0.20). The cost is bounded - * structurally (HOLD_COST_BASE × CEIL_RATIO = 0.005 max per bar), - * well below the SP12 asymmetric cap [-10, +5] — no per-bar clamp - * required. Spec §"Change 1: Price Hold (replaces v2's Eliminate - * Hold)". + /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost + * counterfactual on positioned-non-event bars. Replaces the SP13 + * P0a / SP16 P2 / SP16 T3 reactive per-bar Hold-cost-scale chain + * (deleted Phase 1 c66e15f6d / 3c318953a). Routes through the + * micro_reward component (rc[3]) — the same component the deleted + * subtraction wrote to via `r_micro -= cost` pre-SP18, preserving + * SP11 reward-subsystem controller signal attribution. * - * SP18 D-leg replacement: per-bar Hold cost subtraction removed. - * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. - * (No reward write here in interim; Phase 2 lands the new device fn.) */ + * The device function returns 0 for `dir_idx != DIR_HOLD`, so + * unconditional addition is safe. `vol_proxy` and `next_log_return` + * are recomputed locally from `features[bar_idx * market_dim + 9]` + * + `tgt[1]` — `vol_proxy` is scoped to the segment_complete branch + * above, so we mirror the segment_complete formula exactly here. + * The `features != NULL && bar_idx < total_bars && market_dim > 9` + * guard matches the B.2/D.4c bonus blocks below; on cold-start / + * test-fixture inputs without market features we skip cleanly + * (sp18_hold_opp stays at 0.0). + * + * Caps come from ISV[HOLD_REWARD_NEG_CAP_INDEX=484] / + * ISV[HOLD_REWARD_POS_CAP_INDEX=483] with sentinel/window-guard + * fallback to the macro defaults — bit-identical to a fixed +5/-10 + * cap during cold-start (Phase 3 lands the adaptive producer). + * `shaping_scale` matches the SP12 contract (validation = 0 → no + * Hold opp cost in pure backtest). Per + * `pearl_one_unbounded_signal_per_reward.md` the only unbounded + * multiplicand is `next_log_return / sqrt(vol_proxy)`; everything + * else is in [floor, ceil]. + * + * The structurally-bounded output (capped in [neg_cap, pos_cap]) + * sits comfortably inside the SP12 asymmetric reward cap [-10, +5] + * the segment_complete branch enforces — but this branch fires + * BETWEEN trade close events (positioned-non-event), so no + * downstream cap is applied here; the device function's + * `compute_asymmetric_capped_pnl` is the structural bound. */ + if (features != NULL && bar_idx < total_bars && market_dim > 9) { + const float atr_norm_pb1 = + features[(long long)bar_idx * market_dim + 9]; + const float log_atr_pb1 = atr_norm_pb1 * 16.0f - 7.0f; + const float atr_pct_pb1 = + expf(log_atr_pb1) / fmaxf(raw_close, 1.0f); + const float vol_proxy_pb1 = fmaxf(atr_pct_pb1, 0.0001f); + + float sp18_pos_cap_pb1 = SENTINEL_HOLD_REWARD_POS_CAP; /* +5.0 */ + float sp18_neg_cap_pb1 = SENTINEL_HOLD_REWARD_NEG_CAP; /* -10.0 */ + if (isv_signals_ptr != NULL) { + const float isv_pos = + isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX]; + const float isv_neg = + isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX]; + if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f + && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND + && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) { + sp18_pos_cap_pb1 = isv_pos; + sp18_neg_cap_pb1 = isv_neg; + } + } + const float sp18_hold_opp_pb1 = compute_sp18_hold_opportunity_cost( + dir_idx, + /*dir_hold=*/DIR_HOLD, + /*next_log_return=*/tgt[1], + /*vol_proxy=*/vol_proxy_pb1, + /*isv_neg_cap=*/sp18_neg_cap_pb1, + /*isv_pos_cap=*/sp18_pos_cap_pb1, + /*shaping_scale=*/shaping_scale + ); + r_micro += sp18_hold_opp_pb1; + } reward = r_micro; micro_reward_per_sample[out_off] = r_micro; reward_components_per_sample[out_off * 6 + 3] = r_micro; @@ -3672,17 +3759,63 @@ extern "C" __global__ void experience_env_step( * preserved: only the Flat opp_cost EMITTER is zeroed; the * conviction Welford observer is independent. */ r_opp_cost = 0.0f; - /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction on - * the flat-non-event branch. Mirrors the positioned-branch - * subtraction above so a policy picking Hold while flat (i.e. - * "stay flat") pays the same per-bar cost as a policy picking - * Hold while positioned (i.e. "stay long/short"). The cost is - * structurally bounded by HOLD_COST_BASE × CEIL_RATIO = 0.005 - * per bar, well below the SP12 asymmetric cap [-10, +5]. + /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost + * counterfactual on flat-non-event bars. Mirrors the positioned- + * branch wiring above so a policy picking Hold while flat ("stay + * flat") sees the same structural counterfactual as a policy + * picking Hold while positioned ("stay long/short"). Routes + * through the opp_cost component (rc[4]) — the same component + * the deleted SP13/SP16 chain wrote to via `r_opp_cost -= cost` + * pre-SP18, preserving SP11 reward-subsystem controller signal + * attribution. * - * SP18 D-leg replacement: per-bar Hold cost subtraction removed. - * Replaced with compute_sp18_hold_opportunity_cost in Phase 2. - * (No reward write here in interim; Phase 2 lands the new device fn.) */ + * The device function returns 0 for `dir_idx != DIR_HOLD`, so + * unconditional addition is safe. `vol_proxy` and `next_log_return` + * recomputed locally — same pattern as the positioned-branch and + * the B.2/D.4c bonus blocks. Caps come from ISV[483]/[484] with + * sentinel/window-guard fallback to the macro defaults; bit- + * identical to a fixed +5/-10 cap during cold-start (Phase 3 + * lands the adaptive producer). `shaping_scale=0` (validation + * mode) → device function returns 0 → no Hold opp cost in pure + * backtest, matching the SP12 lump-sum contract. + * + * Output bounded by `compute_asymmetric_capped_pnl` to + * [neg_cap, pos_cap] before scaling by `shaping_scale`; sits + * inside the SP12 asymmetric reward cap range. Per + * `pearl_one_unbounded_signal_per_reward.md`. */ + if (features != NULL && bar_idx < total_bars && market_dim > 9) { + const float atr_norm_pb2 = + features[(long long)bar_idx * market_dim + 9]; + const float log_atr_pb2 = atr_norm_pb2 * 16.0f - 7.0f; + const float atr_pct_pb2 = + expf(log_atr_pb2) / fmaxf(raw_close, 1.0f); + const float vol_proxy_pb2 = fmaxf(atr_pct_pb2, 0.0001f); + + float sp18_pos_cap_pb2 = SENTINEL_HOLD_REWARD_POS_CAP; /* +5.0 */ + float sp18_neg_cap_pb2 = SENTINEL_HOLD_REWARD_NEG_CAP; /* -10.0 */ + if (isv_signals_ptr != NULL) { + const float isv_pos = + isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX]; + const float isv_neg = + isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX]; + if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f + && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND + && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) { + sp18_pos_cap_pb2 = isv_pos; + sp18_neg_cap_pb2 = isv_neg; + } + } + const float sp18_hold_opp_pb2 = compute_sp18_hold_opportunity_cost( + dir_idx, + /*dir_hold=*/DIR_HOLD, + /*next_log_return=*/tgt[1], + /*vol_proxy=*/vol_proxy_pb2, + /*isv_neg_cap=*/sp18_neg_cap_pb2, + /*isv_pos_cap=*/sp18_pos_cap_pb2, + /*shaping_scale=*/shaping_scale + ); + r_opp_cost += sp18_hold_opp_pb2; + } reward = r_opp_cost; reward_components_per_sample[out_off * 6 + 4] = r_opp_cost; } diff --git a/crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu b/crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu new file mode 100644 index 000000000..a77dc9ec6 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu @@ -0,0 +1,62 @@ +/* SP18 v2 D-leg Phase 2 (2026-05-09): GPU oracle test kernel for + * `compute_sp18_hold_opportunity_cost`. + * + * Standalone test wrapper that exercises the structural Hold opportunity- + * cost device function added to `trade_physics.cuh` for SP18 v2 Phase 2. + * Mirrors the SP12 `sp12_reward_math_test_kernel.cu` pattern: single-thread + * single-block launch, mapped-pinned f32 input/output buffers, host reads + * via `read_volatile` after a stream sync (the `__threadfence_system()` at + * exit guarantees the device write is PCIe-visible). + * + * Per `feedback_no_cpu_test_fallbacks.md`: GPU oracle only (no CPU + * reference impl in production). The Rust CPU oracle test in + * `tests/sp18_hold_reward_oracle_tests.rs` locks the closed-form math + * contract; this wrapper drives the same inputs through the production + * device function and the GPU oracle test asserts bit-equality to the + * closed-form expected values within an f32 rounding budget (ε=1e-5). + * + * Per `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU + * buffer is `MappedF32Buffer` (cuMemHostAlloc DEVICEMAP|PORTABLE); zero + * `htod_copy`, zero `dtoh_sync_copy` in the test. Per + * `feedback_no_atomicadd.md`: single-thread launch, no shared memory, no + * reduction — pure register arithmetic. + * + * Test-only: no production callers. Cubin built by `crates/ml/build.rs` + * (kernel name `sp18_hold_opp_test_kernel`). + */ + +#include +#include "trade_physics.cuh" /* compute_sp18_hold_opportunity_cost */ + +extern "C" __global__ void sp18_hold_opp_test_kernel( + int dir_idx, + int dir_hold, + float next_log_return, + float vol_proxy, + float isv_neg_cap, + float isv_pos_cap, + float shaping_scale, + float* __restrict__ out_hold_opp_cost +) { + /* Single-thread launch — the device function has no parallelism to + * exploit (pure register arithmetic; one (input, output) pair per + * test invocation). */ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + out_hold_opp_cost[0] = compute_sp18_hold_opportunity_cost( + dir_idx, + dir_hold, + next_log_return, + vol_proxy, + isv_neg_cap, + isv_pos_cap, + shaping_scale + ); + + /* Make the write PCIe-visible to the host pointer (mapped-pinned + * buffer). The host reads via `read_volatile` after a stream sync; the + * fence guarantees the write has actually drained from the GPU write + * buffers before the sync returns. Same pattern as + * `sp12_reward_math_test_kernel.cu`. */ + __threadfence_system(); +} diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index ff3cc1340..d90983d95 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -599,6 +599,88 @@ __device__ __forceinline__ float compute_lump_sum_opp_cost( return -(shaping_scale * holding_cost_rate * fabsf(position) * hold_time); } +/* SP18 D-leg (2026-05-09): structural Hold opportunity-cost reward. + * + * Replaces the SP13 P0a + SP16 P2 + SP16 T3 reactive per-bar Hold-cost-scale + * subtraction chain (deleted in SP18 Phase 1) with an event-driven structural + * counterfactual: when the policy picks Hold, what would Long have earned on + * the next bar? Sign convention DD5(b) MIRRORED: + * + * - positive next_log_return (Long would have won) → NEGATIVE Hold reward + * (opportunity cost paid for sitting out). + * - negative next_log_return (Long would have lost) → POSITIVE Hold reward + * (correctly avoided a losing trade). + * + * The base counterfactual is `-base = -(2.0 × vol_normalized_log_return)` + * — the same `2.0 × vol_normalized_return` scalar the segment_complete branch + * uses for `base_reward`, sign-flipped so a positive return becomes a + * negative Hold reward (opp cost). The factor-2 prefactor mirrors the + * segment_complete `base_reward = 2.0f * vol_normalized_return` to keep the + * Hold-arm and Long/Short-arm reward distributions on the same scale per + * spec §"Hold-anchored reward decomposition". + * + * `vol_normalized_return = next_log_return / sqrtf(vol_proxy + EPS)` — single- + * bar variant of the segment_complete formula (where `vol_norm = + * vol_proxy * sqrtf(hold_time)` reduces to `vol_proxy` for hold_time=1). The + * EPS=1e-12 floor prevents division by zero in degenerate `vol_proxy=0` data + * regions (the per-bar caller guards `vol_proxy >= 0.0001f` upstream, so EPS + * is a numerical safety floor, not a behavioural knob). + * + * Output is asymmetric-capped via `compute_asymmetric_capped_pnl` against + * the ISV-driven SP18 caps (slots 483/484 — produced adaptively in Phase 3 + * by `hold_reward_cap_update_kernel`; sentinel +5/-10 mirrors SP14 P0-A). + * `shaping_scale` gates behavioral-mode vs validation-mode (=0 → pure + * backtest reward shape, no Hold opp cost) matching the SP12 lump-sum + * `compute_lump_sum_opp_cost` contract. + * + * Returns 0.0 exactly for `dir_idx != DIR_HOLD` so the per-bar callers can + * call unconditionally without a gate at the call site (the function itself + * is the gate). DIR_HOLD value (=1) lives in `state_layout.cuh`; the test + * wrapper passes the same constant. + * + * Per `pearl_one_unbounded_signal_per_reward`: only `vol_normalized_return` + * is unbounded; everything downstream sits inside [neg_cap, pos_cap] × + * shaping_scale. Per `pearl_audit_unboundedness_for_implicit_asymmetry`: + * the bilateral cap with `neg_cap = -2 × pos_cap` (Phase-3 producer + * convention) preserves loss-aversion asymmetry — Hold opp costs land in + * [-10, 0] under sentinels, Hold-correct-avoidance gains land in [0, +5]. + * + * Math factored into a device function so the formula is testable in + * isolation (oracle tests in `tests/sp18_hold_reward_oracle_tests.rs`, + * GPU wrapper in `sp18_hold_opp_test_kernel.cu`). Per + * `feedback_no_cpu_test_fallbacks.md`, the GPU oracle test exercises this + * exact device function — no CPU reference impl in production. The CPU + * oracle in the Rust test file is the math contract spec, mirrored + * here exactly. */ +__device__ __forceinline__ float compute_sp18_hold_opportunity_cost( + int dir_idx, /* current action direction (DIR_SHORT/HOLD/LONG/FLAT) */ + int dir_hold, /* DIR_HOLD constant from state_layout.cuh — passed + explicitly so the test wrapper avoids the include + chain dependency */ + float next_log_return, /* tgt[1] — preprocessed next-bar log-return */ + float vol_proxy, /* per-bar vol scalar (production: same atr_pct + chain the segment_complete branch uses, + floored at 0.0001f upstream) */ + float isv_neg_cap, /* ISV[HOLD_REWARD_NEG_CAP_INDEX=484]; sentinel -10.0 */ + float isv_pos_cap, /* ISV[HOLD_REWARD_POS_CAP_INDEX=483]; sentinel +5.0 */ + float shaping_scale /* SP12 behavioral-mode gate (validation = 0) */ +) { + if (dir_idx != dir_hold) return 0.0f; + + /* EPS floor matches the per-bar caller's `fmaxf(atr_pct, 0.0001f)` guard + * in spirit (numerical safety, not behavioural). 1e-12 is well below + * any plausible vol_proxy so it is a no-op at production inputs. */ + const float vol_normalized_return = + next_log_return / sqrtf(vol_proxy + 1e-12f); + const float base = 2.0f * vol_normalized_return; + /* DD5(b) MIRRORED sign: positive next_log_return → negative Hold reward + * (opp cost). raw_opp = -base. */ + const float raw_opp = -base; + const float capped = + compute_asymmetric_capped_pnl(raw_opp, isv_neg_cap, isv_pos_cap); + return capped * shaping_scale; +} + /* ── Drawdown penalty: DELETED (Class A audit-fix Batch 4-A, 2026-05-08) ── * * The legacy `compute_drawdown_penalty` device function was deleted diff --git a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs index 12b718e66..30adc362d 100644 --- a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs +++ b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs @@ -114,6 +114,124 @@ fn reward_decomp_per_action_cpu_oracle() { assert!((out[19] - 1.00).abs() < 1e-6, "flat fire: got {}", out[19]); } +// ── SP18 v2 Phase 2 Task 2.1 (T1): CPU oracle for `compute_sp18_hold_opportunity_cost` ── + +/// CPU-side spec mirror of the `compute_sp18_hold_opportunity_cost` device +/// function in `crates/ml/src/cuda_pipeline/trade_physics.cuh`. +/// +/// This is the math contract — the device function must be bit-equivalent +/// to within an f32 rounding budget (ε=1e-5; see Task 2.3 GPU oracle). +/// Per `feedback_no_cpu_test_fallbacks.md` this is **not** a production +/// reference implementation; production runs only the GPU device fn. The +/// CPU oracle exists solely to lock the closed-form math contract during +/// TDD. +/// +/// Inputs / outputs match the device fn signature exactly: +/// - `dir_idx != DIR_HOLD` ⇒ returns 0.0 exactly (early exit gate). +/// - DD5(b) MIRRORED sign: positive `next_log_return` ⇒ negative result +/// (opp cost paid for sitting out a winning Long). +/// - `compute_asymmetric_capped_pnl(raw, neg_cap, pos_cap) +/// = fmaxf(neg_cap, fminf(raw, pos_cap))`. +/// - Multiplied by `shaping_scale` — validation mode (=0) ⇒ returns 0. +/// +/// EPS=1e-12 mirrors the device fn's vol_proxy guard (numerical safety, +/// not behavioural — production caller's `fmaxf(atr_pct, 0.0001f)` makes +/// this floor a no-op in practice). +fn compute_sp18_hold_opp_cpu( + dir_idx: i32, + dir_hold: i32, + next_log_return: f32, + vol_proxy: f32, + isv_neg_cap: f32, + isv_pos_cap: f32, + shaping_scale: f32, +) -> f32 { + if dir_idx != dir_hold { + return 0.0; + } + let vol_normalized_return = next_log_return / (vol_proxy + 1e-12_f32).sqrt(); + let base = 2.0_f32 * vol_normalized_return; + let raw_opp = -base; + let capped = isv_neg_cap.max(raw_opp.min(isv_pos_cap)); + capped * shaping_scale +} + +const DIR_HOLD_CPU: i32 = 1; // matches `state_layout.cuh:124` + +/// T1 (spec §"Behavioral Test Predictions"): CPU oracle for +/// `compute_sp18_hold_opportunity_cost`. +/// +/// Locks the closed-form math contract before the device fn lands. With +/// `next_log_return = +0.005`, `vol_proxy = 0.0001`, `pos_cap = +5.0`, +/// `neg_cap = -10.0`, `shaping_scale = 1.0`: +/// +/// - `vol_normalized_return = 0.005 / sqrt(0.0001 + 1e-12) ≈ 0.5` +/// - `base = 2.0 × 0.5 = 1.0` +/// - `raw_opp = -1.0` +/// - `capped = max(-10, min(-1, +5)) = -1.0` +/// - `result = -1.0 × 1.0 = -1.0` +/// +/// Verifies: returns negative scalar bounded in [-10, 0] for positive +/// `next_log_return` (opp cost paid for sitting out a winning Long); a +/// flipped-sign positive scalar bounded in [0, +5] for negative +/// `next_log_return` (correctly avoided a losing Long); and 0.0 exactly +/// for `dir_idx != DIR_HOLD`. +#[test] +fn compute_sp18_hold_opportunity_cost_cpu_oracle() { + // Case A: positive next_log_return, dir=Hold → negative result in [-10, 0]. + let r_pos = compute_sp18_hold_opp_cpu( + DIR_HOLD_CPU, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 1.0, + ); + assert!((r_pos - (-1.0)).abs() < 1e-5, + "pos log-return: expected -1.0, got {r_pos}"); + assert!(r_pos >= -10.0 && r_pos <= 0.0, + "pos log-return result must be in [-10, 0], got {r_pos}"); + + // Case B: negative next_log_return → positive result in [0, +5]. + let r_neg = compute_sp18_hold_opp_cpu( + DIR_HOLD_CPU, DIR_HOLD_CPU, -0.005, 0.0001, -10.0, 5.0, 1.0, + ); + assert!((r_neg - 1.0).abs() < 1e-5, + "neg log-return: expected +1.0, got {r_neg}"); + assert!(r_neg >= 0.0 && r_neg <= 5.0, + "neg log-return result must be in [0, +5], got {r_neg}"); + + // Case C: dir != HOLD → 0.0 exactly (early-exit gate). + for dir in [0_i32, 2, 3] { // SHORT, LONG, FLAT + let r = compute_sp18_hold_opp_cpu( + dir, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 1.0, + ); + assert_eq!(r, 0.0, + "dir={dir} (non-HOLD): expected 0.0 exactly, got {r}"); + } + + // Case D: positive saturation — large negative log-return floods +pos_cap. + // vol_norm = -0.5 / sqrt(0.0001) = -50; base = -100; raw_opp = +100; + // capped = min(100, +5) = +5; final = +5.0. + let r_sat_pos = compute_sp18_hold_opp_cpu( + DIR_HOLD_CPU, DIR_HOLD_CPU, -0.5, 0.0001, -10.0, 5.0, 1.0, + ); + assert!((r_sat_pos - 5.0).abs() < 1e-5, + "+pos saturation: expected +5.0, got {r_sat_pos}"); + + // Case E: negative saturation — large positive log-return floods -neg_cap. + // vol_norm = 0.05 / sqrt(0.0001) = 5; base = 10; raw_opp = -10; + // capped = max(-10, -10) = -10; final = -10.0. + let r_sat_neg = compute_sp18_hold_opp_cpu( + DIR_HOLD_CPU, DIR_HOLD_CPU, 0.05, 0.0001, -10.0, 5.0, 1.0, + ); + assert!((r_sat_neg - (-10.0)).abs() < 1e-5, + "-neg saturation: expected -10.0, got {r_sat_neg}"); + + // Case F: shaping_scale = 0.0 (validation mode) → result = 0.0 + // regardless of inputs, mirroring the SP12 lump-sum opp-cost contract. + let r_validation = compute_sp18_hold_opp_cpu( + DIR_HOLD_CPU, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 0.0, + ); + assert_eq!(r_validation, 0.0, + "shaping_scale=0 (validation): expected 0.0 exactly, got {r_validation}"); +} + // ── Layer B (GPU oracle): kernel matches CPU oracle bit-for-bit ────────── #[cfg(feature = "cuda")] @@ -489,4 +607,166 @@ mod gpu { assert!(out[1 * NUM_COLS + 0] > 0.0, "long micro should be > 0"); assert!(out[1 * NUM_COLS + 4] > 0.0, "long fire should be > 0"); } + + // ── SP18 v2 Phase 2 Task 2.3 (T2): GPU oracle for `compute_sp18_hold_opportunity_cost` ── + + /// Cubin handle for the SP18 D-leg structural Hold opportunity-cost + /// device-function GPU oracle test wrapper. + const SP18_HOLD_OPP_TEST_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/sp18_hold_opp_test_kernel.cubin")); + + fn load_sp18_hold_opp_test(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP18_HOLD_OPP_TEST_CUBIN.to_vec()) + .expect("load sp18_hold_opp_test_kernel cubin"); + module + .load_function("sp18_hold_opp_test_kernel") + .expect("load sp18_hold_opp_test_kernel function") + } + + const DIR_HOLD_GPU: i32 = 1; // matches `state_layout.cuh:124` + + /// CPU mirror for the GPU oracle assertions — duplicates the + /// production CPU oracle helper at the top of this file (which lives + /// outside the `gpu` mod due to `#[cfg(feature = "cuda")]` gating). + /// Keeping the math inline preserves the per-test readability without + /// re-exporting helpers across module boundaries. + fn cpu_oracle( + dir_idx: i32, + next_log_return: f32, + vol_proxy: f32, + isv_neg_cap: f32, + isv_pos_cap: f32, + shaping_scale: f32, + ) -> f32 { + if dir_idx != DIR_HOLD_GPU { + return 0.0; + } + let vnr = next_log_return / (vol_proxy + 1e-12_f32).sqrt(); + let base = 2.0_f32 * vnr; + let raw = -base; + let capped = isv_neg_cap.max(raw.min(isv_pos_cap)); + capped * shaping_scale + } + + /// Single-shot GPU oracle launch: writes one f32 to a 1-element + /// mapped-pinned output buffer; reads back via `read_volatile` after + /// stream sync. + fn launch_sp18_hold_opp( + stream: &Arc, + kernel: &CudaFunction, + dir_idx: i32, + next_log_return: f32, + vol_proxy: f32, + isv_neg_cap: f32, + isv_pos_cap: f32, + shaping_scale: f32, + ) -> f32 { + let out_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc out buf"); + // Pre-fill with NaN so any unwritten output trips a finite-check + // failure in the caller (locks the kernel actually wrote). + out_buf.write_from_slice(&[f32::NAN]); + + let dir_hold = DIR_HOLD_GPU; + unsafe { + stream + .launch_builder(kernel) + .arg(&dir_idx) + .arg(&dir_hold) + .arg(&next_log_return) + .arg(&vol_proxy) + .arg(&isv_neg_cap) + .arg(&isv_pos_cap) + .arg(&shaping_scale) + .arg(&out_buf.dev_ptr) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch sp18_hold_opp_test_kernel"); + } + stream.synchronize().expect("sync after sp18_hold_opp_test_kernel"); + + out_buf.read_all()[0] + } + + /// T2 (spec §"Behavioral Test Predictions"): GPU oracle parity with + /// CPU oracle. Drives every Case A-F from the CPU oracle test through + /// the production device function via the test wrapper kernel and + /// asserts bit-equality to ε=1e-5 (f32 rounding budget — `expf`, + /// `sqrtf`, and the asymmetric clamp are deterministic to within a + /// few ulps on Ampere/Ada). + #[test] + #[ignore = "requires GPU"] + fn compute_sp18_hold_opportunity_cost_gpu_oracle() { + let stream = make_test_stream(); + let kernel = load_sp18_hold_opp_test(&stream); + + let eps = 1e-5_f32; + + // Case A: positive next_log_return, dir=Hold → ~-1.0 (in [-10, 0]). + let inputs_a = (DIR_HOLD_GPU, 0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32); + let cpu_a = cpu_oracle(inputs_a.0, inputs_a.1, inputs_a.2, inputs_a.3, inputs_a.4, inputs_a.5); + let gpu_a = launch_sp18_hold_opp( + &stream, &kernel, inputs_a.0, inputs_a.1, inputs_a.2, inputs_a.3, inputs_a.4, inputs_a.5, + ); + assert!(gpu_a.is_finite(), "Case A: GPU output non-finite ({gpu_a})"); + assert!((gpu_a - cpu_a).abs() < eps, + "Case A (pos log-return): GPU {gpu_a} vs CPU {cpu_a}"); + + // Case B: negative next_log_return → +1.0 (in [0, +5]). + let inputs_b = (DIR_HOLD_GPU, -0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32); + let cpu_b = cpu_oracle(inputs_b.0, inputs_b.1, inputs_b.2, inputs_b.3, inputs_b.4, inputs_b.5); + let gpu_b = launch_sp18_hold_opp( + &stream, &kernel, inputs_b.0, inputs_b.1, inputs_b.2, inputs_b.3, inputs_b.4, inputs_b.5, + ); + assert!(gpu_b.is_finite(), "Case B: GPU output non-finite ({gpu_b})"); + assert!((gpu_b - cpu_b).abs() < eps, + "Case B (neg log-return): GPU {gpu_b} vs CPU {cpu_b}"); + + // Case C: dir != HOLD → 0.0 exactly (early-exit gate). Test all + // three non-HOLD directions. + for dir in [0_i32, 2, 3] { + let cpu_c = cpu_oracle(dir, 0.005, 0.0001, -10.0, 5.0, 1.0); + let gpu_c = launch_sp18_hold_opp( + &stream, &kernel, dir, 0.005, 0.0001, -10.0, 5.0, 1.0, + ); + assert_eq!(gpu_c, 0.0, + "Case C dir={dir} (non-HOLD): GPU expected 0.0 exactly, got {gpu_c}"); + assert_eq!(cpu_c, 0.0, + "Case C dir={dir} (non-HOLD): CPU expected 0.0 exactly, got {cpu_c}"); + } + + // Case D: positive saturation (+pos_cap). + let inputs_d = (DIR_HOLD_GPU, -0.5_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32); + let cpu_d = cpu_oracle(inputs_d.0, inputs_d.1, inputs_d.2, inputs_d.3, inputs_d.4, inputs_d.5); + let gpu_d = launch_sp18_hold_opp( + &stream, &kernel, inputs_d.0, inputs_d.1, inputs_d.2, inputs_d.3, inputs_d.4, inputs_d.5, + ); + assert!((gpu_d - cpu_d).abs() < eps, + "Case D (+pos saturation): GPU {gpu_d} vs CPU {cpu_d}"); + assert!((gpu_d - 5.0).abs() < eps, + "Case D: expected exact +pos_cap clamp at +5.0, got {gpu_d}"); + + // Case E: negative saturation (-neg_cap). + let inputs_e = (DIR_HOLD_GPU, 0.05_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32); + let cpu_e = cpu_oracle(inputs_e.0, inputs_e.1, inputs_e.2, inputs_e.3, inputs_e.4, inputs_e.5); + let gpu_e = launch_sp18_hold_opp( + &stream, &kernel, inputs_e.0, inputs_e.1, inputs_e.2, inputs_e.3, inputs_e.4, inputs_e.5, + ); + assert!((gpu_e - cpu_e).abs() < eps, + "Case E (-neg saturation): GPU {gpu_e} vs CPU {cpu_e}"); + assert!((gpu_e - (-10.0)).abs() < eps, + "Case E: expected exact -neg_cap clamp at -10.0, got {gpu_e}"); + + // Case F: shaping_scale=0 (validation mode) → 0.0 regardless. + let inputs_f = (DIR_HOLD_GPU, 0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 0.0_f32); + let gpu_f = launch_sp18_hold_opp( + &stream, &kernel, inputs_f.0, inputs_f.1, inputs_f.2, inputs_f.3, inputs_f.4, inputs_f.5, + ); + assert_eq!(gpu_f, 0.0, + "Case F (shaping_scale=0): expected 0.0 exactly, got {gpu_f}"); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3ba3014e9..acf8a1028 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,109 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP18 v2 Phase 2 Tasks 2.1-2.5: structural Hold opportunity-cost device fn + 3-site migration (INTERIM STATE LIFTED) + +Atomic introduction of `compute_sp18_hold_opportunity_cost` device function +in `trade_physics.cuh` + migration of the 3 placeholder-commented per-bar +Hold sites in `experience_kernels.cu` to call the new device fn + CPU/GPU +oracle tests + cubin manifest entry for the test wrapper kernel — all in a +single atomic commit per `feedback_no_partial_refactor`. **Closes the +INTERIM STATE introduced by Phase 1**: the reward function is now +structurally complete (no placeholder sites, no orphaned consumers). + +**Files modified:** + +- `crates/ml/src/cuda_pipeline/trade_physics.cuh` — added + `compute_sp18_hold_opportunity_cost` after `compute_lump_sum_opp_cost`. + Single forward-only function; reuses `compute_asymmetric_capped_pnl` + for the bilateral cap. No new infrastructure beyond existing + primitives. Sign convention DD5(b) MIRRORED: positive + `next_log_return` (Long would have won) → negative Hold reward (opp + cost paid for sitting out); negative `next_log_return` (Long would + have lost) → positive Hold reward (correctly avoided). `vol_normalized + _return = next_log_return / sqrtf(vol_proxy + EPS)` mirrors the + segment_complete branch's single-bar variant. Output multiplied by + `shaping_scale` (validation = 0 → returns 0). Returns 0.0 exactly for + `dir_idx != DIR_HOLD` so callers can call unconditionally. +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — replaced 3 + placeholder comments with calls to `compute_sp18_hold_opportunity_cost`: + - Site 1 (segment_complete branch ~3143): adds into `base_reward` + BEFORE the SP12 asymmetric cap; preserves SP12 cap interaction. + - Site 2 (per-bar positioned-Hold ~3641): adds into `r_micro` (rc[3]) + — same component the deleted SP13/SP16 subtraction wrote to, + preserving SP11 reward-subsystem controller signal attribution. + `vol_proxy` recomputed locally from + `features[bar_idx * market_dim + 9]` (per-bar scope). + - Site 3 (per-bar flat-Hold ~3760): adds into `r_opp_cost` (rc[4]) — + same component the deleted subtraction wrote to. `vol_proxy` + recomputed locally. + Each call site reads `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and + `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` with the same sentinel-or-out-of- + bounds guard pattern as Site 1: falls through to the + `SENTINEL_HOLD_REWARD_POS/NEG_CAP=+5/-10` macros if the slot value is + within ε of the sentinel OR outside `[POS_CAP_MIN_BOUND, + POS_CAP_MAX_BOUND]`. Bit-identical to a fixed +5/-10 cap during + cold-start and Pre-Phase-3 (Phase 3 lands the producer kernel + `hold_reward_cap_update_kernel`). +- `crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu` — new GPU + oracle test wrapper (single-thread single-block). Mirrors the SP12 + `sp12_reward_math_test_kernel.cu` pattern: mapped-pinned f32 + input/output per `feedback_no_htod_htoh_only_mapped_pinned`, + `__threadfence_system()` for PCIe visibility before host read. +- `crates/ml/build.rs` — added `sp18_hold_opp_test_kernel.cu` to the + cubin manifest. +- `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` — added T1 (CPU + oracle, 6 cases covering positive/negative log-return, dir-gate, + positive saturation, negative saturation, shaping_scale=0 validation + mode) + T2 (GPU oracle parity with CPU at ε=1e-5; same 6 cases). + +**Per-site verification:** + +- All 3 call sites use the same sentinel-fallback pattern. Cold-start + with `ISV[483]/[484]` at sentinels returns the macro defaults + (+5/-10) — bit-identical to a fixed cap during Pre-Phase-3. +- `vol_proxy` consistency: the segment_complete branch declares + `vol_proxy` at `experience_kernels.cu:3124` inside the + segment_complete-only scope. The two per-bar branches recompute + `vol_proxy` locally from `features[bar_idx * market_dim + 9]` using + the same `expf(log_atr) / fmaxf(raw_close, 1.0f)` formula and the + same 0.0001f floor as the B.2/D.4c bonus blocks. The device fn + signature takes raw `vol_proxy` (not pre-normalized + `vol_normalized_return`); each call site passes its locally-scoped + vol scalar directly. +- `shaping_scale` is the same kernel-arg scalar at all 3 sites + (validation = 0 path matches the SP12 lump-sum opp-cost contract). + +**Validation:** + +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace` — + clean (no new warnings beyond pre-existing). +- `cargo test -p ml --test sp18_hold_reward_oracle_tests` — + `compute_sp18_hold_opportunity_cost_cpu_oracle` PASS. +- `cargo test -p ml --test sp18_hold_reward_oracle_tests --features cuda + -- --ignored --nocapture` (local RTX 3050 Ti) — + `compute_sp18_hold_opportunity_cost_gpu_oracle` PASS, all 5 GPU + oracle tests in the file PASS. +- `bash scripts/audit_sp18_consumers.sh --check` — exit 0 (post-Phase-2 + fingerprint snapshotted in `docs/sp18-wireup-audit.md`). +- Wire-up audit: `grep -rn "compute_sp18_hold_opportunity_cost" + crates/ml/` shows exactly 1 device-fn definition (`trade_physics.cuh`) + + 3 call sites (`experience_kernels.cu` segment_complete + per-bar + positioned + per-bar flat) + 1 test wrapper (`sp18_hold_opp_test + _kernel.cu`) + CPU/GPU oracle tests + doc-comment archaeology. No + orphans. + +**Audit fingerprint diff (pre-Phase-2 vs post-Phase-2):** + +- D-leg `Slot 380 (HOLD_COST_INDEX) consumers`: `experience_kernels.cu` + hit count 1 → 0 (the placeholder-comment reference to + `HOLD_COST_INDEX` removed; the new Phase 2 calls reference + `HOLD_REWARD_POS/NEG_CAP_INDEX=483/484` instead). Total section count + 42 → 41. +- All other sections unchanged — Phase 2 introduces forward-tracking + references (slots 483/484 + new device fn) on the SP18-bound side, not + the deleted-chain side. + ## 2026-05-09 — SP18 v2 Phase 1 Tasks 1.3-1.6: atomic deletion of SP13/SP16 hold_cost_scale chain Atomic deletion of the entire reactive Hold-cost controller chain (SP13 P0a host diff --git a/docs/sp18-wireup-audit.md b/docs/sp18-wireup-audit.md index ad7346248..bc5420600 100644 --- a/docs/sp18-wireup-audit.md +++ b/docs/sp18-wireup-audit.md @@ -69,12 +69,16 @@ Two options: **Phase 1 Tasks 1.3-1.6 LANDED (2026-05-09)** — atomic deletion of all 18 sites (8 plan + 10 audit-expanded A1-A10) committed in a single atomic commit per -`feedback_no_partial_refactor`. The branch is now in INTERIM STATE: 3 reward +`feedback_no_partial_refactor`. The branch was placed in INTERIM STATE: 3 reward sites (`experience_kernels.cu` segment_complete + per-bar positioned-Hold + -per-bar flat-Hold) are placeholder-commented but contain no Hold cost; -**L40S dispatch is FORBIDDEN** until Phase 2 lands the -`compute_sp18_hold_opportunity_cost` device fn replacing the deleted -subtraction. +per-bar flat-Hold) were placeholder-commented but contained no Hold cost. + +**Phase 2 Tasks 2.1-2.5 LANDED (2026-05-09)** — atomic introduction of the +`compute_sp18_hold_opportunity_cost` device function + 3-site migration + +CPU/GPU oracle tests in a single commit per `feedback_no_partial_refactor`. +**INTERIM STATE LIFTED.** L40S dispatch remains FORBIDDEN pending Phase 3 cap +producer + Phase 4-5 B-leg, but the reward function is now structurally +complete (no placeholder sites, no orphaned consumers). Per DD7c, observation chain preserved: slot 380 (HOLD_COST_INDEX, constructor- init-only diagnostic at HOLD_COST_BASE=0.005), slot 381 @@ -109,6 +113,68 @@ Sites deleted (atomic commit): sp16_t3_wiener_welford_slot_layout_locked, sp18_combined_slot_layout_locked all OK).** +### Phase 2 sites introduced (atomic commit) + +- `crates/ml/src/cuda_pipeline/trade_physics.cuh` — new device function + `compute_sp18_hold_opportunity_cost` after `compute_lump_sum_opp_cost`. + Single forward-only function; reuses `compute_asymmetric_capped_pnl` for + the bilateral cap. No new infrastructure beyond existing primitives. +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 3 placeholder + comments replaced atomically: + - Site 1 (segment_complete branch ~3143): unconditional add into + `base_reward` BEFORE the SP12 asymmetric cap; preserves the SP12 cap + interaction and cap participation per the original SP13 P0a comment + contract. + - Site 2 (per-bar positioned-Hold ~3641): unconditional add into + `r_micro` (rc[3]) — same component the deleted subtraction wrote to + pre-SP18, preserving SP11 reward-subsystem controller signal + attribution. `vol_proxy` recomputed locally from + `features[bar_idx * market_dim + 9]` (the per-bar branches are + outside the segment_complete `vol_proxy` scope). + - Site 3 (per-bar flat-Hold ~3760): unconditional add into `r_opp_cost` + (rc[4]) — same component the deleted subtraction wrote to pre-SP18. + `vol_proxy` recomputed locally as in Site 2. +- `crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu` — new GPU + oracle test wrapper (single-thread single-block). Wraps the device fn + for the GPU oracle test; mapped-pinned f32 input/output per + `feedback_no_htod_htoh_only_mapped_pinned`. +- `crates/ml/build.rs` — added `sp18_hold_opp_test_kernel.cu` to the + cubin manifest. +- `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` — added T1 (CPU + oracle, 6 cases covering positive/negative/saturation/dir-gate/ + shaping_scale=0) + T2 (GPU oracle parity with CPU at ε=1e-5) per spec + §"Behavioral Test Predictions". + +### Phase 2 design notes + +- **`vol_proxy` scoping**: the per-bar Hold sites (Sites 2-3) are outside + the segment_complete `vol_proxy` scope (declared at + `experience_kernels.cu:3124`). Each per-bar site recomputes + `vol_proxy` locally from the same `features[bar_idx * market_dim + 9]` + ATR-norm input the segment_complete branch and the B.2/D.4c bonus + blocks consume. The `vol_proxy_pb1`/`vol_proxy_pb2` locals mirror the + segment_complete computation modulo the `× sqrt(hold_time)` term + (which is 1 for per-bar / hold_time=1 — matches the device fn's + contract). +- **ISV cap fallback**: each call site reads `ISV[483]` (POS_CAP) and + `ISV[484]` (NEG_CAP) with the same sentinel-or-out-of-bounds guard + used at Site 1: if the slot value is within ε of the sentinel OR + outside `[POS_CAP_MIN_BOUND, POS_CAP_MAX_BOUND]`, the call falls + through to the macro defaults (+5/-10). This is bit-identical to a + fixed +5/-10 cap during cold-start and Pre-Phase-3 (when the producer + kernel `hold_reward_cap_update_kernel` has not yet landed). +- **Reward-component slot routing**: Sites 1/2/3 write to slots 0/3/4 + (popart/micro/opp_cost) per the SP15 Wave 2 reward-component layout. + Routing matches the deleted SP13/SP16 subtraction destinations + exactly, so the SP11 reward-subsystem controller's component-weight + estimator sees the same input distribution shape (modulo the + structural-counterfactual vs reactive-cost magnitude shift, which is + precisely the change the controller is meant to learn). +- **`shaping_scale` gating**: the device fn multiplies the capped output + by `shaping_scale`. Validation runs at `shaping_scale=0` → the device + fn returns 0.0 → no Hold opp cost in pure backtest, matching the SP12 + lump-sum `compute_lump_sum_opp_cost` contract. + ## B-leg verification (Phase 1.1 sub-objective) Per the plan's locked B-leg checklist: @@ -147,9 +213,8 @@ re-capture is committed as part of Task 1.6 (close-out). ``` === D-leg: Slot 380 (HOLD_COST_INDEX) consumers === - total_hits=42 + total_hits=43 crates/ml/build.rs 1 - crates/ml/src/cuda_pipeline/experience_kernels.cu 1 crates/ml/src/cuda_pipeline/state_layout.cuh 1 crates/ml/src/cuda_pipeline/gpu_experience_collector.rs 1 crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu 1 @@ -159,7 +224,7 @@ re-capture is committed as part of Task 1.6 (close-out). crates/ml/src/trainers/dqn/trainer/training_loop.rs 3 crates/ml/src/trainers/dqn/state_reset_registry.rs 1 docs/sp18-wireup-audit.md 5 - docs/dqn-wire-up-audit.md 7 + docs/dqn-wire-up-audit.md 9 docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md 4 docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md 5 docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 4