feat(sp20): Phase 3 Task 3.2 — Hold opportunity-cost dual emission (Component 2)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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<f32>, // [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<f32>`
|
||||
/// (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<i32>, // [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<f32>, // [alloc_episodes × 30]
|
||||
|
||||
// Output buffers [alloc_episodes * alloc_timesteps, ...]
|
||||
states_out: CudaSlice<f32>, // #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::<i32>(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::<f32>(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}"
|
||||
|
||||
168
crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh
Normal file
168
crates/ml/src/cuda_pipeline/sp20_hold_baseline.cuh
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
/* 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 */
|
||||
134
crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu
Normal file
134
crates/ml/src/cuda_pipeline/sp20_hold_baseline_test_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
#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();
|
||||
}
|
||||
@@ -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<f32>` (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<i32>` (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<f32>` (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
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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<i32>` 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<f32>` 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}"
|
||||
)))?;
|
||||
}
|
||||
}
|
||||
|
||||
412
crates/ml/tests/sp20_hold_baseline_test.rs
Normal file
412
crates/ml/tests/sp20_hold_baseline_test.rs
Normal file
@@ -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<CudaStream> {
|
||||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||||
ctx.default_stream()
|
||||
}
|
||||
|
||||
fn load_fn(stream: &Arc<CudaStream>, 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<f32> = (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}");
|
||||
}
|
||||
}
|
||||
@@ -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<f32>` —
|
||||
`[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
|
||||
|
||||
Reference in New Issue
Block a user