gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(sp18 v2 P2.T1-T5): structural Hold opportunity-cost device fn + 3-site migration (INTERIM STATE LIFTED)

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 10:16:49 +02:00
parent 3c318953a7
commit 8da8e2e584
7 changed files with 783 additions and 44 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
crates/ml/build.rs
View File

@@ -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)

205
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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;
}

62
crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu Normal file
View File

@@ -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 <cuda_runtime.h>
#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();
}

82
crates/ml/src/cuda_pipeline/trade_physics.cuh
View File

@@ -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

280
crates/ml/tests/sp18_hold_reward_oracle_tests.rs
View File

@@ -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<CudaStream>) -> 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<CudaStream>,
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}");
}
}

103
docs/dqn-wire-up-audit.md
View File

@@ -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

81
docs/sp18-wireup-audit.md
View File

@@ -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).
<!-- BEGIN audit_sp18_consumers.sh FINGERPRINT -->
```
=== 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