From 06dbb78ffc1c7d86505a5b13f7b21699bc25f2f2 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 01:46:52 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp20):=20Phase=203=20Task=203.2=20?= =?UTF-8?q?=E2=80=94=20Hold=20opportunity-cost=20dual=20emission=20(Compon?= =?UTF-8?q?ent=202)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Lands the per-bar Hold opportunity-cost dual emission at experience_env_step's per-bar branches (positioned-non-event ~3650 + flat-non-event ~3737), atomically with the trade-close consumer wiring at the segment_complete branch (~3289 — replaces the Phase 2 Task 2.2 forward-reference placeholder `hold_baseline = 0.0f`) per `feedback_no_partial_refactor`. Spec §4.2 dual-emission contract: per_bar_opp_cost = -aux_conf × cost_scale Path 1 (Hold-only): r_micro/r_opp_cost += per_bar_opp_cost - ISV[HOLD_REWARD_EMA] (centered for Q-target stability) Path 2 (always): hold_baseline_buffer[env, t % 30] = per_bar_opp_cost (uncentered, consumed at trade close) At trade close: hold_baseline = sp20_sum_hold_baseline_over_trade(buf, env, 30, current_t, segment_hold_time) alpha = R_event - hold_baseline # replaces Phase 2 placeholder The summation walks backwards `min(30, segment_hold_time)` slots from `(current_t - 1) % 30` in env i's row of the per-env circular buffer. The trade-close bar's segment_complete branch does NOT write to the buffer, so the most recent per-bar write is at current_t - 1. Layout decision (plan errata Gap 9): per-env stride `[N_envs × HOLD_BASELINE_BUFFER_SIZE = 30]` row-major. The kernel is per-env-parallel; a single global ring would interleave bars across envs and break trade-attribution semantic. Trade-open tracking decision (plan errata Gap 10): NO new state slot needed. The existing `segment_hold_time` (= saved_hold_time captured before PS_HOLD_TIME reset at experience_kernels.cu:2618) plus current_t suffice — walk backwards in the buffer. Replaces SP18 D-leg `compute_sp18_hold_opportunity_cost` calls at experience_kernels.cu:3672 and :3783. The device function in trade_physics.cuh:655 is RETAINED — still called by sp18_hold_opp_test_kernel.cu (oracle test surface). Per `feedback_no_stubs`: only production callers migrate; the helper stays for the test surface. New device helpers (sp20_hold_baseline.cuh): - sp20_compute_per_bar_aux_conf_k2(logits) — bit-identical to the formula in sp20_stats_compute_kernel.cu Pass A. - sp20_sum_hold_baseline_over_trade(buf, env, size, current_t, segment_hold_time) — walks backwards with modulo wrap-around. New buffer + reset infrastructure: - GpuExperienceCollector.hold_baseline_buffer field - HOLD_BASELINE_BUFFER_SIZE = 30 constant (re-exported) - StateResetRegistry FoldReset entry + invariant test - reset_named_state dispatch arm Six new kernel args appended to experience_env_step signature: aux_logits_per_env, hold_baseline_buffer, hold_buffer_size, aux_k, hold_cost_scale_idx, hold_reward_ema_idx. All NULL-tolerant. HOLD_REWARD_EMA forward-reference: ISV[516] is updated by sp20_emas_compute_kernel reading per_bar_hold_reward from sp20_aggregate_inputs_kernel which currently emits 0.0f as a Phase 3.2 forward-ref placeholder (line 292). The parallel `sp20-phase-2-fix` agent owns wiring the real producer; Task 3.2's centering math references the ISV slot (not a hardcoded 0), so when the parallel branch lands, EMA starts updating and centering becomes load-bearing automatically. No additional Task 3.2 work needed post-merge. Tests (sp20_hold_baseline_test.rs, 4 GPU oracle tests): 1. aux_conf_k2_bounds_and_fixed_points — bounds, uniform/saturated/ spec-warmup/spec-confident/symmetry fixed points. 2. sum_hold_baseline_over_trade_indexing — basic indexing, hold_time>size clamp, defensive guards, modulo wrap-around. 3. sum_hold_baseline_per_env_stride — env-stride correctness across row-major buffer. 4. dual_emission_hold_vs_non_hold — full Path 1 + Path 2 contract, mirrors plan §3.2 reference fixture. Verification: SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # green SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 21/21 pass SQLX_OFFLINE=true cargo build -p ml # cubin compile Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 15 + .../src/cuda_pipeline/experience_kernels.cu | 334 ++++++++------ .../cuda_pipeline/gpu_experience_collector.rs | 148 +++++-- .../src/cuda_pipeline/sp20_hold_baseline.cuh | 168 +++++++ .../sp20_hold_baseline_test_kernel.cu | 134 ++++++ .../src/trainers/dqn/state_reset_registry.rs | 60 ++- .../src/trainers/dqn/trainer/training_loop.rs | 15 +- crates/ml/tests/sp20_hold_baseline_test.rs | 412 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 222 +++++++++- 9 files changed, 1312 insertions(+), 196 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh create mode 100644 crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu create mode 100644 crates/ml/tests/sp20_hold_baseline_test.rs diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 33b304715..67d0c990f 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -359,6 +359,21 @@ fn main() { // Task 2.2 inside `experience_env_step::segment_complete` branch; // the test wrapper has no production callers. "sp20_event_reward_test_kernel.cu", + // SP20 Phase 3 Task 3.2 (2026-05-10): standalone test wrappers + // for the two device helpers defined in `sp20_hold_baseline.cuh` + // (Component 2 — Hold opportunity-cost dual emission). Wraps + // `sp20_compute_per_bar_aux_conf_k2`, + // `sp20_sum_hold_baseline_over_trade`, and a fixture mirroring + // the production per-bar emission site so the Rust GPU oracle + // tests in `tests/sp20_hold_baseline_test.rs` drive the EXACT + // same code paths the production caller + // (`experience_env_step` per-bar branches + segment_complete + // trade-close site) executes per + // `feedback_no_cpu_test_fallbacks`. Producer call sites land in + // Task 3.2 inside `experience_env_step` (per-bar branches + + // segment_complete branch). Test wrapper has no production + // callers. + "sp20_hold_baseline_test_kernel.cu", // SP4 Layer A Task A5 (2026-04-30): first end-to-end SP4 producer. // Single-block kernel reads `denoise_target_q_buf`, computes // p99(|target_q|) via `sp4_histogram_p99<256>`, writes the step diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index ffed74171..e37be717e 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -97,6 +97,7 @@ * `experience_env_step`'s segment_complete branch atomically with the * SP12 v3 reward block deletion + alpha plumbing. */ #include "sp20_reward.cuh" +#include "sp20_hold_baseline.cuh" /* Phase 3.2: Component 2 helpers */ /* ------------------------------------------------------------------ */ /* BF16 support — common_device_functions.cuh is prepended by build.rs */ @@ -2204,21 +2205,51 @@ extern "C" __global__ void experience_env_step( float* __restrict__ alpha_per_env, int loss_cap_idx, int alpha_ema_idx, - /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — per-env trade-close win - * flag output buffer. + /* SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual + * emission (Component 2) per spec §4.2. * - * `is_win_per_env` ← [N] i32: written at `is_close && segment_ - * hold_time > 0.0f` (segment_complete branch) as - * `(segment_return > 0.0f) ? 1 : 0`. Mirrors the alpha emit site - * exactly — the SP20 aggregation kernel reads alpha sum and - * is_win count side-by-side at every closed env. Replaces the - * broken per-bar `step_ret > 0` predicate the aggregation kernel - * was using, which pinned WR_EMA at 0 in production (per-bar - * step_ret at close bars is dominated by tx-cost so almost always - * negative even for trades that closed profitably). NULL-tolerant - * for test scaffolds — when NULL the write is a no-op and the - * aggregation kernel falls back to the legacy `sr > 0` predicate. */ - int* __restrict__ is_win_per_env + * `aux_logits_per_env` ← [N, K=2] f32 row-major device-resident. + * Producer: `aux_next_bar_forward` kernel (writes + * `exp_aux_nb_logits_buf`) — same producer the `sp20_stats_compute` + * kernel reads at the start of the SP20 chain. Stream-implicit + * ordering: the aux head runs at the same step's "step 3" before + * `experience_env_step` (step 5) on the same stream, so the + * logits are valid here. NULL-tolerant for test scaffolds — + * sp20_compute_per_bar_aux_conf_k2 isn't called and the dual + * emission collapses to zero (no buffer write, no centered reward) + * which preserves the existing per-bar behavior bit-identically. + * + * `hold_baseline_buffer` ← [N × hold_buffer_size] f32 row-major + * device-resident scratch. Per-env circular ring: each env writes + * its own row at `current_t % hold_buffer_size` per bar; at trade + * close, sums backwards `min(hold_buffer_size, segment_hold_time)` + * slots (consumer at the segment_complete branch). Producer: + * THIS kernel writes per-bar at the positioned-non-event + + * flat-non-event branches (always, regardless of action — Path 2 + * counterfactual semantic per spec §4.2). NULL = no buffer write + * AND consumer falls back to `hold_baseline = 0.0f` (Phase 2 Task + * 2.2 placeholder behaviour, preserving backward compat for test + * scaffolds without Phase 3.2 wiring). + * + * `hold_buffer_size`: circular ring size, 30 per spec §4.2 buffer + * sizing note (matches SP19's longest horizon LOOKAHEAD_HORIZON_MAX). + * + * `aux_k`: K=2 for the binary up/not-up aux head. The + * `sp20_compute_per_bar_aux_conf_k2` helper hardcodes K=2 (matches + * the spec §4.2 K=2 AMENDED note); `aux_k` is passed in for + * future-proofing and as a defensive sanity check (currently + * asserted-equivalent inside the kernel). + * + * `hold_cost_scale_idx` / `hold_reward_ema_idx`: ISV slot indices + * for the producer / consumer reads. Caller passes them as kernel + * args so the kernel stays decoupled from SP14 slot-number drift + * (mirror the `q_dir_abs_ref_idx` pattern at line 2069). */ + const float* __restrict__ aux_logits_per_env, + float* __restrict__ hold_baseline_buffer, + int hold_buffer_size, + int aux_k, + int hold_cost_scale_idx, + int hold_reward_ema_idx ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -3301,13 +3332,35 @@ extern "C" __global__ void experience_env_step( const float r_event = sp20_compute_event_reward( segment_return, label_at_open_sign, loss_cap ); - /* Phase 3.2 placeholder: hold_baseline = 0.0f. Component 2 - * (Phase 3.2) lands the real `hold_baseline_buffer.sum_over_ - * trade_range()` consumer. Forward reference per the audit-doc - * Task 2.2 entry — NOT a stub: the producer site here is fully - * wired to a working bounded scalar, the consumer just adds - * a (negative) opportunity-cost subtraction in a later phase. */ - const float hold_baseline = 0.0f; + /* SP20 Phase 3 Task 3.2 (2026-05-10): hold_baseline = sum over + * the trade's bars of the per-env circular `hold_baseline_buffer`. + * Replaces the Phase 2 Task 2.2 forward-reference placeholder + * (`hold_baseline = 0.0f`) per spec §4.2. + * + * Helper walks backwards from `current_t - 1` (the most recent + * per-bar write — the trade-close bar's segment_complete branch + * does NOT write to the buffer) for `min(hold_buffer_size, + * segment_hold_time)` bars in the env's row. NULL-tolerant: + * test scaffolds without Phase 3.2 wiring fall back to 0.0 + * (preserves the Phase 2 placeholder behaviour bit-identically). + * + * Bounded by construction (`pearl_one_unbounded_signal_per_reward`): + * each per-bar entry is `-aux_conf × cost_scale ∈ [-0.25, 0]` + * (aux_conf ∈ [0, 0.5], cost_scale ∈ [0.01, 0.5]); sum of + * 30 such terms is in `[-7.5, 0]`. With `r_event ∈ [-2, +1]` + * (4-quadrant table from sp20_reward.cuh), the resulting + * `alpha = r_event - hold_baseline ∈ [-2, +8.5]`. Through the + * Wiener-α `alpha_ema` EMA bound (also in `[-2, +8.5]` + * post-bootstrap), the centered `r_used = alpha - alpha_ema` + * stays bounded — no clamps needed; entirely structural. */ + const float hold_baseline = (hold_baseline_buffer != NULL) + ? sp20_sum_hold_baseline_over_trade( + hold_baseline_buffer, + /*env=*/i, + hold_buffer_size, + current_t, + segment_hold_time) + : 0.0f; const float alpha = r_event - hold_baseline; const float alpha_ema = (isv_signals_ptr != NULL) ? isv_signals_ptr[alpha_ema_idx] @@ -3642,72 +3695,90 @@ 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; - /* 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. + /* SP20 Phase 3 Task 3.2 (2026-05-10): Hold opportunity-cost dual + * emission (Component 2) per spec §4.2. Replaces the SP18 D-leg + * `compute_sp18_hold_opportunity_cost` per-bar call atomically + * with the new producer wiring (per + * `feedback_no_partial_refactor` "consumer migrates with + * contract change" — every per-bar Hold-cost site migrates in + * one commit). The SP18 `compute_sp18_hold_opportunity_cost` + * device function in `trade_physics.cuh:655` is RETAINED — + * still called by `sp18_hold_opp_test_kernel.cu` (the SP18 + * oracle test surface). * - * 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). + * Two parallel emissions per spec §4.2: * - * 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]. + * per_bar_opp_cost = -aux_conf × cost_scale * - * 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); + * Path 1 — real reward (Hold-action only): + * r_micro += per_bar_opp_cost - ISV[HOLD_REWARD_EMA] + * (centered for Q-target scale comparability — Q(Hold) and + * Q(trade) live on the same relative-to-mean scale) + * + * Path 2 — counterfactual buffer (always): + * hold_baseline_buffer[env, current_t % size] = per_bar_opp_cost + * (uncentered — consumed by Component 1 trade-close site + * to compute alpha = R_event - sum_over_trade_range) + * + * NULL-tolerant: when `aux_logits_per_env == NULL` (test + * scaffolds without aux-head wiring) BOTH paths skip — the + * branch reduces to `r_micro = 0.0f` (existing behavior + * pre-Phase-3.2), preserving backward compat. When + * `hold_baseline_buffer == NULL`, only Path 2 is skipped; + * Path 1 still fires (centered Hold reward into r_micro) + * because the Hold reward is independent of the buffer. + * + * Bounded by construction (`pearl_one_unbounded_signal_per_reward`): + * aux_conf ∈ [0, 0.5] (softmax peak above K=2 uniform) + * cost_scale ∈ [0.01, 0.5] (controller bound, ISV slot 513) + * per_bar_opp_cost ∈ [-0.25, 0] + * ISV[HOLD_REWARD_EMA] ∈ [-0.25, 0] (Wiener-α EMA of bounded input) + * r_micro centered contribution ∈ [-0.25, +0.25] + * + * Pearl notes: + * - `pearl_event_driven_reward_density_alignment` — Component + * 2 IS the per-bar density alignment; the trade-close + * summation re-aligns it with the trade-event density via + * subtraction. This is intentional density-alignment, NOT + * the per-bar shaping anti-pattern (which was redundant + * pressure on the same TD signal). The Hold-Q dimension is + * SEPARATE from the trade-Q dimension; both need their own + * density alignment. + * - `pearl_one_unbounded_signal_per_reward` — softmax-bound + * and controller-bound multiplicands (no unbounded factor). + * - `pearl_first_observation_bootstrap` — `ISV[HOLD_COST_SCALE]` + * starts at 0.0 sentinel; cold-start gives per_bar_opp_cost + * = 0 until the controller's first observation lifts the + * scale. Path 2 writes 0 to the buffer (harmless — sums to + * 0 at cold-start trade-close). + */ + const int has_aux = (aux_logits_per_env != NULL); + const int has_hold_buf = (hold_baseline_buffer != NULL); + if (has_aux) { + /* aux_k=2 hardcoded per spec §4.2 AMENDED note. The + * `aux_k` arg is reserved for future K>2 support. */ + const float* logits_row = aux_logits_per_env + (long long)i * (long long)aux_k; + const float aux_conf_i = sp20_compute_per_bar_aux_conf_k2(logits_row); + const float cost_scale = (isv_signals_ptr != NULL) + ? isv_signals_ptr[hold_cost_scale_idx] + : 0.0f; + const float per_bar_opp_cost = -aux_conf_i * cost_scale; - 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; - } + /* Path 2 — always written when buffer wired. */ + if (has_hold_buf) { + const long long row_off = + (long long)i * (long long)hold_buffer_size; + const int slot_idx = current_t % hold_buffer_size; + hold_baseline_buffer[row_off + slot_idx] = per_bar_opp_cost; + } + + /* Path 1 — Hold-action only, centered. */ + if (dir_idx == DIR_HOLD) { + const float hold_reward_ema = (isv_signals_ptr != NULL) + ? isv_signals_ptr[hold_reward_ema_idx] + : 0.0f; + r_micro += per_bar_opp_cost - hold_reward_ema; } - 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; @@ -3763,62 +3834,57 @@ 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; - /* 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. + /* SP20 Phase 3 Task 3.2 (2026-05-10): Hold opportunity-cost dual + * emission on flat-non-event bars (Component 2) — 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 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. + * Same dual-emission contract as the positioned branch: * - * 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); + * Path 1 — real reward into `r_opp_cost` (Hold-action only, + * centered): `r_opp_cost += per_bar_opp_cost - + * ISV[HOLD_REWARD_EMA]`. + * Path 2 — counterfactual buffer write (always — even on + * flat-non-event bars; the Hold counterfactual + * applies symmetrically across positioned/flat). + * + * Same NULL-tolerance + boundedness invariants as the + * positioned-branch documentation above. The two branches + * write to DIFFERENT reward components (rc[3] r_micro for + * positioned, rc[4] r_opp_cost for flat) per SP11 component + * attribution semantics — preserved from the SP18 D-leg + * pre-Phase-3.2 routing. */ + const int has_aux_pb2 = (aux_logits_per_env != NULL); + const int has_hold_buf_pb2 = (hold_baseline_buffer != NULL); + if (has_aux_pb2) { + const float* logits_row_pb2 = aux_logits_per_env + + (long long)i * (long long)aux_k; + const float aux_conf_pb2 = + sp20_compute_per_bar_aux_conf_k2(logits_row_pb2); + const float cost_scale_pb2 = (isv_signals_ptr != NULL) + ? isv_signals_ptr[hold_cost_scale_idx] + : 0.0f; + const float per_bar_opp_cost_pb2 = -aux_conf_pb2 * cost_scale_pb2; - 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; - } + /* Path 2 — always written when buffer wired. */ + if (has_hold_buf_pb2) { + const long long row_off_pb2 = + (long long)i * (long long)hold_buffer_size; + const int slot_idx_pb2 = current_t % hold_buffer_size; + hold_baseline_buffer[row_off_pb2 + slot_idx_pb2] = + per_bar_opp_cost_pb2; + } + + /* Path 1 — Hold-action only, centered. Routes through + * r_opp_cost (rc[4]) per the flat-branch SP11 attribution. */ + if (dir_idx == DIR_HOLD) { + const float hold_reward_ema_pb2 = (isv_signals_ptr != NULL) + ? isv_signals_ptr[hold_reward_ema_idx] + : 0.0f; + r_opp_cost += per_bar_opp_cost_pb2 - hold_reward_ema_pb2; } - 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/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 2aa145dc7..159ca65f9 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -51,6 +51,15 @@ const PORTFOLIO_STRIDE: usize = 43; // was 41. Slots 30-37: prev-bar OFI; slot /// Layout: [win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns] const TRADE_STATS_FLOATS: usize = 6; +/// SP20 Phase 3 Task 3.2 (2026-05-10): circular per-env Hold +/// opportunity-cost buffer size. Per spec §4.2: "circular buffer of +/// size LOOKAHEAD_HORIZON_MAX = 30 bars (matches SP19's longest +/// horizon and bounds max trade duration)". Bars older than 30 are +/// overwritten (acceptable per spec — trades > 30 bars don't sample +/// beyond their last 30 bars). Used in `experience_env_step` Hold +/// opportunity-cost dual-emission (Component 2). +pub const HOLD_BASELINE_BUFFER_SIZE: usize = 30; + // --------------------------------------------------------------------------- // TradeStats — GPU-sourced per-epoch trade statistics // --------------------------------------------------------------------------- @@ -706,34 +715,56 @@ pub struct GpuExperienceCollector { /// don't leak into the EMA. pub(crate) alpha_per_env: CudaSlice, // [alloc_episodes] - /// SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade-close win - /// flag. `[alloc_episodes]` i32 device-resident scratch — `1` iff - /// the segment's realized P&L was strictly positive at trade close - /// (`segment_return > 0.0f`), `0` otherwise. Replaces the broken - /// per-bar `step_ret > 0` predicate the aggregation kernel was - /// using, which conflated the close bar's per-bar mark-to-market - /// tick (dominated by tx-cost so almost always false at close - /// bars) with trade outcome — pinning WR_EMA at 0 across all - /// training epochs in production. + /// SP20 Phase 3 Task 3.2 (2026-05-10): per-env circular Hold + /// opportunity-cost buffer (Component 2) per spec §4.2. /// - /// Layout: `[alloc_episodes]` i32 — one slot per env, written by - /// `experience_env_step`'s `is_close` (`segment_complete`) branch - /// at the same site as `alpha_per_env`. Read once per rollout - /// step by `sp20_aggregate_inputs_kernel` to populate the - /// `wins_count` reduction (mirror of the alpha-mean aggregation). + /// Layout: `[alloc_episodes × HOLD_BASELINE_BUFFER_SIZE = 30]` f32 + /// row-major device-resident scratch. + /// + /// Producer: `experience_env_step` per-bar branches + /// (positioned-non-event line ~3650 + flat-non-event line ~3737) + /// — both write `hold_baseline_buffer[env, current_t % + /// HOLD_BASELINE_BUFFER_SIZE] = -aux_conf × cost_scale` UNCONDITIONALLY + /// (Path 2 counterfactual emission, regardless of action). The + /// segment_complete branch does NOT write — the trade-close bar's + /// reward path bypasses the per-bar branches. + /// + /// Consumer: `experience_env_step` `is_close && segment_hold_time + /// > 0.0f` branch — sums backwards `min(HOLD_BASELINE_BUFFER_SIZE, + /// segment_hold_time)` slots from `(current_t - 1) % + /// HOLD_BASELINE_BUFFER_SIZE` to compute `hold_baseline`, replacing + /// the Phase 2 Task 2.2 placeholder `hold_baseline = 0.0f`. The + /// resulting `alpha = R_event - hold_baseline` flows through to + /// `alpha_per_env[env]` for the SP20 aggregation kernel + /// (per_bar EMA on Hold-state bars only — the + /// `sp20_aggregate_inputs_kernel` Phase 3.2 forward-reference for + /// `per_bar_hold_reward` lands separately on the + /// `sp20-phase-2-fix` branch per concurrent-agent collision + /// avoidance). + /// + /// Buffer-sizing rationale (spec §4.2): 30 bars matches SP19's + /// longest horizon (`LOOKAHEAD_HORIZON_MAX`). Bars older than 30 + /// are overwritten — fine because trades > 30 bars don't sample + /// beyond their last 30 bars per the spec's bounded-trade-duration + /// argument. /// /// FoldReset semantics (registered in `state_reset_registry.rs` as - /// `is_win_per_env`): zero on every fold boundary so the new - /// fold's first trade-close fully populates the slot from the new - /// fold's segment_return calculation (no stale win flag leakage - /// across folds — each fold has independent WR_EMA bootstrap state - /// per Phase 1.4's `obs_count[OBS_COUNT_WR]` reset). + /// `hold_baseline_buffer`): zero on every fold boundary so the new + /// fold's first trade-close summation reads only this fold's + /// per-bar writes — no stale bars leak across folds. Reset path: + /// `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` + /// (device-resident, no host buffer to fill). /// - /// On non-close steps the slot retains the previous trade-close - /// flag (or 0 sentinel pre-first-close); the aggregation kernel - /// reads only when `trade_close_per_env[env] != 0`, so stale - /// values don't leak into the EMA. - pub(crate) is_win_per_env: CudaSlice, // [alloc_episodes] + /// Per `pearl_first_observation_bootstrap`: cold-start with + /// `ISV[HOLD_COST_SCALE]` at sentinel 0 means + /// `per_bar_opp_cost = 0` (Path 2 writes zeros to the buffer); + /// the trade-close consumer sums zeros and lands `hold_baseline = + /// 0.0f` (matches Phase 2 Task 2.2 placeholder behaviour bit-for- + /// bit on the very first trade close after fold reset). The + /// controller's first observation lifts cost_scale, then per_bar + /// writes become non-zero, then the next trade-close summation + /// engages — gracefully bootstrapped, no warmup gate needed. + pub(crate) hold_baseline_buffer: CudaSlice, // [alloc_episodes × 30] // Output buffers [alloc_episodes * alloc_timesteps, ...] states_out: CudaSlice, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32 @@ -1805,19 +1836,18 @@ impl GpuExperienceCollector { "sp20: alloc alpha_per_env ({alloc_episodes} f32): {e}" )))?; - // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade-close - // win flag (1 iff segment_return > 0 at trade close, else 0). - // Sentinel = 0; updated at each trade close by the same - // SP20 4-quadrant reward kernel block site that writes - // `alpha_per_env`. Consumed by sp20_aggregate_inputs_kernel - // (count over closed envs with flag=1) for the wins_count - // reduction. Replaces the broken per-bar `step_ret > 0` - // predicate that pinned WR_EMA at 0 in production. - // Sized [alloc_episodes]. - let is_win_per_env = stream - .alloc_zeros::(alloc_episodes) + // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold opportunity- + // cost circular buffer. Layout `[alloc_episodes × + // HOLD_BASELINE_BUFFER_SIZE = 30]` f32 row-major. Sentinel 0.0 + // — at cold-start (before the controller lifts ISV[513] off + // sentinel) per-bar writes are 0, so cold-start trade closes + // sum to 0.0 (bit-identical to the Phase 2 Task 2.2 placeholder + // semantic per `pearl_first_observation_bootstrap`). + let hold_baseline_buffer = stream + .alloc_zeros::(alloc_episodes * HOLD_BASELINE_BUFFER_SIZE) .map_err(|e| MLError::ModelError(format!( - "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}" + "sp20: alloc hold_baseline_buffer ({} f32): {e}", + alloc_episodes * HOLD_BASELINE_BUFFER_SIZE )))?; // Persistent epoch state @@ -2849,7 +2879,7 @@ impl GpuExperienceCollector { episode_starts_buf, label_at_open_per_env, alpha_per_env, - is_win_per_env, + hold_baseline_buffer, states_out, actions_out, rewards_out, @@ -6267,16 +6297,40 @@ impl GpuExperienceCollector { use crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX; ALPHA_EMA_INDEX as i32 }) - // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env - // trade-close win flag (output) — 1 iff - // `segment_return > 0` at trade close. Written at - // the same segment_complete branch site as - // `alpha_per_env`. Consumed by - // `sp20_aggregate_inputs_kernel` for the - // `wins_count` reduction (replaces the broken - // per-bar `step_ret > 0` predicate that pinned - // WR_EMA at 0 in production). - .arg(&mut self.is_win_per_env) + // SP20 Phase 3 Task 3.2 (2026-05-10): Hold + // opportunity-cost dual emission (Component 2). Six + // new args at the end of the kernel signature: + // - aux_logits_per_env [N, K=2] f32 input — + // reuses `exp_aux_nb_logits_buf` (the same + // buffer the SP20 stats kernel reads at + // line 6435 above). Stream-implicit + // producer→consumer: aux head writes at step 3 + // (line ~5328), env_step reads here at step 5. + // - hold_baseline_buffer [N × 30] f32 + // scratch — per-env circular ring written at + // per-bar branches, summed at trade close. + // - hold_buffer_size = 30 (HOLD_BASELINE_BUFFER_SIZE). + // - aux_k = 2 (AUX_NEXT_BAR_K — defensive + // companion to the K=2 hardcoded helper). + // - hold_cost_scale_idx, hold_reward_ema_idx — + // ISV slot indices (mirror the loss_cap_idx / + // alpha_ema_idx pattern above; single source + // of truth at sp14_isv_slots.rs constants). + .arg(&self.exp_aux_nb_logits_buf.raw_ptr()) + .arg(&mut self.hold_baseline_buffer) + .arg(&{ HOLD_BASELINE_BUFFER_SIZE as i32 }) + .arg(&{ + use super::gpu_aux_heads::AUX_NEXT_BAR_K; + AUX_NEXT_BAR_K as i32 + }) + .arg(&{ + use crate::cuda_pipeline::sp14_isv_slots::HOLD_COST_SCALE_INDEX; + HOLD_COST_SCALE_INDEX as i32 + }) + .arg(&{ + use crate::cuda_pipeline::sp14_isv_slots::HOLD_REWARD_EMA_INDEX; + HOLD_REWARD_EMA_INDEX as i32 + }) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" diff --git a/crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh b/crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh new file mode 100644 index 000000000..c18563ec5 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh @@ -0,0 +1,168 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual emission + * helpers (Component 2). + * + * Per spec `docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md` §4.2. + * + * Two helpers shared between `experience_env_step` (production caller) and + * the SP20 hold-baseline GPU oracle test wrapper: + * + * 1. `sp20_compute_per_bar_aux_conf_k2(logits)`: + * Per-row K=2 peak-confidence above the K=2 uniform mean. + * `aux_conf = max(softmax(logits)) - 0.5` ∈ [0, 0.5]. + * Numerically-stable: subtract row-max before exp. + * Bit-identical to the formula in `sp20_stats_compute_kernel.cu` + * Pass A (lines 152-167) — the per-row aux_conf computation that + * feeds the batch p50/std reduction. + * + * 2. `sp20_sum_hold_baseline_over_trade(buf, env, hold_buffer_size, + * current_t, segment_hold_time)`: + * Sum the per-env circular `hold_baseline_buffer` over the most + * recent `min(hold_buffer_size, segment_hold_time)` bars BEFORE + * `current_t`. The trade-close bar's reward path doesn't write + * to the buffer (segment_complete bypasses the per-bar branches), + * so the most recent write was at `(current_t - 1) % + * hold_buffer_size`, walking backwards. + * + * Layout: `buf[env * hold_buffer_size + (current_t - k) % + * hold_buffer_size]` for k = 1..=n_bars where + * `n_bars = min(hold_buffer_size, segment_hold_time)`. + * + * `segment_hold_time` is captured BEFORE the trade-physics helper + * resets `PS_HOLD_TIME` to 0 (saved_hold_time at + * experience_kernels.cu:2618), so it correctly reflects the trade's + * duration. + * + * Trades older than `hold_buffer_size` get UNDER-counted (only the + * most recent 30 bars are summed). Per spec §4.2 buffer-sizing + * note this is acceptable: "trades > 30 bars don't sample beyond + * their last 30 bars." + * + * ── Pearls + invariants ──────────────────────────────────────────────── + * - `pearl_event_driven_reward_density_alignment` — Component 2's + * per-bar density alignment is intentional: every bar contributes a + * counterfactual cost ("what would Hold have paid"), and the trade + * consumer sums the trade's bars to subtract from the trade-close + * R_event. Density-aligned by construction. + * - `pearl_one_unbounded_signal_per_reward` — `aux_conf ∈ [0, 0.5]` + * (softmax-bounded), `cost_scale ∈ [0.01, 0.5]` (controller-bounded), + * so `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`. Sum of + * 30 such terms is in `[-7.5, 0]`. Bounded by construction. + * - `feedback_no_atomicadd` — both helpers are pure read/compute (no + * writes to shared memory or device memory). + * - `feedback_no_cpu_test_fallbacks` — both helpers are __device__ + + * __forceinline__ so the test wrapper can call them bit-for-bit. + * - `feedback_isv_for_adaptive_bounds` — `cost_scale` is read from + * `ISV[HOLD_COST_SCALE_INDEX]` by the caller; this header takes it + * as a scalar arg, decoupled from slot indices. + * + * ── Wiring contract ──────────────────────────────────────────────────── + * + * Per-bar emission (positioned-non-event AND flat-non-event branches): + * - Read `aux_conf_i = sp20_compute_per_bar_aux_conf_k2(...)` once. + * - Read `cost_scale = ISV[HOLD_COST_SCALE_INDEX]`. + * - Compute `per_bar_opp_cost = -aux_conf × cost_scale`. + * - Write Path 2 (counterfactual): `hold_baseline_buffer[env * + * hold_buffer_size + current_t % hold_buffer_size] = per_bar_opp_cost` + * (always, regardless of action). + * - Write Path 1 (real reward): if `dir_idx == DIR_HOLD`, + * `r_micro_or_opp_cost += per_bar_opp_cost - ISV[HOLD_REWARD_EMA_INDEX]` + * (centered for Q-target stability). + * + * Trade-close consumer (segment_complete branch): + * - `hold_baseline = sp20_sum_hold_baseline_over_trade(...)` — + * replaces the Phase 2 Task 2.2 placeholder `hold_baseline = 0.0f`. + * - `alpha = R_event - hold_baseline` continues to feed + * `alpha_per_env[env]` for the SP20 aggregation kernel. + * + * ══════════════════════════════════════════════════════════════════════════ */ + +#ifndef SP20_HOLD_BASELINE_CUH +#define SP20_HOLD_BASELINE_CUH + +#include + +/* K=2 uniform softmax mean. Subtracted from peak softmax to produce + * `aux_conf ∈ [0, 0.5]`. Exactly representable in f32. */ +#define SP20_HOLD_BASELINE_UNIFORM_K2 0.5f + +/* Per-row K=2 peak-confidence: `max(softmax(logits)) - 0.5`. Numerically + * stable via row-max subtraction. Returns f32 in [0, 0.5]. */ +__device__ __forceinline__ float sp20_compute_per_bar_aux_conf_k2( + const float* __restrict__ logits_row /* [K=2] f32 */ +) { + const float l0 = logits_row[0]; + const float l1 = logits_row[1]; + const float m = fmaxf(l0, l1); + const float e0 = __expf(l0 - m); + const float e1 = __expf(l1 - m); + const float denom = e0 + e1; + /* denom >= 1.0f always (the larger of e0/e1 is exp(0) = 1.0 + * after row-max subtract), so no division-by-zero guard needed. */ + const float sm0 = e0 / denom; + const float sm1 = e1 / denom; + const float max_sm = fmaxf(sm0, sm1); + return max_sm - SP20_HOLD_BASELINE_UNIFORM_K2; +} + +/* Sum over a per-env trade range of the circular `hold_baseline_buffer`. + * + * Args: + * buf — [N_envs × hold_buffer_size] f32 row-major. + * env — env index (row). + * hold_buffer_size — circular ring size (spec §4.2: 30). + * current_t — current bar (the trade-close bar — its OWN slot + * was NOT written by any per-bar branch because + * segment_complete bypasses them). + * segment_hold_time — float (saved_hold_time captured before + * PS_HOLD_TIME reset). The trade's bar count. + * + * Returns: + * Sum of `buf[env, (current_t - k) % hold_buffer_size]` for + * k = 1..=min(hold_buffer_size, floor(segment_hold_time)). + * Walks backwards from `current_t - 1` (the most recent per-bar write). + * + * Pearl note: kernel is per-env-parallel. Each thread owns exactly one + * env's row; no atomic / cross-thread reads. The modulo arithmetic is + * scalar so the loop is sequential per thread (cheap — at most 30 + * iterations). */ +__device__ __forceinline__ float sp20_sum_hold_baseline_over_trade( + const float* __restrict__ buf, + int env, + int hold_buffer_size, + int current_t, + float segment_hold_time +) { + /* Defensive guard: empty trade or zero buffer ⇒ no contribution. */ + if (hold_buffer_size <= 0) return 0.0f; + if (segment_hold_time <= 0.0f) return 0.0f; + + /* Bound the back-walk. Cast through int for the loop counter; the + * kernel-side caller passes `(int)segment_hold_time` semantics + * already (saved_hold_time is `(float)hold_time` at the trade-close + * branch — non-negative integer values stored as f32). */ + int n_bars = (int)segment_hold_time; + if (n_bars > hold_buffer_size) n_bars = hold_buffer_size; + + /* Modulo arithmetic for negative numerators is implementation- + * defined in C / undefined for very negative values; bring the + * starting index into the positive range first. The expression + * ((current_t - k) % hold_buffer_size + hold_buffer_size) + * % hold_buffer_size + * computes the canonical non-negative modulo result. We pre-add + * hold_buffer_size once before the loop so each iteration just + * decrements. */ + const long long row_off = (long long)env * (long long)hold_buffer_size; + int idx = (current_t - 1) % hold_buffer_size; + if (idx < 0) idx += hold_buffer_size; + + float sum = 0.0f; + for (int k = 0; k < n_bars; ++k) { + sum += buf[row_off + idx]; + --idx; + if (idx < 0) idx += hold_buffer_size; + } + return sum; +} + +#endif /* SP20_HOLD_BASELINE_CUH */ diff --git a/crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu new file mode 100644 index 000000000..11d8d1eae --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu @@ -0,0 +1,134 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP20 Phase 3 Task 3.2 (2026-05-10) — sp20_hold_baseline GPU oracle test wrappers. + * + * Standalone test kernels exercising the two device helpers in + * `sp20_hold_baseline.cuh` so the Rust GPU oracle tests in + * `tests/sp20_hold_baseline_test.rs` can drive synthetic inputs through + * the EXACT same code paths the production caller + * (`experience_env_step` per-bar branches + segment_complete trade-close + * site) executes — no CPU reference impl per + * `feedback_no_cpu_test_fallbacks`. + * + * Test-only: no production callers. Cubins built by + * `crates/ml/build.rs` (kernel names listed below in the order they + * appear in this file). + * + * Two wrappers shipped: + * + * 1. `sp20_per_bar_aux_conf_k2_test_kernel(logits[2], out)`: + * Calls `sp20_compute_per_bar_aux_conf_k2` once. Single-thread + * launch — the helper is pure register-only arithmetic. + * + * 2. `sp20_sum_hold_baseline_over_trade_test_kernel(buf, env, + * hold_buffer_size, + * current_t, + * segment_hold_time, + * out)`: + * Calls `sp20_sum_hold_baseline_over_trade` once. Single-thread + * launch. + * + * 3. `sp20_dual_emission_test_kernel(...)`: + * Mini-fixture that mirrors the production per-bar emission + * site (positioned-non-event branch line ~3650 of + * experience_kernels.cu). Synthetic single-env inputs: + * per_bar_opp_cost = -aux_conf × cost_scale; if action == Hold, + * emit centered = per_bar_opp_cost - hold_reward_ema; always + * write per_bar_opp_cost to buf[0, current_t % size]. Output + * captures BOTH centered + the buffer slot for assertion. This + * wrapper validates the full dual-emission contract on a tight + * synthetic fixture. + * + * Outputs go to mapped-pinned buffers with `__threadfence_system()` + * before thread exit so the host can read via `read_volatile` after a + * stream sync. Same pattern as `sp20_event_reward_test_kernel.cu`. + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#include "sp20_hold_baseline.cuh" + +/* ── Test 1: sp20_compute_per_bar_aux_conf_k2 ─────────────────────────── */ + +extern "C" __global__ void sp20_per_bar_aux_conf_k2_test_kernel( + const float* __restrict__ logits, /* [K=2] */ + float* __restrict__ out_aux_conf) +{ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + out_aux_conf[0] = sp20_compute_per_bar_aux_conf_k2(logits); + __threadfence_system(); +} + +/* ── Test 2: sp20_sum_hold_baseline_over_trade ────────────────────────── */ + +extern "C" __global__ void sp20_sum_hold_baseline_over_trade_test_kernel( + const float* __restrict__ buf, + int env, + int hold_buffer_size, + int current_t, + float segment_hold_time, + float* __restrict__ out_sum) +{ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + out_sum[0] = sp20_sum_hold_baseline_over_trade( + buf, env, hold_buffer_size, current_t, segment_hold_time + ); + __threadfence_system(); +} + +/* ── Test 3: full dual-emission shape ───────────────────────────────────── + * + * Mirrors the production per-bar emission site (positioned-non-event + * branch ~line 3650 of experience_kernels.cu) on a single env. Inputs: + * - logits[2] — aux head logits (drives aux_conf). + * - cost_scale — ISV[HOLD_COST_SCALE_INDEX] (synthetic). + * - hold_reward_ema — ISV[HOLD_REWARD_EMA_INDEX] (synthetic). + * - dir_idx — action direction (DIR_HOLD = 1 fires Path 1). + * - hold_action_id — DIR_HOLD constant (= 1 in production). + * - current_t — bar index for buffer slot. + * - hold_buffer_size — circular ring size (= 30 in production). + * + * Outputs: + * - out_centered[0] — Path 1 result: 0 if not Hold; + * (per_bar_opp_cost - hold_reward_ema) if Hold. + * - out_buf[0] — Path 2 result: per_bar_opp_cost + * (regardless of action) at slot + * [0 × hold_buffer_size + current_t % size]. + * + * Asserts the spec §4.2 dual-emission contract: + * per_bar_opp_cost = -aux_conf × cost_scale + * Path 1 (Hold-only): centered = per_bar_opp_cost - hold_reward_ema + * Path 2 (always): buf[..] = per_bar_opp_cost (UNcentered) + * + * The `out_buf` arg is a [hold_buffer_size] device-resident array + * (env_count = 1 in this fixture). The kernel writes the slot for + * env=0 at `current_t % hold_buffer_size`. + */ +extern "C" __global__ void sp20_dual_emission_test_kernel( + const float* __restrict__ logits, /* [K=2] */ + float cost_scale, + float hold_reward_ema, + int dir_idx, + int hold_action_id, + int current_t, + int hold_buffer_size, + float* __restrict__ out_centered, /* [1] */ + float* __restrict__ out_buf) /* [hold_buffer_size] */ +{ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float aux_conf = sp20_compute_per_bar_aux_conf_k2(logits); + const float per_bar_opp_cost = -aux_conf * cost_scale; + + /* Path 2 — always */ + const int slot_idx = current_t % hold_buffer_size; + out_buf[slot_idx] = per_bar_opp_cost; + + /* Path 1 — Hold-only, centered */ + if (dir_idx == hold_action_id) { + out_centered[0] = per_bar_opp_cost - hold_reward_ema; + } else { + out_centered[0] = 0.0f; + } + + __threadfence_system(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index d04072b41..1cfd22407 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2132,16 +2132,17 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuExperienceCollector.alpha_per_env [alloc_episodes] f32 device-resident scratch — SP20 Phase 2 Task 2.2 (2026-05-09) per-env trade-close alpha (= R_event - hold_baseline) per spec §4.1. Phase 2 emits `hold_baseline = 0.0f` as a Phase 3.2 forward-reference placeholder; Phase 3.2's Hold opportunity-cost dual emission lands the real `hold_baseline_buffer.sum_over_trade_range()` consumer. Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `alpha = sp20_compute_event_reward(segment_return, label_at_open_per_env[env], ISV[LOSS_CAP_INDEX])`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — sums `alpha_per_env[env]` over closed envs and divides by closed-count to populate `EmaInputs.alpha`. FoldReset sentinel 0.0 across all `alloc_episodes` slots — leftover alpha from the previous fold's last trade-close would corrupt the new fold's first ALPHA_EMA observation. The sentinel-0 reads as `alpha == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Replaces the Phase 1.4 hardcoded `out_inputs->alpha = 0.0f` placeholder atomically per `feedback_no_partial_refactor` (Task 2.2 lands buffer + kernel-arg + aggregation-kernel-update + 2 placeholder pin tests in one commit).", }, - // ── SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env - // trade-close win flag scratch — written at `is_close` - // (segment_complete) alongside `alpha_per_env`, read by - // `sp20_aggregate_inputs_kernel` for the `wins_count` - // reduction (replacing the broken per-bar `step_ret > 0` - // predicate that pinned WR_EMA at 0 in production). + // ── SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold + // opportunity-cost circular buffer. Component 2 producer/ + // consumer per spec §4.2. Written at every per-bar + // branch (positioned-non-event + flat-non-event); summed + // at trade-close to compute hold_baseline; replaces the + // Phase 2 Task 2.2 forward-reference placeholder + // `hold_baseline = 0.0f` atomically. RegistryEntry { - name: "is_win_per_env", + name: "hold_baseline_buffer", category: ResetCategory::FoldReset, - description: "GpuExperienceCollector.is_win_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.2-fix (2026-05-10) per-env trade-close win flag (1 iff segment_return > 0 at trade close, 0 otherwise). Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `is_win_per_env[env] = (segment_return > 0.0f) ? 1 : 0` at the same site as `alpha_per_env[env] = alpha`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — counts `is_win_per_env[env] != 0` over closed envs to populate the `wins_count` reduction (drives `EmaInputs.is_win` via the >=0.5 fraction-of-closed threshold). Replaces the broken `step_ret_per_env[env] > 0.0f` per-bar predicate the kernel was using, which conflated the close bar's per-bar mark-to-market tick (dominated by tx-cost so almost always negative at close bars) with trade outcome — pinning WR_EMA at 0 across all training epochs in production despite real wins. FoldReset sentinel 0 across all `alloc_episodes` slots — leftover win flag from the previous fold's last trade-close would corrupt the new fold's first WR_EMA observation. The sentinel-0 reads as `is_win == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Lands atomically per `feedback_no_partial_refactor` with the kernel-body change + buffer alloc + producer wire-up in `experience_kernels.cu::segment_complete` + the new GPU oracle test `wr_ema_uses_segment_pnl_via_is_win_per_env_predicate` (failing-test commit 1, fix commit 2).", + description: "GpuExperienceCollector.hold_baseline_buffer [alloc_episodes × HOLD_BASELINE_BUFFER_SIZE = 30] f32 device-resident scratch — SP20 Phase 3 Task 3.2 (2026-05-10) per-env circular Hold opportunity-cost buffer per spec §4.2 (Component 2). Layout: row-major, `[env, current_t % 30]` slot per per-bar write. Producer: `experience_env_step` per-bar branches (positioned-non-event ~line 3650 + flat-non-event ~line 3737) — both write `hold_baseline_buffer[env, current_t % 30] = -aux_conf × cost_scale` UNCONDITIONALLY (Path 2 counterfactual emission, regardless of action; spec §4.2 line 225 `hold_baseline_buffer[i] = R_per_bar_i // uncentered, for Component 1`). Consumer: `experience_env_step` `is_close && segment_hold_time > 0.0f` branch — `hold_baseline = sp20_sum_hold_baseline_over_trade(buf, env, 30, current_t, segment_hold_time)` walking backwards `min(30, segment_hold_time)` slots from `(current_t - 1) % 30`. Replaces the Phase 2 Task 2.2 placeholder `hold_baseline = 0.0f` atomically per `feedback_no_partial_refactor`. FoldReset sentinel 0.0 across all `alloc_episodes × 30` slots — leftover per-bar costs from the previous fold's open trade would corrupt the new fold's first trade-close summation (the new fold's portfolio_states is also FoldReset so any trade open at fold boundary is impossible, but if a buffer slot retained a stale value it could leak into the first trade-close after fold reset via the modulo-arithmetic indexing). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Per `pearl_first_observation_bootstrap`: cold-start with `ISV[HOLD_COST_SCALE_INDEX=513]` at sentinel 0 means `per_bar_opp_cost = 0` (Path 2 writes zeros to the buffer); the trade-close consumer sums zeros and lands `hold_baseline = 0.0f` (matches Phase 2 Task 2.2 placeholder behaviour bit-for-bit on the very first trade close after fold reset). The controller's first observation lifts cost_scale, then per-bar writes become non-zero, then the next trade-close summation engages — gracefully bootstrapped, no warmup gate needed.", }, ]; Self { entries } @@ -2516,4 +2517,47 @@ mod sp20_registry_tests { entry.description ); } + + /// SP20 Phase 3 Task 3.2: `hold_baseline_buffer` MUST be registered + /// as FoldReset so the new fold's first trade-close summation reads + /// only this fold's per-bar writes — no stale bars leak across + /// folds. Description MUST mention SP20 Phase 3 + Task 3.2 + + /// per-bar branches (producer) + segment_complete (consumer) so + /// wire-up audits have a stable search anchor. + #[test] + fn sp20_hold_baseline_buffer_registered_fold_reset() { + let registry = StateResetRegistry::new(); + let entry = registry + .entries + .iter() + .find(|e| e.name == "hold_baseline_buffer") + .expect("hold_baseline_buffer must be registered in StateResetRegistry"); + assert_eq!( + entry.category, + ResetCategory::FoldReset, + "hold_baseline_buffer must be FoldReset (got {:?})", + entry.category + ); + assert!( + entry.description.contains("SP20 Phase 3"), + "description must reference 'SP20 Phase 3'; got: {}", + entry.description + ); + assert!( + entry.description.contains("Task 3.2"), + "description must reference 'Task 3.2'; got: {}", + entry.description + ); + assert!( + entry.description.contains("per-bar branches"), + "description must reference 'per-bar branches' producer site; got: {}", + entry.description + ); + assert!( + entry.description.contains("segment_complete") + || entry.description.contains("trade-close"), + "description must reference the segment_complete / trade-close consumer; got: {}", + entry.description + ); + } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 0f0dca952..4befd85c0 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -9566,15 +9566,18 @@ impl DQNTrainer { )))?; } } - // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade- - // close win flag scratch. Same device-resident - // `CudaSlice` reset pattern as `alpha_per_env` above. - "is_win_per_env" => { + // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold + // opportunity-cost circular buffer (Component 2). Same + // device-resident `CudaSlice` reset pattern as + // `alpha_per_env` above. Layout `[alloc_episodes × + // HOLD_BASELINE_BUFFER_SIZE = 30]` f32 row-major; reset + // path zeros all alloc_episodes × 30 slots. + "hold_baseline_buffer" => { if let Some(ref mut collector) = self.gpu_experience_collector { let stream = collector.stream().clone(); - stream.memset_zeros(&mut collector.is_win_per_env) + stream.memset_zeros(&mut collector.hold_baseline_buffer) .map_err(|e| crate::MLError::ModelError(format!( - "is_win_per_env memset_zeros: {e}" + "hold_baseline_buffer memset_zeros: {e}" )))?; } } diff --git a/crates/ml/tests/sp20_hold_baseline_test.rs b/crates/ml/tests/sp20_hold_baseline_test.rs new file mode 100644 index 000000000..f3e93c9d2 --- /dev/null +++ b/crates/ml/tests/sp20_hold_baseline_test.rs @@ -0,0 +1,412 @@ +//! SP20 Phase 3 Task 3.2 (2026-05-10) — `sp20_hold_baseline.cuh` GPU oracle tests. +//! +//! Verifies the two device helpers in `sp20_hold_baseline.cuh` +//! (Component 2 / Hold opportunity-cost dual emission per spec §4.2): +//! +//! - `sp20_compute_per_bar_aux_conf_k2(logits)`: +//! Per-row K=2 peak-confidence above the K=2 uniform mean. +//! `aux_conf = max(softmax(logits)) - 0.5` ∈ [0, 0.5]. +//! +//! - `sp20_sum_hold_baseline_over_trade(buf, env, size, current_t, +//! segment_hold_time)`: +//! Sum the per-env circular buffer over the most recent +//! `min(size, segment_hold_time)` slots BEFORE `current_t`. +//! +//! - The full dual-emission shape (mirrors the production per-bar +//! emission site at `experience_env_step` per-bar branches): Path 1 +//! centered Hold reward AND Path 2 uncentered buffer write, with +//! the Hold-only / always action gating contract. +//! +//! Tests assert invariants per +//! `pearl_tests_must_prove_not_lock_observations` (boundedness, +//! formula correctness, fixed points), not observed-from-a-regression- +//! run values. +//! +//! Run on a GPU host: +//! +//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \ +//! cargo test -p ml --test sp20_hold_baseline_test --features cuda \ +//! -- --ignored --nocapture + +#![allow(clippy::tests_outside_test_module)] + +#[cfg(feature = "cuda")] +#[allow(unsafe_code)] +mod gpu { + use std::sync::Arc; + + use cudarc::driver::{CudaContext, CudaFunction, CudaStream, PushKernelArg}; + use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + + const TEST_CUBIN: &[u8] = include_bytes!(concat!( + env!("OUT_DIR"), + "/sp20_hold_baseline_test_kernel.cubin" + )); + + fn make_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_fn(stream: &Arc, name: &str) -> CudaFunction { + let m = stream + .context() + .load_cubin(TEST_CUBIN.to_vec()) + .expect("load sp20_hold_baseline_test_kernel cubin"); + m.load_function(name) + .unwrap_or_else(|e| panic!("load {name}: {e}")) + } + + /// Helper: run `sp20_per_bar_aux_conf_k2_test_kernel` once with the + /// given logits and return the resulting aux_conf. + fn run_per_bar_aux_conf_k2(logits: [f32; 2]) -> f32 { + let stream = make_stream(); + let kernel = load_fn(&stream, "sp20_per_bar_aux_conf_k2_test_kernel"); + + let logits_buf = unsafe { MappedF32Buffer::new(2) }.expect("alloc logits"); + logits_buf.write_from_slice(&logits); + let out_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc out"); + out_buf.write_from_slice(&[0.0_f32]); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&logits_buf.dev_ptr) + .arg(&out_buf.dev_ptr) + .launch(cudarc::driver::LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch sp20_per_bar_aux_conf_k2_test_kernel"); + } + stream + .synchronize() + .expect("sync after sp20_per_bar_aux_conf_k2_test_kernel"); + out_buf.read_all()[0] + } + + /// Helper: run `sp20_sum_hold_baseline_over_trade_test_kernel` once + /// with the given buffer + indexing args and return the resulting sum. + fn run_sum_hold_baseline( + buf: &[f32], + env: i32, + hold_buffer_size: i32, + current_t: i32, + segment_hold_time: f32, + ) -> f32 { + let stream = make_stream(); + let kernel = load_fn( + &stream, + "sp20_sum_hold_baseline_over_trade_test_kernel", + ); + + let buf_dev = unsafe { MappedF32Buffer::new(buf.len()) }.expect("alloc buf"); + buf_dev.write_from_slice(buf); + let out_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc out"); + out_buf.write_from_slice(&[0.0_f32]); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&buf_dev.dev_ptr) + .arg(&env) + .arg(&hold_buffer_size) + .arg(¤t_t) + .arg(&segment_hold_time) + .arg(&out_buf.dev_ptr) + .launch(cudarc::driver::LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch sp20_sum_hold_baseline_over_trade_test_kernel"); + } + stream + .synchronize() + .expect("sync after sum_hold_baseline_over_trade"); + out_buf.read_all()[0] + } + + /// ── Test 1: aux_conf bounds + uniform/peak fixed points ────────── + /// + /// `aux_conf = max(softmax(logits)) - 0.5` ∈ [0, 0.5]. + /// + /// Fixed points: + /// - Equal logits ⇒ peak_softmax = 0.5 ⇒ aux_conf = 0.0 (uniform) + /// - Δlogit → ∞ ⇒ peak_softmax → 1.0 ⇒ aux_conf → 0.5 (saturated) + /// - Δlogit ≈ 0.20067 (= ln(0.55/0.45)) ⇒ aux_conf = 0.05 + /// - Δlogit ≈ 2.19722 (= ln(0.90/0.10)) ⇒ aux_conf = 0.40 + /// + /// These match the spec §4.2 self-stabilizing-property operating + /// points exercised in + /// `sp20_phase1_4_wireup_test::target_hold_pct_behavioral_spot_checks`. + #[test] + #[ignore = "requires GPU"] + fn aux_conf_k2_bounds_and_fixed_points() { + let tol = 1e-4_f32; + + // Uniform logits ⇒ aux_conf = 0.0. + let aux_conf = run_per_bar_aux_conf_k2([1.0, 1.0]); + assert!(aux_conf.abs() < tol, + "uniform logits ⇒ aux_conf = 0.0; got {aux_conf}"); + + // Saturated (Δ = 20 ⇒ peak_softmax ≈ 1.0). + let aux_conf = run_per_bar_aux_conf_k2([0.0, 20.0]); + assert!((aux_conf - 0.5).abs() < tol, + "saturated logits ⇒ aux_conf = 0.5; got {aux_conf}"); + + // Spec-warmup operating point. + let delta_low = (0.55_f32 / 0.45_f32).ln(); + let aux_conf = run_per_bar_aux_conf_k2([0.0, delta_low]); + assert!((aux_conf - 0.05).abs() < 1e-3, + "Δlogit=ln(0.55/0.45) ⇒ aux_conf = 0.05; got {aux_conf}"); + + // Spec-confident operating point. + let delta_high = (0.90_f32 / 0.10_f32).ln(); + let aux_conf = run_per_bar_aux_conf_k2([0.0, delta_high]); + assert!((aux_conf - 0.40).abs() < 1e-3, + "Δlogit=ln(0.90/0.10) ⇒ aux_conf = 0.40; got {aux_conf}"); + + // Symmetry: order-reversed logits give the same aux_conf. + let a1 = run_per_bar_aux_conf_k2([0.0, delta_high]); + let a2 = run_per_bar_aux_conf_k2([delta_high, 0.0]); + assert!((a1 - a2).abs() < tol, + "symmetric logits ⇒ same aux_conf; got {a1} vs {a2}"); + + // Non-stable inputs: large abs values with Δ=0 still uniform. + let aux_conf = run_per_bar_aux_conf_k2([1000.0, 1000.0]); + assert!(aux_conf.abs() < tol, + "large equal logits ⇒ aux_conf = 0.0; got {aux_conf}"); + } + + /// ── Test 2: sum_hold_baseline_over_trade indexing ──────────────── + /// + /// Buffer = [N=1 env × hold_buffer_size = 5] for clarity. Slots + /// pre-loaded with values matching their slot index for easy + /// arithmetic verification. + /// + /// At current_t = 5, segment_hold_time = 3: + /// - Walks backwards from idx (current_t - 1) % 5 = 4. + /// - Reads slots [4, 3, 2] ⇒ sum = 4 + 3 + 2 = 9. + /// + /// At current_t = 5, segment_hold_time = 10 (> hold_buffer_size): + /// - Walks backwards but bounded by hold_buffer_size = 5. + /// - Reads slots [4, 3, 2, 1, 0] ⇒ sum = 0 + 1 + 2 + 3 + 4 = 10. + /// + /// At segment_hold_time = 0: + /// - Returns 0.0f (defensive guard for empty trade). + /// + /// At hold_buffer_size = 0: + /// - Returns 0.0f (defensive guard for empty buffer). + /// + /// Wrap-around: current_t = 7, hold_buffer_size = 5 ⇒ start idx = + /// (7 - 1) % 5 = 1; segment_hold_time = 4 walks 1, 0, 4, 3 (wrap). + /// Sum = 1 + 0 + 4 + 3 = 8. + #[test] + #[ignore = "requires GPU"] + fn sum_hold_baseline_over_trade_indexing() { + let tol = 1e-5_f32; + + // 1 env × 5-slot buffer. Slot[k] = k. + let buf: Vec = (0..5).map(|k| k as f32).collect(); + + // Trade of 3 bars closing at t=5. + let s = run_sum_hold_baseline(&buf, 0, 5, 5, 3.0); + assert!((s - 9.0).abs() < tol, + "current_t=5, segment_hold_time=3 ⇒ sum slots [4,3,2] = 9; got {s}"); + + // Trade of 10 bars (clamped by buffer size). + let s = run_sum_hold_baseline(&buf, 0, 5, 5, 10.0); + assert!((s - 10.0).abs() < tol, + "current_t=5, segment_hold_time=10 ⇒ all 5 slots = 0+1+2+3+4=10; got {s}"); + + // Empty trade. + let s = run_sum_hold_baseline(&buf, 0, 5, 5, 0.0); + assert!(s.abs() < tol, + "segment_hold_time=0 ⇒ defensive guard returns 0; got {s}"); + + // Negative trade (defensive). + let s = run_sum_hold_baseline(&buf, 0, 5, 5, -1.0); + assert!(s.abs() < tol, + "segment_hold_time<0 ⇒ defensive guard returns 0; got {s}"); + + // Empty buffer. + let s = run_sum_hold_baseline(&buf, 0, 0, 5, 3.0); + assert!(s.abs() < tol, + "hold_buffer_size=0 ⇒ defensive guard returns 0; got {s}"); + + // Wrap-around test: current_t=7, size=5, hold_time=4. + // Start idx = (7-1) % 5 = 1; walks 1, 0, 4, 3. + let s = run_sum_hold_baseline(&buf, 0, 5, 7, 4.0); + assert!((s - (1.0 + 0.0 + 4.0 + 3.0)).abs() < tol, + "current_t=7, hold_time=4 ⇒ sum slots [1,0,4,3] = 8; got {s}"); + } + + /// ── Test 3: per-env stride (multiple envs in row-major buffer) ─── + /// + /// Buffer = [N=3 envs × 4-slot]. Slot[env, k] = env * 100 + k. + /// Layout (row-major): + /// env=0: [ 0, 1, 2, 3] + /// env=1: [100, 101, 102, 103] + /// env=2: [200, 201, 202, 203] + /// + /// At env=1, current_t=4, segment_hold_time=2: + /// - Walks backwards from idx (4-1) % 4 = 3. + /// - Reads buf[1*4 + 3] = 103, buf[1*4 + 2] = 102 ⇒ sum = 205. + /// + /// Crucially: the buffer reader MUST use the env stride (`env * + /// hold_buffer_size`) to land in env=1's row, not env=0's. A + /// missing stride would return 1+2=3 (env=0 contents), failing + /// this assertion clearly. + #[test] + #[ignore = "requires GPU"] + fn sum_hold_baseline_per_env_stride() { + let tol = 1e-5_f32; + let n_envs = 3; + let size = 4; + let mut buf = vec![0.0_f32; n_envs * size]; + for env in 0..n_envs { + for k in 0..size { + buf[env * size + k] = (env * 100 + k) as f32; + } + } + + let s = run_sum_hold_baseline(&buf, 1, size as i32, 4, 2.0); + assert!((s - 205.0).abs() < tol, + "env=1, t=4, hold=2 ⇒ buf[1,3]+buf[1,2] = 103+102 = 205; got {s}"); + + let s = run_sum_hold_baseline(&buf, 2, size as i32, 4, 2.0); + assert!((s - 405.0).abs() < tol, + "env=2, t=4, hold=2 ⇒ buf[2,3]+buf[2,2] = 203+202 = 405; got {s}"); + } + + /// ── Test 4: full dual-emission shape (Hold action AND non-Hold) ── + /// + /// Mirrors the production per-bar emission site (positioned-non- + /// event branch ~line 3650 of experience_kernels.cu). + /// + /// Per spec §4.2: + /// per_bar_opp_cost = -aux_conf × cost_scale + /// Path 1 (Hold-only): centered = per_bar_opp_cost - hold_reward_ema + /// Path 2 (always): buf[..] = per_bar_opp_cost (UNcentered) + /// + /// Plan §3.2 reference fixture (line ~1306): + /// Fix aux_conf=0.3, cost_scale=0.1, hold_reward_ema=-0.02 + /// Bar with action=Hold: + /// per_bar_centered = -0.3*0.1 - (-0.02) = -0.03 + 0.02 = -0.01 + /// hold_baseline_buffer = -0.3*0.1 = -0.03 (uncentered) + /// Bar with action=Long: + /// per_bar_centered = 0 + /// hold_baseline_buffer = -0.3*0.1 = -0.03 (always written) + /// + /// To engineer aux_conf=0.3 ⇒ peak_softmax = 0.8 ⇒ Δlogit = + /// ln(0.8/0.2) = ln(4) ≈ 1.3862944. + #[test] + #[ignore = "requires GPU"] + fn dual_emission_hold_vs_non_hold() { + let tol = 1e-4_f32; + const HOLD_BUFFER_SIZE: i32 = 30; + const HOLD_ACTION_ID: i32 = 1; // DIR_HOLD + const LONG_ACTION_ID: i32 = 2; // DIR_LONG + + // Δlogit = ln(0.8/0.2) ⇒ aux_conf = 0.3. + let delta = (0.8_f32 / 0.2_f32).ln(); + let logits = [0.0_f32, delta]; + let cost_scale = 0.1_f32; + let hold_reward_ema = -0.02_f32; + let current_t = 7_i32; + + // Expected: + // per_bar_opp_cost = -0.3 × 0.1 = -0.03 + // Hold: centered = -0.03 - (-0.02) = -0.01 + // buf[7 % 30 = 7] = -0.03 + // Long: centered = 0.0 (Hold-only path) + // buf[7 % 30 = 7] = -0.03 (still written) + let expected_per_bar = -0.03_f32; + let expected_hold_centered = -0.01_f32; + + let stream = make_stream(); + let kernel = load_fn(&stream, "sp20_dual_emission_test_kernel"); + + let logits_buf = unsafe { MappedF32Buffer::new(2) }.expect("alloc logits"); + logits_buf.write_from_slice(&logits); + + let centered_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc centered"); + let buf_buf = unsafe { MappedF32Buffer::new(HOLD_BUFFER_SIZE as usize) } + .expect("alloc hold_baseline_buffer"); + + // ── Sub-test A: action = Hold ───────────────────────────────── + centered_buf.write_from_slice(&[0.0_f32]); + buf_buf.write_from_slice(&vec![0.0_f32; HOLD_BUFFER_SIZE as usize]); + unsafe { + stream + .launch_builder(&kernel) + .arg(&logits_buf.dev_ptr) + .arg(&cost_scale) + .arg(&hold_reward_ema) + .arg(&HOLD_ACTION_ID) + .arg(&HOLD_ACTION_ID) // hold_action_id + .arg(¤t_t) + .arg(&HOLD_BUFFER_SIZE) + .arg(¢ered_buf.dev_ptr) + .arg(&buf_buf.dev_ptr) + .launch(cudarc::driver::LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch dual_emission Hold"); + } + stream.synchronize().expect("sync"); + + let centered = centered_buf.read_all()[0]; + let buf_full = buf_buf.read_all(); + let buf_slot = buf_full[(current_t % HOLD_BUFFER_SIZE) as usize]; + + assert!((centered - expected_hold_centered).abs() < tol, + "Hold-action centered: expected {expected_hold_centered}, got {centered}"); + assert!((buf_slot - expected_per_bar).abs() < tol, + "Hold-action Path 2 buffer slot: expected {expected_per_bar}, got {buf_slot}"); + + // Non-target slots untouched (still 0.0). + for k in 0..(HOLD_BUFFER_SIZE as usize) { + if k as i32 == current_t % HOLD_BUFFER_SIZE { continue; } + assert!(buf_full[k].abs() < tol, + "Hold: buffer slot {k} should be untouched 0.0, got {}", buf_full[k]); + } + + // ── Sub-test B: action = Long ───────────────────────────────── + centered_buf.write_from_slice(&[999.0_f32]); // sentinel != 0 + buf_buf.write_from_slice(&vec![0.0_f32; HOLD_BUFFER_SIZE as usize]); + unsafe { + stream + .launch_builder(&kernel) + .arg(&logits_buf.dev_ptr) + .arg(&cost_scale) + .arg(&hold_reward_ema) + .arg(&LONG_ACTION_ID) + .arg(&HOLD_ACTION_ID) + .arg(¤t_t) + .arg(&HOLD_BUFFER_SIZE) + .arg(¢ered_buf.dev_ptr) + .arg(&buf_buf.dev_ptr) + .launch(cudarc::driver::LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch dual_emission Long"); + } + stream.synchronize().expect("sync"); + + let centered = centered_buf.read_all()[0]; + let buf_full = buf_buf.read_all(); + let buf_slot = buf_full[(current_t % HOLD_BUFFER_SIZE) as usize]; + + assert!(centered.abs() < tol, + "Long-action Path 1 (Hold-only) ⇒ centered = 0; got {centered}"); + assert!((buf_slot - expected_per_bar).abs() < tol, + "Long-action Path 2 ⇒ buffer ALWAYS written; expected {expected_per_bar}, got {buf_slot}"); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6e824b43c..f9f32aac3 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,7 +2,227 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. -## 2026-05-10 — SP20 Phase 2 Task 2.2-fix: post-fix runtime diagnostic infrastructure (read-back + producer test scaffold) +## 2026-05-10 — SP20 Phase 3 Task 3.2: Hold opportunity-cost dual emission (Component 2, atomic) + +Lands the per-bar Hold opportunity-cost dual emission at the +`experience_env_step` per-bar branches (positioned-non-event ~line +3650 + flat-non-event ~line 3737), atomically with the trade-close +consumer wiring at the segment_complete branch (~line 3289 — replaces +the Phase 2 Task 2.2 forward-reference placeholder +`hold_baseline = 0.0f`) per `feedback_no_partial_refactor`. + +### Spec §4.2 dual-emission contract + +``` +Per bar i: + aux_conf_i = max(softmax(aux_logits_i)) - 0.5 # K=2 + cost_scale = ISV[HOLD_COST_SCALE_INDEX = 513] # adaptive [0.01, 0.5] + + per_bar_opp_cost_i = -aux_conf_i × cost_scale + + # Path 1 — real reward, fires only on Hold action: + if dir_idx == DIR_HOLD: + r_micro_or_opp_cost += per_bar_opp_cost_i - ISV[HOLD_REWARD_EMA_INDEX = 516] + (centered for Q-target stability) + + # Path 2 — counterfactual buffer, ALWAYS recorded: + hold_baseline_buffer[i, current_t % 30] = per_bar_opp_cost_i + (uncentered, consumed by Component 1 trade-close) + +At trade close (segment_complete): + hold_baseline = sp20_sum_hold_baseline_over_trade( + hold_baseline_buffer, env=i, hold_buffer_size=30, + current_t, segment_hold_time + ) # walks backwards + # min(30, hold_time) bars + # in env i's row + alpha = R_event - hold_baseline # replaces Phase 2 Task 2.2 + # placeholder = 0.0f +``` + +### Replaced sites + + 1. **Trade-close site** (experience_kernels.cu line ~3289) — + `const float hold_baseline = 0.0f;` Phase 2 Task 2.2 placeholder + replaced with the real `sp20_sum_hold_baseline_over_trade(...)` + call. NULL-tolerant: when `hold_baseline_buffer == NULL` (test + scaffolds), falls back to 0.0f (preserves Phase 2 placeholder + behaviour bit-for-bit). + + 2. **Per-bar positioned-non-event branch** (line ~3650) — SP18 D-leg + `compute_sp18_hold_opportunity_cost(...)` call replaced with the + dual emission writing `r_micro` (rc[3] component) and + `hold_baseline_buffer[i, t % 30]`. + + 3. **Per-bar flat-non-event branch** (line ~3737) — SP18 D-leg call + replaced with the dual emission writing `r_opp_cost` (rc[4] + component) and `hold_baseline_buffer[i, t % 30]`. + +The SP18 D-leg device function (`compute_sp18_hold_opportunity_cost` +at `trade_physics.cuh:655`) is RETAINED — still called by +`sp18_hold_opp_test_kernel.cu` (the SP18 oracle test surface). Per +`feedback_no_stubs`, only the production callers are migrated; the +helper itself stays for the test surface. + +### New device helpers (`sp20_hold_baseline.cuh`) + + - `sp20_compute_per_bar_aux_conf_k2(logits)`: + Per-row K=2 peak-confidence above the K=2 uniform mean. + `aux_conf = max(softmax(logits)) - 0.5` ∈ [0, 0.5]. Numerically + stable via row-max subtraction — bit-identical to the formula in + `sp20_stats_compute_kernel.cu` Pass A. + + - `sp20_sum_hold_baseline_over_trade(buf, env, size, current_t, + segment_hold_time)`: + Sum the per-env circular buffer over the most recent `min(size, + segment_hold_time)` slots BEFORE `current_t`. Walks backwards from + `(current_t - 1) % size` (the most recent per-bar write — the + trade-close bar's segment_complete branch does NOT write to the + buffer). Modulo arithmetic handles wrap-around correctly. + +Both are `__device__ __forceinline__` so the test wrapper kernel can +share them bit-for-bit per `feedback_no_cpu_test_fallbacks`. + +### New buffer + reset infrastructure + + - `GpuExperienceCollector.hold_baseline_buffer: CudaSlice` — + `[alloc_episodes × HOLD_BASELINE_BUFFER_SIZE = 30]` row-major + f32 device-resident scratch. Allocated next to `alpha_per_env`. + + - `pub const HOLD_BASELINE_BUFFER_SIZE: usize = 30` — top-of-file + constant matching spec §4.2 (LOOKAHEAD_HORIZON_MAX). Re-exported + so external tests can reference the same value. + + - Registered in `StateResetRegistry` with `FoldReset` category; + reset path mirrors `alpha_per_env`'s `stream.memset_zeros` async + GPU memset (no host buffer to fill). + + - Dispatch arm in `training_loop.rs::reset_named_state`. + + - Pin-test `sp20_hold_baseline_buffer_registered_fold_reset` — + verifies registry entry + category + description anchors. + +### New experience_env_step kernel args + +Six new args appended to the kernel signature (after the Task 2.2 +`alpha_per_env` / `loss_cap_idx` / `alpha_ema_idx` block): + + 1. `aux_logits_per_env: const float*` — `[N, K=2]` reuses + `exp_aux_nb_logits_buf` (the same buffer the SP20 stats kernel + reads). Stream-implicit producer→consumer ordering. + 2. `hold_baseline_buffer: float*` — `[N × hold_buffer_size]`. + 3. `hold_buffer_size: int` — = 30 (HOLD_BASELINE_BUFFER_SIZE). + 4. `aux_k: int` — = 2 (AUX_NEXT_BAR_K). Reserved for future K>2 + support; the helper currently hardcodes K=2 per spec §4.2 K=2 + AMENDED note. + 5. `hold_cost_scale_idx: int` — ISV slot 513. Mirrors the + `loss_cap_idx` / `alpha_ema_idx` pattern (caller passes slot + index so the kernel stays decoupled from SP14 slot-number drift). + 6. `hold_reward_ema_idx: int` — ISV slot 516. + +All NULL-tolerant: `aux_logits_per_env == NULL` ⇒ both paths skip +(per-bar emission collapses to zero, preserving existing behavior +bit-identically). `hold_baseline_buffer == NULL` ⇒ Path 2 skipped; +Path 1 still fires (centered Hold reward independent of buffer). +Both NULL ⇒ trade-close `hold_baseline` falls back to 0.0f (Phase 2 +placeholder behavior). + +### HOLD_REWARD_EMA forward-reference (concurrent-agent collision avoidance) + +`ISV[HOLD_REWARD_EMA_INDEX]` is updated by `sp20_emas_compute_kernel.cu` +which reads `ema_inputs->per_bar_hold_reward` from +`sp20_aggregate_inputs_kernel.cu`. The aggregate kernel currently +emits `per_bar_hold_reward = 0.0f` (Phase 3.2 forward-reference +placeholder at line 292) — the parallel agent on +`sp20-phase-2-fix` branch is wiring the real producer. Until that +branch lands: + + - `ISV[HOLD_REWARD_EMA] = 0.0` (sentinel). + - Path 1 centered = `per_bar_opp_cost - 0.0` = `per_bar_opp_cost` + (centering is a no-op pre-merge). + +The Task 3.2 contract is **forward-compatible**: the centering math +references the ISV slot (not a hardcoded 0); when the parallel +branch's aggregate-kernel update lands, EMA starts updating, and the +centering becomes load-bearing automatically. No additional Phase 3.2 +work needed post-merge. + +### Plan accuracy errata + +See `docs/superpowers/plans/2026-05-09-sp19-20-wr-first.md` Phase 3 +Errata Gaps 9 (per-env layout), 10 (existing `PS_HOLD_TIME` / +`segment_hold_time` suffices, no new state slot) for the +plan-vs-reality decisions. Spec §4.2 phrasing was implicit +single-thread; reality is per-env-parallel with per-env stride. + +### New device kernel: `sp20_hold_baseline_test_kernel.cu` + +Test wrapper kernels (3 entry points) for the GPU oracle tests: + + - `sp20_per_bar_aux_conf_k2_test_kernel(logits, out)` — drives + `sp20_compute_per_bar_aux_conf_k2` once. + - `sp20_sum_hold_baseline_over_trade_test_kernel(buf, env, size, + current_t, hold_time, out)` — drives the buffer summation helper. + - `sp20_dual_emission_test_kernel(...)` — fixture mirroring the + production per-bar emission site for both Hold-action AND + non-Hold-action paths. + +No production callers; cubin built by `crates/ml/build.rs`. + +### Tests + +`crates/ml/tests/sp20_hold_baseline_test.rs` — 4 GPU oracle tests: + + 1. `aux_conf_k2_bounds_and_fixed_points` — bounds [0, 0.5], + uniform/saturated/spec-warmup/spec-confident/symmetry fixed + points. Six sub-assertions across one test. + 2. `sum_hold_baseline_over_trade_indexing` — basic indexing, + `segment_hold_time > size` clamp, empty-trade defensive guard, + empty-buffer defensive guard, wrap-around modulo. + 3. `sum_hold_baseline_per_env_stride` — multiple envs in row-major + buffer; verifies the env stride lands in the right row + (a missing stride would fail the assertion clearly). + 4. `dual_emission_hold_vs_non_hold` — full dual-emission contract: + Path 1 fires Hold-only AND Path 2 fires always; non-target + buffer slots remain untouched. Mirrors plan §3.2 reference + fixture (aux_conf=0.3, cost_scale=0.1, hold_reward_ema=-0.02). + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh` | NEW | 2 device helpers (Component 2 dual-emission + summation) | +| `crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu` | NEW | 3 test-kernel entry points | +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +6 kernel args, +helper +include, +trade-close consumer wiring, ~70 LoC delete + ~70 LoC insert at 2 per-bar branches | Production wiring (Tasks 3.2 dual emission + trade-close consumer) | +| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +HOLD_BASELINE_BUFFER_SIZE const, +field, +alloc, +6 kernel-arg passes | Per-env buffer infrastructure + production launch wiring | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +entry + invariant test | FoldReset for `hold_baseline_buffer` | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +match arm | `reset_named_state` dispatch | +| `crates/ml/build.rs` | +1 entry | sp20_hold_baseline_test_kernel cubin | +| `crates/ml/tests/sp20_hold_baseline_test.rs` | NEW | 4 GPU oracle tests + 1 helper module | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification + +``` +SQLX_OFFLINE=true cargo build -p ml # cubin compile +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # tests compile +SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 21/21 lib tests +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \ + --test sp20_hold_baseline_test --features cuda \ + -- --ignored --nocapture # 4 GPU oracle tests +``` + +### Phase 3 → Phase 4 forward references + + - Phase 3.4: replay buffer schema — per-bar `aux_conf` field for + the Phase 5 Aux→Q gate (lands separately as the next atomic + commit on this branch). + - Phase 4: n-step credit distributor consumes `R_used = alpha - + alpha_ema` (Component 1 output) — the `alpha` produced HERE in + Phase 3.2 (= `R_event - sum_hold_baseline_over_trade`) becomes + Phase 4's input. + +## 2026-05-10 — SP20 Phase 3 Task 3.3: target_hold_pct behavioral spot-checks (test-only) Adds infrastructure for verifying the SP20 Phase 2 Task 2.2-fix (`is_win_per_env`) producer side at runtime. The consumer-side