eaaab152fc4368e856363eea581e291a9704f37a
4877 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
eaaab152fc |
feat(sp20): Phase 2 Task 2.1 — sp20_compute_event_reward + 8 GPU oracle tests
Lands the SP20 Component 1 reward math as a header-only `__device__
__forceinline__` function in a new `sp20_reward.cuh` header, plus the
GPU oracle test scaffold for the 4-quadrant truth table per spec §4.1.
* `sp20_reward.cuh` — `sp20_compute_event_reward(close_pnl,
label_at_open_sign, loss_cap)` function. Returns bounded scalar
`R_event ∈ [loss_cap, +1.0]`. 4-quadrant table:
+1, +1 → +1.0 (right reason, right outcome)
+1, -1 → +0.5 (wrong reason, right outcome)
-1, -1 → -0.5 (right reason, wrong outcome)
-1, +1 → loss_cap (wrong reason, wrong outcome — ISV-driven)
0, * → 0.0 (no information — close_pnl == 0)
Sentinel `label_at_open_sign == 0` (Task 2.0 contract) maps to
dir_match=false ⇒ wrong-reason quadrant. Documented as the safer
default in the header comment.
* `sp20_event_reward_test_kernel.cu` — standalone single-thread test
wrapper, mirrors the `sp12_reward_math_test_kernel.cu` /
`thompson_test_kernel.cu` pattern. Outputs to a mapped-pinned [1]
f32 with `__threadfence_system()` for PCIe-visible coherence.
* `tests/sp20_event_reward_test.rs` — 8 GPU oracle tests
(`#[ignore = "requires GPU"]`-gated):
1. quadrant_right_reason_right_outcome
2. quadrant_wrong_reason_right_outcome
3. quadrant_right_reason_wrong_outcome
4. quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_1
5. quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_2
6. zero_pnl_returns_zero
7. sentinel_label_winning_trade_lands_shoulder (Task 2.0 contract)
8. sentinel_label_losing_trade_lands_loss_cap (Task 2.0 contract)
* `experience_kernels.cu` — `#include "sp20_reward.cuh"` so Task 2.2's
trade-close site replacement has the function in scope.
* `build.rs` — `sp20_event_reward_test_kernel.cu` cubin entry.
No production callers in this commit. Task 2.2 atomically:
- replaces the SP12 v3 reward block at experience_kernels.cu:3216-3380
with the SP20 4-quadrant reward;
- adds the `is_close` consumer of `label_at_open_per_env[i]` (Task
2.0's per-env scratch);
- threads per-env alpha through the SP20 aggregation kernel,
making `alpha_ema` non-zero post-trade-close (replacing the
Phase 1.4 0.0 forward-reference placeholder).
Per `feedback_no_partial_refactor.md` the device function ships
BEFORE the production caller so Task 2.1's GPU oracle tests can
exercise the math in isolation; the partial refactor that breaks
the no-partial-refactor rule is "feature behind the function with
no caller", which this commit's tests resolve in scope.
Per `feedback_no_cpu_test_fallbacks.md` GPU oracle only — no CPU
reference impl. Per `feedback_no_htod_htoh_only_mapped_pinned.md`
all CPU↔GPU buffers are mapped-pinned (`MappedF32Buffer`).
Verification:
SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # clean
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
29a2615f6a |
feat(sp20): Phase 2 Task 2.0 — per-env trade-open label-sign infra
Lands the producer-side infrastructure for SP20 Phase 2's 4-quadrant
reward kernel (Task 2.2):
* New `[alloc_episodes]` `i32` device buffer `label_at_open_per_env`
on `GpuExperienceCollector`, allocated alongside `episode_starts_buf`.
* Two new `experience_env_step` kernel args (`aux_label_per_env` input,
`label_at_open_per_env` output), threaded into the launch site.
* Write site at the `entering_trade` branch maps the upstream aux
next-bar label (`{0,1,-1}` from `aux_sign_label_per_step_kernel`) to
the SP20 spec sign convention (`{-1,+1,0}`) per spec §4.1.
* StateResetRegistry entry `label_at_open_per_env` (FoldReset) +
`reset_named_state` dispatch arm via `stream.memset_zeros`.
* Registry invariant test `sp20_label_at_open_per_env_registered_fold_reset`
pins the FoldReset category + key description anchors.
* Audit-doc entry documents the design rationale (why per-env buffer
vs. derived-at-trade-close read), sign mapping, FoldReset semantics,
kernel-arg threading, and Task 2.1/2.2 forward references.
Producer-only this commit. The consumer (`is_close` branch reading
`label_at_open_per_env[i]` and passing to `sp20_compute_event_reward`)
lands atomically with the SP12 v3 reward block replacement in Task 2.2,
per `feedback_no_partial_refactor`.
NULL-tolerant kernel arg pair lets test scaffolds without aux-head
wiring continue to work; the FoldReset-zeroed buffer reads as sentinel
0 at the consumer (which Task 2.2 maps to "no info" / wrong-reason
quadrant — safer default than over-rewarding a no-signal trade).
Verification:
SQLX_OFFLINE=true cargo check -p ml # clean
SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 18/18 (+1 new)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
abd7e533bc |
fix(architectural): volume_bar_size in cache key + OFI front-month filter
## Two architectural cleanups, both surfaced by the wgdc8 experiment ### Part 1: volume_bar_size in cache key Mirrors the imbalance_bar_threshold/ewma_alpha fix from `f7718b376`. The volume bar size constant (100 contracts/bar) was previously hardcoded and not in the fxcache key. Tuning it would have hit the same fossilization bug as imbalance_bar_threshold did pre-fix. Changes: - `Hyperparams.volume_bar_size: u64` field added (default 100, matches `DEFAULT_VOLUME_BAR_SIZE` for backwards compat). - TrainingProfile loader reads `volume_bar_size` TOML key. - `calculate_dbn_cache_key_full` signature 7 → 8 args. Hashed via `to_le_bytes()`. Test `test_cache_key_includes_volume_bar_size` added; passes alongside the 5 existing tests. - 4 callers updated atomically (per `feedback_no_partial_refactor`): `discover_and_load`, `data_loading.rs:146`, `train_baseline_rl.rs:599`, `precompute_features.rs:259,720`. - `data_loading.rs:279` now passes `self.hyperparams.volume_bar_size` to `build_volume_bars` instead of the hardcoded `DEFAULT_VOLUME_BAR_SIZE`. - New `--volume-bar-size` CLI arg on both binaries (default 100). - New Argo workflow params `volume-bar-size` (default "100") and `data-source` (default "mbp10") on both `train-template.yaml` and `train-multi-seed-template.yaml`. Threaded into precompute + trainer invocations. - `scripts/argo-train.sh` exposes `--volume-bar-size <n>` and `--data-source <s>` for ad-hoc overrides. ### Part 2: OFI front-month filter (latent bug fix) `crates/ml/examples/precompute_features.rs:539-557` (the OFI/VPIN/Kyle's Lambda computation branch when MBP-10 + trades data is available) was loading trades unfiltered for per-bar microstructure feature computation. The volume bar formation path filters front-month per-file (line 354), but the OFI path did not. Effect pre-fix: during contract rollover windows (e.g., ESZ24 → ESH25), OFI per-bar microstructure features included trades from BOTH contracts simultaneously, distorting VPIN, Kyle's Lambda, and trade imbalance signals. Severity in production: small (front-month dominates ES.FUT volume by 10-100×) but real and present in every prior MBP-10+trades production run. Fix: mirror the per-file `filter_front_month` call from the volume bar path. Volume bar formation and OFI computation now both see the same in-month trade tape. Added log line shows raw vs filtered count per file for transparency. ## Why bundled Both fixes touch trade-data plumbing in `precompute_features.rs` and the fxcache key contract. Per `feedback_no_partial_refactor`, related architectural cleanups land atomically. Both surfaced from the same wgdc8 audit; bundling avoids two cache-key-invalidating commits in sequence (each would force full fxcache regen). ## Compatibility - `volume_bar_size` defaults to 100 → existing wgdc7-equivalent runs reproduce, but with a *new* fxcache key (the f7718b376-era cache file is unreachable; harmless, can GC manually). - OFI fix is strictly more correct; no opt-out needed. Existing models trained on contaminated OFI features may show slight feature distribution drift on first cache regen — expected, not a regression. - `data_source = "ohlcv"` Argo param now possible; routes precompute through volume bar branch directly. wgdc8 experiment uses this to test bar resolution sensitivity at volume_bar_size=500 (5× DEFAULT). Tests: 6/6 feature_cache tests pass. Workspace + examples compile clean. Audit-doc: `docs/dqn-wire-up-audit.md` updated. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
f7718b3761 |
fix(architectural): include bar formation params in fxcache key + actually USE imbalance bars
## The bug (audit 2026-05-09) `crates/ml/src/feature_cache.rs:calculate_dbn_cache_key_full` hashed only `(symbol, data_source, dbn_filenames+sizes)` — NOT `imbalance_bar_threshold` or `imbalance_bar_ewma_alpha`. Combined with `precompute_features.rs:346` unconditionally calling `build_volume_bars` regardless of `data_source`, this meant: 1. 14 audited production runs (Apr 11–May 8) all collided on the same fxcache key (`a3f933aa...` / `c07c960a...`) regardless of TOML `imbalance_bar_threshold` value 2. The imbalance-bar code path was reachable only via fxcache MISS, which never happens in production because `ensure-fxcache` always populates first 3. Every "tuning" of `imbalance_bar_threshold` across 16+ SP runs was a silent no-op — the system was actually running volume bars at DEFAULT_VOLUME_BAR_SIZE (100 contracts/bar) ## The fix (this commit) **Part A — cache key includes bar formation params:** - `calculate_dbn_cache_key_full` signature: 5 args → 7 args. Two new f64 params hashed via `to_le_bytes()`. - 4 callers updated atomically (per `feedback_no_partial_refactor`). - 2 new unit tests (`test_cache_key_includes_bar_threshold`, `test_cache_key_includes_bar_alpha`) pin the contract. **Part B — precompute_features actually USES data_source:** - New CLI args `--imbalance-bar-threshold` (default 0.5) and `--imbalance-bar-ewma-alpha` (default 0.1) on both train_baseline_rl and precompute_features. - `precompute_features.rs:346` now branches: when `data_source == "mbp10"` AND `mbp10_data_dir.is_some()`, calls `mbp10_to_imbalance_bars` instead of `build_volume_bars`. **Argo plumbing:** - `train-template.yaml` + `train-multi-seed-template.yaml`: new workflow parameters threaded into BOTH precompute and trainer invocations so both compute the same fxcache key. - `scripts/argo-train.sh`: new CLI flags for ad-hoc overrides. - ensure-fxcache regen path: removed `rm -f /feature-cache/*.fxcache` (with bar-params now in key, parallel experiments coexist). ## Effects going forward - Tuning `imbalance_bar_threshold` actually changes bar density - Configuring `data_source = "mbp10"` actually produces imbalance bars - Multiple parallel experiments at different thresholds coexist on PVC - `dqn-production.toml: imbalance_bar_threshold = 0.5` no longer ignored Default values match prior production behavior → existing wgdc7-equivalent runs reproduce, just with a *new* fxcache key (the old volume-bar cache file is still on disk but won't be hit; harmless, can GC manually). Audit-doc: `docs/dqn-wire-up-audit.md` updated with full context. Tests: 5/5 feature_cache tests pass, full workspace + examples compile clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
1aaf94306c |
fix(wgdc8): bypass EWMA adaptation in ImbalanceBarSampler — use fixed threshold
Without this, the configured imbalance_bar_threshold is washed out within ~5 bars by the adaptive EWMA recursion (T_new = 0.1·T_old + 0.9·observed; with α=0.1, half-life under 1 bar, fixed point = data's natural imbalance scale). The bar-resolution smoke test would have been comparing two identical equilibria, not two different resolutions. Switch mbp10_to_imbalance_bars from `new_with_ewma()` to `new()` so the configured threshold (= 2.5 on this branch, 5× the production 0.5) actually holds for the entire run. Cache-key continuity preserved (ewma_alpha still hashed and logged). Architectural follow-up: 16+ prior SP runs likely had configured-threshold- washout silently. Whether they were all measuring "ES.FUT MBP-10 equilibrium imbalance scale" regardless of the threshold value in their TOML is worth a separate investigation. Out of scope for this branch. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
7b6a2a63f8 |
experiment(wgdc8): 5× imbalance_bar_threshold (0.5 → 2.5) for resolution smoke
Hypothesis test branch off sp19-20-wr-first HEAD
|
||
|
jgrusewski
|
235e838422 |
feat(sp20): Phase 1.4 (Path C) — kernel-arg refactor + aggregation kernel + wire-up
Atomic Phase 1.4 commit per `feedback_no_partial_refactor`. Wires the
SP20 fused-producer chain (Stats → Aggregate → EMAs → Controllers) into
GpuExperienceCollector's per-rollout-step path with one new aggregation
kernel + Phase 1.2 EMA kernel-signature refactor + production wire-up
+ tests + audit/spec/plan amendments, all in one commit.
## Path C decision rationale
The Phase 1.2 EMA kernel originally took 9 scalar value-args. The Phase
1.4 wire-up site (per-rollout-step in GpuExperienceCollector) needs to
feed per-env GPU-resident trade signals (trade_close_per_sample,
step_ret_per_sample, hold_at_exit_per_sample, packed factored
actions_out) into the kernel — forbidden via host sync
(`feedback_cpu_is_read_only`) and via memcpy_dtoh/htod
(`feedback_no_htod_htoh_only_mapped_pinned`). Path C refactors the EMA
kernel to take a device struct pointer (`SP20EmaInputs* ema_inputs`) +
adds an aggregation kernel that writes the struct on the GPU. Phase 1.2
had zero production consumers yet, so the sig change + Phase 1.4 wire-up
land atomically.
## What this commit contains
1. **EMA kernel signature refactor** (Phase 1.2 → Path C):
`sp20_emas_compute_kernel.cu` swaps 9 scalar args for a single
`const SP20EmaInputs* __restrict__ ema_inputs` device pointer.
Math semantics bit-identical. Rust `EmaInputs` mirrored as
`#[repr(C)]` byte-for-byte; new `pack_inputs_into_f32_view` helper
for tests + collectors that fill the struct via a mapped-pinned
f32-aliased buffer.
2. **New aggregation kernel** (`sp20_aggregate_inputs_kernel.cu`):
Per-env arrays (sliced to current rollout step) + sp20_stats outputs
(p50/std) + aux_dir_acc_reduce_kernel output [0] → SP20EmaInputs
struct. Block tree-reduce (4 stripes × bdim × 4 bytes shmem); no
atomicAdd. Aggregation rules per spec §4.5:
- is_close = OR over envs
- is_win = (≥ 0.5 of closed envs were wins)
- trade_duration = round(mean(hold_at_exit) over closed envs)
- action_is_hold = strict majority (count*2 > n_envs) HOLD
- alpha = 0.0 (Phase 2 forward ref — reward kernel)
- per_bar_hold_reward = 0.0 (Phase 3.2 forward ref — Hold-cost dual)
- aux_logits_p50/std/dir_acc forwarded from upstream
3. **Production wire-up in GpuExperienceCollector**:
4 kernel handles + 4 mapped-pinned buffers (struct fields, alloc,
init); per-rollout-step launch sequence after env_step (step 5c):
`Stats → Aggregate → EMAs → Controllers`. All on the same stream;
stream-implicit producer→consumer ordering. Gated on
`isv_signals_dev_ptr != 0 && trainer_params_ptr != 0` (matches the
existing SP14-β EGF chain pattern).
4. **Fold-boundary reset**:
3 new RegistryEntry records (sp20_ema_inputs_buf,
sp20_emas_internal_buf, sp20_emas_obs_count_buf) + 13 new dispatch
arms in training_loop.rs (10 SP20 ISV slot resets + 3 buffer resets).
Closes the pre-existing `every_fold_and_soft_reset_entry_has_dispatch_arm`
regression (was failing on fresh check after the SP20 ISV slot
registrations landed in commit
|
||
|
jgrusewski
|
e96f7fcdd1 |
fix(sp20): replay buffer records magnitude INTENT (Bellman consistency)
Parallel to the 2026-05-08 Class C direction fix but on the **magnitude axis**.
Class C correctly chose REALIZED for direction because env-gated trades (capital
floor / trail / broker cap → Flat) mean no trade happens — recording as Flat is
honest. For magnitude, Kelly cap clamping does NOT gate the trade — the trade
still happens, just at a smaller size. `actual_mag_core` is the bucketed
magnitude derived in trade_physics.cuh:1099-1117 from |position|/max_position;
when the policy intends Full and Kelly clamps to 0.35× max (bucket → Quarter),
recording Quarter loses the policy's intent. Q(s, mag=Full) never receives the
reward gradient from Full-intent trades that actually executed.
One-line swap inside the Class C encoding block at experience_kernels.cu:2689-
2697: `actual_mag_core * b2_size * b3_size` → `mag_idx * b2_size * b3_size`.
`mag_idx` is the local intent magnitude already decoded at line ~2460 from the
original action_idx BEFORE any Kelly/CVaR/Q-gap clamping. Direction axis keeps
Class C semantics (`actual_dir_core` → gated trades record as Flat). Order/
urgency remain intent passthrough.
`actual_mag_core` remains consumed downstream by Task 2.X per-magnitude trade-
lifecycle instrumentation at the seg_mag_bin site below — direction-branch Q
learns from realized; magnitude-branch Q learns from intent. The two axes have
different env-enforcement semantics so they take different actions at the
replay-write site.
Verification: new CPU oracle test crates/ml/tests/sp20_magnitude_intent_record_
test.rs pins the encoding contract with 4 invariants (Full→Quarter scenario,
all-pairs sweep, order/urgency passthrough, direction Class C unchanged). Per
pearl_tests_must_prove_not_lock_observations the test asserts invariants
("recorded mag == intent mag, regardless of realized") not observed values, so
it cannot become a bug-lock if the kernel's compute path changes — only if the
encoding contract itself is broken. All 4 tests pass; cargo check clean.
Audit-doc entry added to docs/dqn-wire-up-audit.md as the SP20 magnitude intent
fix, parallel to the Class C direction fix above with the asymmetry between
the two axes spelled out.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4e21c38b85 |
fixup(sp20): Phase 1.1 K=3 → K=2 retarget (production aux head is K=2)
The Phase 1.1 sp20_stats_compute kernel (
|
||
|
jgrusewski
|
5ded1cb4b9 |
feat(sp20): Phase 1.3 sp20_controllers_compute kernel
Component 5 / Kernel 2 of the SP20 fused-producer chain — the
deterministic derivation step. Reads the EMAs Phase 1.2's
sp20_emas_compute wrote (4 ISV + 4 internal scratch slots) and
derives 6 controller outputs into ISV slots:
- LOSS_CAP (510) = -1 - clamp((wr-0.50)/0.05, 0, 1)
- HOLD_COST_SCALE (513) two-sided ramp around TARGET_HOLD_PCT ±0.05
- TARGET_HOLD_PCT (514) = clamp(0.8 - aux_p50_ema*1.5, 0.1, 0.8)
- N_STEP (517) = clamp(round(trade_dur_ema), 1, 30) (f32)
- AUX_CONF_THRESHOLD (518) = clamp(aux_dir_acc_ema-0.50, 0.01, 0.20)
- AUX_GATE_TEMP (519) = max(aux_conf_std_ema, 0.01)
TARGET_HOLD_PCT computed before HOLD_COST_SCALE because the
HOLD_COST_SCALE controller reads tgt; implemented as a local f32
written to ISV then re-used for the controller (no redundant ISV
read-back). HOLD_COST_SCALE is the only output that READS its own
previous ISV value (in-place update); deadband path writes the
unchanged previous value verbatim per feedback_no_stubs.
Single-block, single-thread kernel — captureable in the per-step
CUDA Graph per pearl_no_host_branches_in_captured_graph. ISV +
internal pointers are mapped-pinned device pointers (existing
buffers from Phase 1.2); kernel emits __threadfence_system() for
PCIe-visible coherence.
Pearls + invariants honoured:
- pearl_controller_anchors_isv_driven: every output's anchor /
target / cap derives from EMAs; only spec-frozen ramp / clamp
parameters are compile-time constants.
- feedback_isv_for_adaptive_bounds: these 6 outputs ARE the
adaptive bounds.
- feedback_no_atomicadd, feedback_no_cpu_compute_strict,
feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs.
- feedback_no_partial_refactor: kernel + launcher + tests + build
entry land atomically; production wire-up lands in Phase 1.4.
- pearl_tests_must_prove_not_lock_observations: 7 GPU oracle
tests assert clamp boundaries, formula correctness, and
bidirectional ramp behavior — NOT specific observed values.
Tests (all #[ignore = "requires GPU"], pass on RTX 3050 Ti sm_86):
1. loss_cap_ramp_boundaries — 4 wr_ema sweeps incl. clamps.
2. n_step_bounds — round + clamp [1, 30].
3. aux_conf_threshold_bounds — clamp [0.01, 0.20].
4. aux_gate_temp_floor — max(std, 0.01).
5. target_hold_pct_inverse_relation — 0.8 − p50*1.5 with clamps.
6. hold_cost_scale_two_sided_ramp — up/down/deadband + clamps.
7. all_six_outputs_written_in_one_launch — guards missed writes.
Plus 3 launcher unit tests (slot range, slot uniqueness, ISV input
slot range).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --features cuda
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \
--test sp20_controllers_compute_test --features cuda \
-- --ignored --nocapture
→ 7/7 GPU oracle tests pass; 3/3 launcher unit tests pass.
Phase 1.4 will land the production caller atomically alongside
Kernel 1 + Kernel 3 launches on the same stream in order
Stats → EMAs → Controllers per the SP20 design.
|
||
|
jgrusewski
|
71aade18cf |
feat(sp20): Phase 1.2 sp20_emas_compute kernel
Component 5 / Kernel 1 of the SP20 design — central state-tracker
that updates 8 Wiener-α EMAs per training step:
- 4 ISV slots: ALPHA_EMA (511), WR_EMA (512),
HOLD_PCT_EMA (515), HOLD_REWARD_EMA (516)
- 4 internal: trade_duration_ema, aux_conf_p50_ema,
aux_conf_std_ema, aux_dir_acc_ema (private
scratch consumed by Phase 1.3 controllers)
Per-EMA i32 observation counters (mapped-pinned obs_count[8])
handle the pearl_first_observation_bootstrap sentinel transition
correctly even when 0.0 is a legitimate observation (e.g.,
first-loss WR=0). count==0 ⇒ replace, count>0 ⇒ Wiener-blend at
α = 0.4 (WIENER_ALPHA_FLOOR per
pearl_wiener_alpha_floor_for_nonstationary).
Phase 1.2 lands kernel + launcher + 5 tests (4 GPU oracle, 1
floor lock) + build entry + audit doc atomically per
feedback_no_partial_refactor. Production wire-up (Phase 1.4)
deferred — kernel is dead code until then.
Verified on RTX 3050 Ti (sm_86):
- 4 GPU oracle tests pass: first_observation_replaces_sentinel,
wiener_alpha_converges_to_long_run_mean,
hold_reward_ema_gated_on_hold_bars,
per_step_emas_fire_unconditionally
- 4 launcher unit tests pass (constants + struct sanity)
- 1 floor-lock test pass (WIENER_ALPHA_FLOOR == 0.4)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
de922c6a4a |
feat(sp20): Phase 1.1 sp20_stats_compute kernel
Component 5 / Kernel 3 of the SP20 fused-producer chain. Single-block BLOCK=256 kernel reads `aux_logits [B, 3]` (the SP14-C aux head's 3-class direction logits) and emits `[aux_conf_p50, aux_conf_std]` into a `MappedF32Buffer<2>`, where the per-row signal is `aux_conf[i] = max_c softmax(logits[i, *])[c] - 1/3`. p50 uses the inlined `sp4_histogram_p99` pattern (per-warp tile binning + cumulative-from-bottom, no atomicAdd per `feedback_no_atomicadd`); std uses two block tree-reductions sharing one shmem tile sequentially. One fused kernel streams `aux_logits` once for both stats per `pearl_fused_per_group_statistics_oracle`. Phase 1.4 wires the production launch site atomically with the rest of the SP20 reward chain per `feedback_no_partial_refactor`. This commit lands kernel + Rust launcher + GPU oracle tests + build entry + audit-doc entry together so the kernel is independently verifiable on RTX 3050 Ti (sm_86) and L40S (sm_89) before the EMA + controller producers (Phase 1.2 + 1.3) reference its outputs. Tests verify: - uniform logits → aux_conf = 0 → [p50, std] = [0, 0] - varied confidence (logit ramp 0 → 3) → matches CPU oracle - heterogeneous half-hot half-uniform → matches CPU oracle - empty batch → degenerate-guard writes [0, 0] All 4 GPU oracle tests + 4 launcher unit tests pass on RTX 3050 Ti. Test data uses per-row variance to avoid the `pearl_sp4_histogram_warp_tile_undercount` lockstep-uniform trap (concentrated values within one bin_width race the per-warp non-atomic increments). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
ef5e745a13 |
feat(sp20): scaffold 7 behavioral test stubs (Phase 0)
All 7 tests marked #[ignore] until corresponding Phase 6 task implements. Tests gate L40S deployment per spec §4.6. - sp20_pure_trend (Phase 6.1) - sp20_pure_noise (Phase 6.2) - sp20_confidence_correlation (Phase 6.3) - sp20_asymmetry_no_game (Phase 6.4) - sp20_n_step_distribution (Phase 6.5) - sp20_per_regime_wr (Phase 6.6) - sp20_regime_transition (Phase 6.7) |
||
|
jgrusewski
|
4249ebc961 |
feat(sp20): register 10 ISV slots in StateResetRegistry
Sentinel = 0.0 per pearl_first_observation_bootstrap. First observation of each EMA replaces the sentinel directly, no blending. Slots [510..520): - loss_cap (510): adaptive loss cap for reward clamp - alpha_ema (511): Wiener-α EMA for loss_cap producer - wr_ema (512): win-rate EMA driving loss_cap adaptive ramp - hold_cost_scale (513): hold penalty cost multiplier - target_hold_pct (514): hold-engagement target - hold_pct_ema (515): hold-engagement EMA - hold_reward_ema (516): hold-action reward EMA - n_step (517): multi-step TD horizon adapter - aux_conf_threshold (518): auxiliary task confidence threshold - aux_gate_temp (519): auxiliary task gating temperature |
||
|
jgrusewski
|
f5eed1fa79 | spec(sp20): reserve 10 ISV slots [510..520) for WR-first reward optimization | ||
|
jgrusewski
|
99332003a4 |
plan(sp19+20): WR-first reward implementation plan
Implements the spec at docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md
(commit
|
||
|
jgrusewski
|
9d9ca3e6e7 |
spec(sp19+20): apply Q1(b) — Hold-reward EMA for Q-scale comparability
Q1 design decision (b): add explicit hold_reward_ema to center the per-bar Hold reward, so Q(Hold) and Q(trade) targets are scale-comparable. The marginal Q(trade) > Q(Hold) preference now comes from data variance in each state (high-aux states pull Q(Hold) more negative), not from structural scale asymmetry that depends on cost_scale magnitude. Q2 decision: keep 4-quadrant fixed (no ramping partials). Already in spec. Changes: - §4.2 Hold opp-cost: dual emission documented — R_per_bar_centered (= R_per_bar - hold_reward_ema) for Q-target/replay tuple, R_per_bar uncentered for hold_baseline_buffer (Component 1 baseline) - §4.5 Kernel 1: hold_reward_ema added (per-step on Hold-state bars only) - §5 data flow: per-bar reward path shows the centered/uncentered split - ISV slots: 9 → 10 (HOLD_REWARD_EMA_INDEX added) - §8 footprint: Component 2 LoC 70 → 90 (+20 for dual emission) - Total LoC estimate: 1620 |
||
|
jgrusewski
|
defbd0abe1 |
spec(sp19+20): patch 7 review issues
P1 (must-fix bugs/gaps): 1. ASYM_RATIO_INDEX → LOSS_CAP_INDEX with explicit formula in §4.1 and §4.5 (was double source-of-truth — formula in §4.1, ISV slot orphaned) 2. Replay buffer schema change (per-bar aux_conf) added to §8 implementation footprint (~50 LoC additional, was invisible in original spec) 3. hold_baseline_buffer size specified = LOOKAHEAD_HORIZON_MAX = 30 bars (§4.2) 4. "4-tier gate" typo → "5-tier gate" in §7 P2 (clarifications): 5. alpha_ema centering documented as load-bearing for cold-start learning (§4.1) — not just for Q-target stability. EV is slightly negative at WR=46% with uninformed SP19 label; advantage-style centering rescues cold-start. 6. Q-scale asymmetry between Hold (uncentered) and trade (alpha_ema centered) documented as INTENTIONAL design choice (§4.2) — produces marginal Q(trade) > Q(Hold) preference that counteracts the Q(Hold) attractor. Behavioral test sp20_pure_noise validates the no-signal Hold default still works. P3 (tightening): 7. Behavioral test thresholds tightened (§4.6): - sp20_pure_trend: WR > 70% → > 90%, PF > 2.0 → > 3.0 - sp20_pure_noise: Hold% > 80% → > 95%, trades < 50 → < 20 |
||
|
jgrusewski
|
730337375f |
spec(sp19+20): WR-first reward + multi-horizon label utilization
Combines SP19 (multi-horizon labels, already landed) + SP20 (WR-first reward) into one spec per pearl_no_deferrals_for_complementary_fixes. Goal: WR ≥ 55%, textbook PF ≥ 2.0, walk-forward stable, per-regime stable. Six components, atomic ship per feedback_no_partial_refactor: 1. Reward kernel (event-driven, 4-quadrant, asymmetric clamp ramped from wr_ema, multi-horizon directional ground-truth check) 2. Hold opportunity-cost (per-bar, dual emission for real reward + Hold baseline buffer) 3. n-step credit distributor (uniform over trade duration, fixes SP18 B-leg self-bootstrap bug at gpu_experience_collector.rs:4154) 4. Aux→Q confidence gate (sigmoid threshold, mean_a Q baseline avoids Hold-everywhere punishment) 5. 3 fused producer kernels (sp20_emas, sp20_controllers, sp20_stats) driving 9 ISV slots in [510..520) 6. 7 behavioral tests on RTX 3050 Ti, gating L40S deployment ~1,550 LoC total. Spec includes data flow, error handling philosophy, 4-tier test gate, success criteria with explicit failure modes. Cross-references: - pearl_event_driven_reward_density_alignment (design principle) - pearl_audit_unboundedness_for_implicit_asymmetry (asymmetric clamp) - pearl_separate_aux_trunk_when_shared_starves (aux gate leverages SP14-C) - project_metric_pipeline_inflation_audit (WR honest, Sharpe honest, goal grounded) - project_goal_wr_55_pf_2 (the goal this spec implements) |
||
|
jgrusewski
|
d39005c6f4 |
feat(sp19 commit b): producer-side multi-horizon reward blend at fxcache write time
Path (B) Commit B — load-bearing semantic change for the SP19 multi-
horizon reward augmentation. At fxcache-write time, blend 1-bar / 5-bar
/ 30-bar log-returns into `tgt[1]` (`preproc_next`) using equal-thirds
weights with `1/sqrt(N)` vol-scale correction:
blended = (1/3) * r_1
+ (1/3) * r_5 / sqrt(5)
+ (1/3) * r_30 / sqrt(30)
The vol-scale correction applies INSIDE the blend (before weighting) so
each horizon's log-return is at unit-volatility-equivalent under
random-walk assumption — drops naturally into the existing `tgt[1]`
preprocessing pipeline (consumed in `experience_kernels.cu:1556` and
`cuda_pipeline/mod.rs:508`) without double-scaling.
Producer trim: valid range shrinks by `LOOKAHEAD_HORIZON_MAX = 30`
bars since `r_30` needs `close[i + WARMUP + 30]`. Walk-forward fold
construction is dataset-relative, so trimming at producer time shifts
every fold's tail by 30 bars — one-time cost per fxcache regen.
Both producer call sites (`precompute_features.rs` for the
feature-precompute binary, `data_loading.rs` for the DBN-fallback
in-process path) change atomically per `feedback_no_partial_refactor`.
Kernel consumers are unchanged — only `tgt[1]`'s value composition
differs from the consumer's perspective.
ISV slots [507..510) (allocated in Commit A) are NOT consumed at
producer time — `precompute_features` runs BEFORE training so the
ISV bus isn't initialised when the blend happens. Producer hardcodes
1/3 each via `SP19_HORIZON_BLEND_WEIGHTS`. Future Path (A) refactor
would bump TARGET_DIM to carry per-horizon log-returns and read
ISV at training-step time; for the empirical hypothesis test
("does multi-horizon blend lift WR?") equal-thirds is sufficient.
fxcache schema invalidation: `FXCACHE_VERSION` 8 → 9. Existing
`.fxcache` files (1-bar-only `preproc_next`) fail validation at load
and trigger Argo's ensure-fxcache regen step. First L40S run after
this commit takes ~10-15 min longer for the regen — one-time cost,
expected.
Touches:
- crates/ml/examples/precompute_features.rs: SP19_HORIZON_BLEND_WEIGHTS
constant + LOOKAHEAD_HORIZON_MAX trim + blend in tgt[] writer +
early-bail-out check on minimum dataset size.
- crates/ml/src/trainers/dqn/data_loading.rs: same blend +
trim in DBN-fallback path; `last sample targets itself` block
removed (the trimmed range guarantees feature-vector and target
lengths match without a sentinel last row).
- crates/ml/src/fxcache.rs: FXCACHE_VERSION 8 → 9 + v9 docstring entry.
- crates/ml/tests/multi_horizon_reward_blend_test.rs (NEW): CPU-only
oracle behavioural test — 4 cases including known returns, trim
contract, zero-close short-circuit, constant-price blend.
- docs/dqn-wire-up-audit.md: Concerns subsection appended to the
SP19 Commit B entry already drafted in Commit A (pre-existing
test_fxcache_empty + test_dqn_checkpoint_round_trip flakiness
documented for Invariant 7).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean
cargo test -p ml --test multi_horizon_reward_blend_test 4/4 pass
cargo test -p ml --test fxcache_roundtrip_test test_fxcache_f32_roundtrip passes (TARGET_DIM unchanged)
cargo test -p ml --lib ... 13 baseline failures (pre-existing); zero new regressions
bash scripts/audit_sp18_consumers.sh --check exit 0
Pre-existing test_fxcache_empty failure: hand-crafts a 64-byte header
but the actual header is 72 bytes; reader bails early on
"failed to fill whole buffer" instead of "bar_count=0". Test setup
bug, NOT a regression. Pre-Commit B failure count: 13. Post-Commit B
failure count: 13 (the test_dqn_checkpoint_round_trip test is flaky
and toggles independently of this change — verified by stashing and
re-running 3× pre-Commit B).
Atomic-refactor invariant satisfied: Commit A (slot reservations) +
Commit B (producer-side blend) land on the same branch with no L40S
dispatch between. Per task instruction the branch stays unpushed
pending user review.
DBN-fallback path: applied identically in `data_loading.rs:497-528`.
Both producers use the same `SP19_HORIZON_BLEND_WEIGHTS` constant and
the same `LOOKAHEAD_HORIZON_MAX = 30` trim.
Vol-scale correction site: applied inside the blend (before weighting)
in BOTH producers. The existing `tgt[1]` preprocessing pipeline does
NOT double-scale — `experience_kernels.cu:1556` and
`cuda_pipeline/mod.rs:508` consume `tgt[1]` as a unit-scale log-return
exactly as before the blend was added; the blended value drops in
without further scaling.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
37964ae2a4 |
feat(sp19 commit a): ISV slot reservations [507..510) for multi-horizon reward blend
Path (B) Commit A — additive infrastructure for the SP19 producer-side
multi-horizon reward augmentation. Three ISV slots reserve indices for
a future Path (A) refactor that would let a controller drive per-batch
horizon weights; producer (Commit B) hardcodes equal-thirds 1/3 each at
fxcache-write time so the slot reservations are dormant in this branch.
Pure additive — changes NO runtime behaviour. The dispatch arms write
the equal-thirds sentinel at every fold boundary, but no kernel
consumes the slots. This is forward-compatible reservation per
feedback_isv_for_adaptive_bounds, NOT a half-fix; the producer wiring
is complete (next commit), the slot reservations are documented
forward-compatibility for the TARGET_DIM-bumping refactor.
Slot allocation:
- 507 REWARD_HORIZON_WEIGHT_1BAR_INDEX sentinel = 1/3
- 508 REWARD_HORIZON_WEIGHT_5BAR_INDEX sentinel = 1/3
- 509 REWARD_HORIZON_WEIGHT_30BAR_INDEX sentinel = 1/3
Touches:
- crates/ml/src/cuda_pipeline/sp14_isv_slots.rs: 3 slot constants +
sentinel + range markers + sp19_reward_horizon_slot_layout_locked +
all_sp19_slots_fit_within_isv_total_dim lock tests.
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs: ISV_TOTAL_DIM 507 →
510 + layout_fingerprint_seed extension (3 SLOT_* lines + a
SP19_PRODUCER_HARDCODED_HORIZON_BLEND token registering producer-time
consumption).
- crates/ml/src/cuda_pipeline/state_layout.cuh: 3 #define mirrors of
the Rust slot indices + SENTINEL_REWARD_HORIZON_WEIGHT_DEFAULT macro.
- crates/ml/src/trainers/dqn/state_reset_registry.rs: 3 RegistryEntry {
FoldReset } with the "SP19 Path (B) reservation" marker + lock-test
expansion 24 → 27 entries.
- crates/ml/src/trainers/dqn/trainer/training_loop.rs: 3 dispatch arms
in reset_named_state writing SENTINEL_REWARD_HORIZON_WEIGHT_DEFAULT.
- docs/dqn-wire-up-audit.md: SP19 Commit A entry documenting the
reservation rationale + verification steps.
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean
cargo test -p ml --lib sp19_reward_horizon_slot_layout_locked passes
cargo test -p ml --lib all_sp19_slots_fit_within_isv_total_dim passes
cargo test -p ml --lib sp18_fold_reset_entries_present passes (27)
cargo test -p ml --lib every_fold_and_soft_reset_entry_has_dispatch_arm passes
bash scripts/audit_sp18_consumers.sh --check exit 0
Atomic-refactor invariant (HARD — feedback_no_partial_refactor): NO
L40S DISPATCH between Commit A and Commit B. The producer-side blend
+ fxcache version bump + behavioural test land in the next commit on
this branch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
ce841eb56e |
fix(audit): lookahead housekeeping — purge gap + per-fold NormStats
Closes recs 2 + 3 in lookahead-bias-audit-2026-04-28.md.
Fix 2A (rec 2): add walk_forward::PURGE_BARS = 5 constant and insert
val_start = train_end + PURGE_BARS in all 4 fold-construction sites:
generate_walk_forward_indices, _from_timestamps, _windows (date-based,
drops first PURGE_BARS bars of val), and gpu_walk_forward::generate_folds
(both stratified variants preserve the gap on boundary shift). 5-bar
width matches max per-bar feature lookback (autocorr lags 1/5/10);
compile-time per feedback_isv_for_adaptive_bounds since the
feature-pipeline lookback is itself compile-time.
Fix 2B (rec 3): per-fold NormStats fit from train-only data.
- walk_forward.rs: add from_features_slice + denormalize/denormalize_batch
helpers (inverse of normalize_batch for clamp-bounded round-trip).
- train_baseline_rl.rs fold loop: load fxcache sidecar norm_stats.json,
denormalise back to RAW, refit per-fold via from_features_slice,
renormalise full dataset, re-upload via init_from_fxcache, save
per-fold norm_stats_fold{N}.json. DBN-fallback path keeps legacy
behaviour (no sidecar to denormalise from) with a warn! log.
Ensemble trainer block is documented follow-up (separate code path).
Behavioral tests:
- test_norm_stats_per_fold_fit_train_only: synthetic train-mean=1.0 /
val-mean=9.0 dataset; from_features_slice recovers train-only mean to
ε=1e-5 while from_features returns global 5.0
- test_norm_stats_denormalize_roundtrip: non-clamped features round-trip
bit-identically
- test_walk_forward_purge_gap_indices_from_timestamps: every fold has
val_start - train_end == PURGE_BARS
- test_fold_generation_basic (gpu_walk_forward): updated to expect the
purge gap
Validation: cargo check --workspace clean (12.30s); 20/20 walk_forward
+ gpu_walk_forward tests pass; audit_sp18_consumers.sh --check exit 0.
Pre-existing test failures on local RTX 3050 Ti are unrelated SIGSEGVs
(VRAM exhaustion in chunked Thompson tests, pre-existing on baseline).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
e140392f86 |
feat(audit): per-regime val WR instrumentation (T/R/V buckets)
Adds 6 output slots to compute_backtest_metrics_kernel — per-bucket
(win_rate, trade_count) for the {Trending, Ranging, Volatile} regime
split — so the val backtest surfaces whether the long-running ~46%
aggregate WR hides regime-conditional edge. Trades are bucketed at
trade-OPEN by feature[40] (ADX-norm) per the structural thresholds
(T:ADX>0.4, R:ADX<0.2, V:otherwise), mirroring
gpu_walk_forward.rs::classify_regime_from_features. Block tree-reduce
only per feedback_no_atomicadd. Observability-only emission via new
HEALTH_DIAG[N]: val_regime [wr_T=... n_T=... wr_R=... n_R=... wr_V=...
n_V=...] line; thresholds remain kernel constants per
feedback_isv_for_adaptive_bounds (no controller consumer yet).
Implementation atomic (kernel + launcher + WindowMetrics + HEALTH_DIAG
emit + 3 GPU oracle tests + audit doc):
- backtest_metrics_kernel.cu: per-thread per-regime trade counters,
2-slot boundary buffer extension carrying open-bar regime through
block stitch, output stride 13 → 19, shmem 5 → 11 reduction tiles
- gpu_backtest_evaluator.rs: WindowMetrics +6 fields, metrics_buf
size 13 → 19, launcher passes features_buf + feature_dim, consume
populates per-regime fields
- metrics.rs: val_regime HEALTH_DIAG line in consume_validation_loss
- regime_wr_oracle_tests.rs (NEW): 1 CPU sanity + 3 GPU oracle tests
(bit-exact match ε=1e-5 vs CPU oracle on stratified 30/40/30 batch)
Validation: cargo check --workspace clean; 17/17 gpu_backtest_evaluator
unit tests pass; 1+3/4 regime WR tests pass on local RTX 3050 Ti
(1.89s); audit_sp18_consumers.sh --check exit 0.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5ee5a1b655 |
chore(sp18): refresh audit fingerprint after Phase 4 docstring update
Off-by-one count drift in the td_lambda_kernel launch-sites bucket (19→20) from a single additional reference in sp18_td_lambda_q_next_oracle_tests.rs's docstring (the docstring now mentions td_lambda_kernel one extra time per the T_SKEL.3 atomic-refactor guard test). Pure documentation update; no source-code change. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
834eaec4bd |
feat(sp18 v2 P4.T1+T2): B-leg next_states hoist + compute_q_next_target_bootstrap skeleton (additive dead code)
Phase 4 lands the additive infrastructure for the B-leg target-Q DDQN
bootstrap that Phase 5 will swap into td_lambda_kernel's q_next argument
at gpu_experience_collector.rs:~4313 (post-Task-4.1 hoist).
Tasks landed:
- **Task 4.1**: build_next_states_f32 invocation hoisted to BEFORE
the TD(λ) launch block. Pure ordering change — verified via the
SP18 audit grep flagging 0 post-TD(λ) consumers of next_states.
- **Task 4.2**: compute_q_next_target_bootstrap skeleton method on
GpuExperienceCollector with full plan-aligned signature
(next_states, target_params_ptr, online_params_ptr, batch_size →
Result<CudaSlice<f32>, MLError>). Body is a hard-error early-return
per feedback_no_stubs — no production caller in Phase 4 (the
caller is wired by Phase 5 Task 5.1, replacing the self-bootstrap
clone). Returning Err satisfies Invariant 9 (no deferred-work
markers); any accidental pre-Phase-5 call hard-errors with a
clear pointer to the open sub-tasks.
Tasks 4.3 (online forward + DDQN per-branch argmax), 4.4 (target
forward + compute_expected_q gather), 4.5 (replace error early-return
with full implementation + GPU oracle test gates) are deferred to a
follow-up subagent dispatch per their plan-defined per-task TDD cycle
scope. Each requires substantial new infrastructure on the collector
(target-net trunk forward chain duplicating gpu_dqn_trainer.rs Pass 2's
VSN+BN+OFI+trunk+branch-head sequence).
Tests:
- crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs (NEW):
3 introspection tests (no GPU required):
* compute_q_next_target_bootstrap_method_exists
* next_states_built_before_td_lambda (Task 4.1 ordering invariant)
* td_lambda_still_consumes_self_bootstrap_q_next_in_phase4
(Phase 4→Phase 5 atomic-refactor guard)
Atomic-refactor invariant (HARD — feedback_no_partial_refactor):
NO L40S DISPATCH between this Phase 4 close-out commit and the
Phase 5 Task 5.1 commit that replaces the q_next clone with
self.compute_q_next_target_bootstrap(&next_states, ...).
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace → clean
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
cargo test -p ml --test sp18_td_lambda_q_next_oracle_tests
→ 3/3 pass
bash scripts/audit_sp18_consumers.sh --check → exit 0
Files:
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs (Task 4.1+4.2)
crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs (NEW; 3 tests)
docs/sp18-wireup-audit.md (Phase 4 close-out section + fingerprint)
docs/dqn-wire-up-audit.md (Invariant 7 entry)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5ea5aa9b8e |
feat(sp18 v2 P3.T1-T5): adaptive HOLD_REWARD_POS/NEG_CAP producer kernel
Lifts the Phase 2 caps from sentinel-driven cold-start (5.0/-10.0
fixed) to p99(|step_ret|) over Long/Short trade closes × 1.5 safety
factor with Wiener-optimal alpha blend. The Phase 2 consumer
(compute_sp18_hold_opportunity_cost) sees producer-driven slots
[483]/[484] from epoch 1 onward; sentinel branch in the consumer
remains as the cold-start fallback path (zero-trade-close epoch ⇒
producer early-returns ⇒ slots stay at sentinel ⇒ consumer falls
through to the +5/-10 macro defaults — bit-identical to pre-Phase-3).
Mirrors SP14 P0-A reward_cap_update_kernel structural template with
three differences: (1) filter (is_close && step_ret != 0) — both
winners AND losers (the consumer needs the magnitude scale of *all*
Long/Short closes); (2) Welford-derived Wiener-α (slots [487..493))
replaces fixed α=0.01, with floor at WELFORD_ALPHA_MIN=0.4 per
pearl_wiener_alpha_floor_for_nonstationary (the policy-realised
distribution is intrinsically non-stationary as the policy adapts);
(3) bounds [0.5, 50.0] (vs. position-side [1.0, 50.0]).
Atomic single-commit per feedback_no_partial_refactor:
- crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu (NEW)
- crates/ml/build.rs cubin manifest entry
- HoldRewardCapUpdateOps in gpu_aux_trunk.rs (new struct + impl)
- HOLD_REWARD_CAP_UPDATE_CUBIN static + struct field +
launch_hold_reward_cap_update method + constructor instantiation +
field-init in gpu_dqn_trainer.rs (5 sites)
- Per-epoch boundary launch in training_loop.rs right AFTER
launch_reward_cap_update (shared step_ret/trade_close source buffers,
independent ISV slot pairs)
- HEALTH_DIAG[N]: hold_reward_cap [pos={:.4} neg={:.4} fire_rate={:.4}]
- 3 GPU oracle tests (T5 producer-drives-slots, Pearl-A REPLACE,
no-closes preserves-isv) — all pass on local RTX 3050 Ti
- Phase 3 close-out sections in docs/sp18-wireup-audit.md and
docs/dqn-wire-up-audit.md
Pearls applied: feedback_no_atomicadd, pearl_first_observation_bootstrap,
pearl_wiener_optimal_adaptive_alpha, pearl_wiener_alpha_floor_for_nonstationary,
pearl_no_host_branches_in_captured_graph, pearl_symmetric_clamp_audit,
pearl_audit_unboundedness_for_implicit_asymmetry (NEG = -2 × POS at
producer time, single source of truth), feedback_isv_for_adaptive_bounds,
pearl_fused_per_group_statistics_oracle.
Validation: cargo check --workspace clean; 3 GPU oracle tests pass on
local RTX 3050 Ti; scripts/audit_sp18_consumers.sh --check exits 0
(no fingerprint drift in tracked sections).
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Phase 4-5 (B-leg target-net forward + q_next replacement) follows.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
1f4cc0f207 |
chore(sp18): refresh audit fingerprint for archaeology cleanup self-references
The cleanup commit's own additions to docs/dqn-wire-up-audit.md (the 2026-05-09 archaeology-cleanup-pass entry) cite HOLD_COST_SCALE_INDEX, hold_cost_scale_update_kernel, and HOLD_COST_INDEX by name in the prose describing what was stripped. The audit script's per-section grep counts those legitimately, producing a +1/+1/+4 hit-count drift on docs/dqn-wire-up-audit.md across the slot 380, slot 461, and hold_cost_scale_update_kernel sections respectively. Regenerate the fingerprint snapshot in docs/sp18-wireup-audit.md so `scripts/audit_sp18_consumers.sh --check` passes against the new post-cleanup baseline. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
e877e8aefe |
chore(sp18): archaeology cleanup — strip post-deletion markers + ISV_TOTAL_DIM giant docstring
Pure-comment follow-up to SP18 Phase 1 atomic deletion ( |
||
|
jgrusewski
|
8da8e2e584 |
feat(sp18 v2 P2.T1-T5): 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` + 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 (
|
||
|
jgrusewski
|
3c318953a7 |
feat(sp18 v2 P1.T3-T6): atomic delete SP13/SP16 hold_cost_scale chain
Deletes the entire reactive Hold-cost controller chain from SP13 P0a / SP16
P2 / SP16 T3 atomically — replaced by the SP18 D-leg structural opportunity-
cost reward landing in Phase 2.
Per DD7(c): observation chain preserved (slots 380, 382). Slot 380 becomes
constructor-init-only diagnostic at HOLD_COST_BASE=0.005, never updated.
18 consumer sites deleted (8 from plan checklist + 10 from A1-A10 audit
expansion):
Production code:
- 3 reward subtraction sites (replaced with Phase 2 placeholder comments)
- hold_cost_scale_update_kernel.cu (file deleted)
- HoldCostScaleUpdateOps struct in gpu_aux_trunk.rs (~120 lines)
- launch_hold_cost_scale_update in gpu_dqn_trainer.rs (forwarder)
- Host controller block in training_loop.rs (~60 lines)
- Slot constants [461..468) in sp14_isv_slots.rs
- HOLD_COST_CONTROLLER_GAIN/FLOOR/CEIL constants in sp13_isv_slots.rs
- state_layout.cuh mirror constants for [461..468)
- build.rs cubin manifest entry
- 7 fold-reset registry entries + 7 dispatch arms
Tests:
- sp14_oracle_tests.rs lines 2173-2926 (11 tests + helpers, ~750 lines)
— these include_bytes! the deleted cubin and cannot survive
- lock_sp18_v2_pp4_retired_chain test (deleted, no inverse-contract
replacement — locking a deletion is pointless ceremony per user call)
Layout fingerprint: range [461..468) is RESERVED gap (SP14-C.1 pattern;
ISV_TOTAL_DIM stays at 507 for checkpoint compatibility). Layout fingerprint
seed updated atomically: HOLD_COST_SCALE + HCS_* slots replaced with
RESERVED_GAP_461_TO_468=SP18_P1_RETIRED marker.
INTERIM STATE: 3 reward sites are now missing the per-bar Hold cost. This
is forbidden to L40S-dispatch until Phase 2 lands the
compute_sp18_hold_opportunity_cost device fn that replaces the deleted
subtraction. The audit doc records this hold.
Validation: cargo check --workspace clean; cargo test -p ml --lib passes
(state_reset_registry::every_fold_and_soft_reset_entry_has_dispatch_arm,
sp16_t3_wiener_welford_slot_layout_locked, sp18_combined_slot_layout_locked,
sp18_shrink_perturb_slot_layout_locked, sp18_fold_reset_entries_present —
all OK). Audit fingerprint: per-bar Hold-cost subtraction sites count
13 → 0 (smoking gun for complete deletion).
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Audit: docs/sp18-wireup-audit.md (full A1-A10 expansion + post-deletion
fingerprint)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
b43413e4d3 |
feat(sp18 v2 P1.T1+T2): pre-dispatch consumer-audit script + pre-commit hook
Phase 1 Tasks 1.1 + 1.2 of docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md.
Per the plan's audit-first design: surface every live reference to the
SP13/SP16 Hold-cost-scale chain (D-leg) and the TD(λ) `q_next = rewards`
self-bootstrap (B-leg) BEFORE atomic deletion. This catches the SP17-style
"missed quantile_q_select consumer" failure mode where a sweeping refactor
left a dangling reference that compiled but broke at runtime.
What lands
- `scripts/audit_sp18_consumers.sh` — 17-section grep that walks every
consumer pattern from the plan's locked checklist (D-leg slot 380,
461, [462..468) HCS_*, hold_cost_scale_update kernel, hold_rate_observer
kernel retained chain, build.rs cubin manifest, state_layout.cuh
mirror constants, state_reset_registry entries, plus B-leg
td_lambda_kernel launch sites, q_next origin, rewards_out consumers,
PER priority sites, replay buffer schema, c51_loss target-Q origin,
target_params_buf consumers, PopArt slot 63 references). Three modes:
default (full grep output), `--fingerprint` (per-section per-file hit
count for diff-able snapshots), `--check` (diff fingerprint vs locked
snapshot in docs/sp18-wireup-audit.md, exit 1 on drift).
- `docs/sp18-wireup-audit.md` — Phase 1 Task 1.1 outcome with two
sections of import:
1. The plan's locked checklist (8 entries the plan author explicitly
identified as deletion targets).
2. The 10 ADDITIONAL CONSUMERS surfaced by the audit (A1-A10) that
the plan-author missed. Per the task input's halt-on-drift
directive, Phase 1 Tasks 1.3-1.5 (atomic deletion) are HALTED
pending human review of the expanded scope. The audit doc is the
spec-amendment record.
3. A locked fingerprint snapshot the pre-commit hook uses for drift
detection on subsequent commits.
B-leg verification confirms B-DD4 (no PER migration) + B-DD1 (target_
params_buf reusable) + the single q_next bootstrap site at
gpu_experience_collector.rs:4143.
- `scripts/pre-commit-hook.sh` — Invariant 7 list extended with the new
audit doc; new `check_sp18_consumer_audit` step runs the audit script
in `--check` mode whenever a commit touches the chain files
(experience_kernels.cu, hold_*_kernel.cu, state_layout.cuh, sp1[3-8]_
isv_slots.rs, gpu_dqn_trainer.rs, gpu_aux_trunk.rs, gpu_experience_
collector.rs, state_reset_registry.rs, training_loop.rs, build.rs,
sp1[3-8]_oracle_tests.rs). Drift triggers a hook failure with a
pointer to the regeneration command. This is a generalisation of
Invariant 7 and addresses Open Q-B from the spec (audit-as-pre-commit
hook for SP-chain consumer drift).
Findings (HALT trigger)
The audit found 10 consumers NOT in the plan's locked 8-site checklist:
A1 — gpu_aux_trunk.rs:1240-1323 HoldCostScaleUpdateOps struct + impl
(84 lines; the plan said launcher was in gpu_dqn_trainer but the
real struct/impl lives in gpu_aux_trunk; trainer just has a thin
forwarding method)
A2 — sp14_oracle_tests.rs:2174-2920 11 GPU oracle tests (~750 lines)
directly exercising the deleted kernel via include_bytes!
(sp16_phase2_hold_cost_scale_climbs_with_overrun + 10 others)
A3 — training_loop.rs:8971-9043 7 dispatch arms in reset_named_state
(slot 461 + HCS_* slots 462-467); contract test
every_fold_and_soft_reset_entry_has_dispatch_arm requires
atomic deletion alongside registry entries
A4 — gpu_dqn_trainer.rs:23200-23216 constructor block writing
HOLD_COST_BASE to slot 380 (RETAINED per DD7c, comment requires
update)
A5 — gpu_dqn_trainer.rs:617, 2245-2256 doc-block prose
A6 — sp13_isv_slots.rs:75-77 HOLD_COST_CONTROLLER_GAIN/FLOOR/CEIL
constants (HOLD_COST_BASE retained for A4)
A7 — gpu_experience_collector.rs:5579-5585 stale comment doc-ref
A8 — sp5_isv_slots.rs:325-327 comment doc-ref
A9 — state_reset_registry.rs:1138-1271 7 already-RETIRED entries
promote to FULL DELETION
A10 — state_reset_registry.rs:2256-2308 lock_sp18_v2_pp4_retired_chain
contract test asserts retired entries STILL EXIST; must be
deleted/rewritten when entries are removed
None of these are architectural surprises — they're straight extensions
of the atomic-deletion scope. But per the task input's halt-on-drift
directive ("If your audit surfaces ANY additional consumer, halt and
report — do NOT proceed with deletion ad-hoc"), Phase 1 Tasks 1.3-1.5
HALT pending human sign-off on the expanded scope.
Path forward (next phase)
Either expand the atomic-deletion commit scope to cover all 18 sites
(8 plan + 10 audit-surfaced) — recommended per `feedback_no_partial_
refactor`; the audit doc serves as the spec-amendment record — or
human-review the audit doc and explicitly approve/reject each A* entry.
The audit doc + script + hook are the pre-condition for either path
and ship together as Tasks 1.1 + 1.2 of Phase 1. Tasks 1.3-1.6 await
human go-ahead.
Branch is at INTERIM STATE: NOT runnable for L40S smoke between this
commit and the post-Phase-1.6 close-out. The interim is purely-additive
(audit script + hook + doc); the actual deletion that creates the
"3 reward sites missing Hold cost" interim from the plan's Task 1.5
has NOT yet been performed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
0b737ec30a |
feat(sp18-v2): ISV-adaptive shrink-and-perturb at fold transition
Replaces the hardcoded `α=0.8 / σ=0.01` constants in `reset_for_fold`'s
`shrink_and_perturb` call with ISV-driven adaptive bounds per
`feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`.
The driving signal is the relative weight drift `||current - best||₂ /
||best||₂` already computed each epoch by the SP18 v2 weight-drift
diagnostic kernel (commit
|
||
|
jgrusewski
|
aab13a83f2 |
feat(sp18-v2): weight-drift HEALTH_DIAG diagnostic
Add per-epoch HEALTH_DIAG line measuring how far current online params
have drifted from the best-Sharpe checkpoint. The fold-transition
restore-best fix preserves peak weights across folds, but within-fold
edge-decay still happens (Adam keeps stepping after peak Sharpe). The
drift diagnostic surfaces the trajectory so operators can correlate
Sharpe-peak decay with parameter movement.
HEALTH_DIAG line:
HEALTH_DIAG[N]: weight_drift [norm_l2={:.6} relative={:.6} branch_max={:.6}]
Where:
- norm_l2 = ||params_flat - best_params||₂ (absolute L2)
- relative = norm_l2 / max(||best_params||₂, EPS) (relative drift)
- branch_max — currently mirrors `relative` (single-scalar form per spec).
Per-branch breakdown is a deferred follow-up: branch heads are
protected from S&P (skip_start..skip_end), so trunk drift dominates
the L2 in any case.
Cold-start: when best_params_snapshot is None (first fold or any fold
without a Sharpe improvement yet), launcher emits [0.0, 0.0] directly
(kernel skipped, mapped-pinned host slice written from CPU).
Distinguishable from "snapshot matches current" via best_epoch elsewhere.
Kernel (weight_drift_diag_kernel.cu): single block × 256 threads. Two
block tree-reductions sharing one shmem tile sequentially:
Pass 1: ||params - best||₂² over (params - best)
Pass 2: ||best||₂² over best
Thread 0 finalizes sqrt + EPS-floored ratio + writes via
__threadfence_system() for PCIe-visible coherence.
Per feedback_no_atomicadd — block tree-reduce only.
Per feedback_no_htod_htoh_only_mapped_pinned — output is
MappedF32Buffer<2>; cold-start bypass uses host_slice_mut.
Wire-up (atomic):
- crates/ml/build.rs: cubin manifest entry
- crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu (NEW)
- crates/ml/src/trainers/dqn/fused_training.rs: field + constructor
+ launch_weight_drift_diag() + read_weight_drift_diag()
- crates/ml/src/trainers/dqn/trainer/mod.rs:
DQNTrainer::read_weight_drift_diag() public wrapper (CPU-only path
emits zeros)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs: per-epoch
HEALTH_DIAG emit adjacent to SP17 dueling block
GPU oracle test (crates/ml/tests/sp18_weight_drift_test.rs, NEW):
4 cases, all #[ignore = "requires GPU"]:
1. weight_drift_matches_cpu_oracle_4096 — N=4096, ε=1e-5
2. weight_drift_is_zero_when_params_equals_best
3. weight_drift_eps_floor_when_best_is_zero (catches NaN/Inf edge)
4. weight_drift_n_zero_emits_zeros (degenerate guard)
All 4 pass on local RTX 3050 Ti.
Pre-commit Invariant 7: docs/dqn-wire-up-audit.md updated with kernel
algorithm, wire-up table, cross-pearl invariants, and oracle-test
inventory.
Per:
- feedback_no_atomicadd (block tree-reduce in drift kernel)
- feedback_no_htod_htoh_only_mapped_pinned (MappedF32Buffer<2>)
- feedback_wire_everything_up (kernel + launcher + emit + test atomic)
- pearl_no_host_branches_in_captured_graph (cold-path, outside graph)
- feedback_no_partial_refactor (single atomic commit)
|
||
|
jgrusewski
|
3e9cefdbd9 |
fix(sp18-v2): restore-best precedes shrink-and-perturb at fold transition
`train_walk_forward` called `reset_for_fold().await?` directly at the fold boundary, which applies `shrink_and_perturb(α=0.8, σ=0.01)` to whatever `params_flat` holds. At end-of-fold, `params_flat` carries the LATEST-EPOCH decayed weights, NOT the best-Sharpe snapshot saved mid-fold. Fold N+1 then started from `0.8 × decayed + noise` and carried the within-fold edge-decay forward. Insert `restore_best_gpu_params()` BEFORE `reset_for_fold` so S&P operates on the best-Sharpe snapshot. Cold-start guard swallows the "no snapshot saved" error on the first fold (or any fold without a Sharpe improvement) — matches pre-fix behavior on cold start, no regression. State interaction with `reset_for_fold`: restore-best writes `params_flat` ONLY (online weights); reset_for_fold then runs S&P on the restored online, hard-syncs target ← restored online, and resets Adam state on every optimizer. Non-conflicting. `best_params_snapshot` is constructor-init `None` on FusedTrainingCtx and is NEVER cleared at fold boundary. The trainer's `self.best_sharpe` IS reset to NEG_INFINITY in `reset_for_fold` so the first improvement in fold N+1 will overwrite the snapshot. Until then, the snapshot holds whichever fold most recently saved a peak — intended training-scoped behavior. Strict within-fold semantics flagged as a follow-up consideration. Behavioral test (CPU oracle) pins the math contract: - shrink_and_perturb(α, σ) on X = α × X + (1−α) × N(0, σ) - with-fix vs without-fix differ by α × (W_best − W_curr) - cold-start (best == current) is bit-identical between orderings GPU-level kernel coverage stays in compile_training_kernels smoke and the L40S 30-epoch validation gating SP18 Phase 1. Pre-commit Invariant 7: docs/dqn-wire-up-audit.md updated with full rationale, state-interaction analysis, lifecycle notes, and the preserved cross-pearl invariants. Per: - feedback_no_partial_refactor (call-ordering + test + audit-doc atomic) - feedback_no_legacy_aliases (reuses existing API unchanged) - feedback_wire_everything_up (restore_best_gpu_params gains second cold-path consumer) - pearl_no_host_branches_in_captured_graph (runs outside graph capture) |
||
|
jgrusewski
|
662bf7eb36 |
plan(sp18 v2): Phase 0 Task 0.5 — pearl candidate draft
Drafts `pearl_diagnostic_decomposition_before_reward_intervention` per
the plan's Task 0.5 deliverable. Captures the discipline pattern
applied in Tasks 0.1–0.2:
Before changing reward weights or Bellman target arithmetic to fight
a Q-attractor, instrument a per-action decomposition of the existing
reward components AND a trajectory observable on the Q-target
distribution. The instrumentation is observability-only and lands
atomically. One epoch of dispatch surfaces whether the suspected
pathology is real BEFORE any production-path code changes.
Pearl is DRAFT — pending the Task 0.3 (reviewer L40S 1-epoch
dispatch) + Task 0.4 (KILL CRITERION evaluation) outcomes. Promotion
to user memory + cross-reference to MEMORY.md is gated on both legs
returning PROCEED.
Cross-references: pearl_canary_input_freshness_launch_order,
pearl_first_observation_bootstrap, pearl_wiener_alpha_floor_for_
nonstationary, feedback_no_partial_refactor, feedback_wire_everything_up.
No production code changes — single doc-only commit.
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
§ Phase 0 Task 0.5.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
27f8e332da |
plan(sp18 v2): Phase 0 Task 0.2 — B-leg V_SHARE trajectory + TD-error magnitude diagnostic
Adds the SP18 v2 Phase 0 B-leg observability scaffold per the plan's
Task 0.2 spec:
- New kernel `td_error_mag_ema_kernel.cu` (single-block × 256 threads):
reads `td_errors_buf [B]` (already populated by C51/MSE loss for PER
priority recomputation, post-train-step), block tree-reduces
`mean(|td_errors[b]|)` (no atomicAdd), and EMA-blends into
`ISV[TD_ERROR_MAG_EMA_INDEX=493]` via Pearl-A first-observation
bootstrap (sentinel 0.0 → REPLACE) + fixed α=`WELFORD_ALPHA_MIN=0.4`
per `pearl_wiener_alpha_floor_for_nonstationary`. The TDB_* Welford
accumulators in slots [498..504) are RESERVED for the Phase 4
q_next_target Wiener-α chain — not used in Phase 0.
- Cubin manifest entry in `crates/ml/build.rs` + `TD_ERROR_MAG_EMA_
CUBIN` re-export in `gpu_dqn_trainer.rs`.
- `sp18_td_error_mag_ema_kernel` field on `GpuDqnTrainer` + cubin load
on the trainer's stream + `launch_sp18_td_error_mag_ema_update()`
cold-path launcher + `read_sp18_td_error_mag_ema()` convenience
wrapper.
- New `sp18_v_share_history: [[f32; 4]; 5]` field on `DQNTrainer` —
fixed-size ring buffer of the last 5 epochs of per-branch V_SHARE
EMA readings (slots [478..482) per SP17 Phase 3.2). Initialised to
`[[NaN; 4]; 5]`; epochs 0–3 emit `nan` as the slope and skip the
ISV write; epoch 4 onward computes `(EMA[now] - EMA[now-4]) / 4`
per branch and writes the dir-branch slope to
`ISV[V_SHARE_TREND_DIAG_INDEX=496]`.
- Two new HEALTH_DIAG lines in `training_loop.rs` at the per-epoch
boundary (right after the SP18 reward_decomp line):
HEALTH_DIAG[N]: v_share_traj [dir_slope=X mag_slope=Y ord_slope=Z urg_slope=W]
HEALTH_DIAG[N]: td_error_pre [magnitude_ema=X]
V_SHARE slope is host-side computation against the ring buffer
(`(now - now_m4) / 4` per branch). TD-error magnitude is post-blend
ISV slot 493 read (producer fires inside `read_sp18_td_error_mag_
ema()`). Pre-fix baseline for the B-DD9 ratio gate
(`avg(|TD-error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`).
- New GPU oracle tests in `crates/ml/tests/sp18_hold_reward_oracle_
tests.rs`:
* `td_error_mag_ema_pearl_a_bootstrap` — synthetic td_errors with
closed-form mean(|td|)=1.125; pre-populate slot at sentinel;
assert post-launch slot equals the mean (Pearl-A direct-replace).
* `td_error_mag_ema_blend_post_bootstrap` — synthetic td_errors
with mean(|td|)=0.5; pre-populate slot at non-sentinel 1.0;
assert blend equals `(1 - 0.4) × 1.0 + 0.4 × 0.5 = 0.8`.
Pure observability — no production-path consumer in this commit. No
reward changes, no Bellman target changes, no kernel modifications to
the action-selection or training paths. Per `feedback_no_partial_
refactor` the kernel + cubin manifest + buffer + launcher + ring
buffer + HEALTH_DIAG emit + GPU oracle tests all land atomically.
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
All 4 GPU oracle tests pass on RTX 3050 Ti (2.09s):
reward_decomp_per_action_gpu_oracle, reward_decomp_empty_bin,
td_error_mag_ema_pearl_a_bootstrap, td_error_mag_ema_blend_post_bootstrap.
Existing slot lock + state_reset_registry tests still pass.
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
§ Phase 0 Task 0.2.
Audit: docs/dqn-wire-up-audit.md § "SP18 v2 Phase 0 Task 0.2".
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
15b50ac38f |
plan(sp18 v2): Phase 0 Task 0.1 — D-leg per-action reward decomposition diagnostic
Adds the SP18 v2 Phase 0 D-leg observability scaffold per the plan's
"Phase 0 — diagnostic emit (NO functional change)" task:
- New kernel `reward_decomp_diag_kernel.cu`: block tree-reduce
(4 blocks × 256 threads, one block per direction-axis bin) reading
`reward_components_per_sample [N×6]` + `actions_out [N]` and emitting
5 per-bin stats (mean r_micro / mean r_opp_cost / mean r_popart /
mean |reward| / fire_rate) into a 20-float row-major output. Bin
order: Hold(0)→Long(1)→Short(2)→Flat(3); col order: micro→opp→
popart→abs→fire. Empty-bin guard emits 0.0 (NOT NaN) per the
consumer-side KILL CRITERION arithmetic contract.
- Cubin manifest entry in `crates/ml/build.rs` + `REWARD_DECOMP_DIAG_
CUBIN` re-export in `gpu_dqn_trainer.rs`.
- 20-float `MappedF32Buffer sp18_reward_decomp_diag_buf` field on
`GpuDqnTrainer` + accessor pair (`sp18_reward_decomp_diag_dev_ptr`
for the writer-side launcher; `read_sp18_reward_decomp_diag` for the
HEALTH_DIAG reader). Buffer is constructor-zeroed so cold-start
HEALTH_DIAG emits a deterministic zero block.
- `sp18_reward_decomp_diag_kernel` field on `GpuExperienceCollector` +
cubin load on the collector's stream + `launch_sp18_reward_decomp_
diag(n, b1, b2, b3, out_dev_ptr)` launcher. Wired in
`training_loop.rs` at the per-step boundary, BEFORE
`launch_reward_component_ema_inplace` (which `memset_zeros` the
source buffer after consuming it) per `pearl_canary_input_freshness_
launch_order`.
- New per-epoch HEALTH_DIAG line emit at the existing per-epoch
boundary (after the SP17 dueling line):
HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z abs=W
fire=F) long(...) short(...) flat(...)]
Reads the mapped-pinned 20-float diag buffer directly via the
collector→trainer host_ptr — no DtoH copy.
- New `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
* `reward_decomp_per_action_cpu_oracle` (CPU oracle pinning the
per-bin reduction math against a 4-sample synthetic batch).
* `reward_decomp_per_action_gpu_oracle` (GPU oracle, ignored unless
`--ignored`; asserts kernel matches CPU oracle bit-for-bit within
1e-6 f32 budget).
* `reward_decomp_empty_bin_emits_zero_not_nan` (empty-bin contract
guard).
Pure observability — no production-path consumer in this commit. No
reward changes, no Bellman target changes, no kernel modifications to
the action-selection or training paths. Per
`feedback_no_partial_refactor` the kernel + cubin manifest + buffer +
launcher + production wire-up + HEALTH_DIAG emit + GPU oracle test all
land atomically.
Verification:
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
CPU oracle test passes; GPU oracle + empty-bin guard both pass on
RTX 3050 Ti (2.13s).
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
§ Phase 0 Task 0.1.
Audit: docs/dqn-wire-up-audit.md § "SP18 v2 Phase 0 Task 0.1".
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
52100ae7a1 |
docs(sp18 v2): PP.5 H100 collection-time profiling — deferred-to-Phase-4 marker
PP.5 is the B-DD6 headroom-check: profile current
gpu_experience_collector::collect_experience end-to-end ms/cycle on
H100 to confirm budget for the new target-net forward pass (B-leg
Phase 4 estimated ~30% slowdown). Plan explicitly marks this as
reviewer-only ("subagent CANNOT run this step (reviewer-only). Hands
off to reviewer at PP.5.").
Pre-Phase subagent dispatch context lacks H100/Argo deploy access, so
the actual nsys profile is deferred to Phase 4 reviewer pickup. NO
production code modified — this is observability infrastructure for
the Phase 4 atomic refactor's regression check.
Pre-Phase exit criteria still satisfied (22 ISV slots locked, registry
entries with dispatch arms, retirement markers, audit doc current);
only the H100 measurement is deferred per the plan's reviewer-only
gating.
Phase 4 pickup instructions documented in audit doc (4-step procedure:
push → argo-train profile → read nsys → fallback rules + human-review
gate).
Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
c66e15f6d0 |
plan(sp18 v2): mark SP16 P2/T3 HOLD_COST_SCALE chain RETIRED (PP.4)
Per DD7=c, the SP18 D-leg structural Hold opportunity-cost reward (compute_sp18_hold_opportunity_cost in trade_physics.cuh + ISV[483..493) cap producer chain) replaces the reactive per-bar Hold-cost chain (SP13 P0a + SP16 P2 + SP16 T3). This Pre-Phase commit is slot-level documentation only — Phase 1 of SP18 deletes the producer kernel + 3 consumer sites + host controller atomically. Retired slot descriptions (sentinel 0.0 preserved for checkpoint compat): 461 sp16_phase2_hold_cost_scale_adaptive 462 sp16_t3_hcs_target_mean (Welford accumulator) 463 sp16_t3_hcs_target_m2 464 sp16_t3_hcs_diff_mean 465 sp16_t3_hcs_diff_m2 466 sp16_t3_hcs_prev_target 467 sp16_t3_hcs_sample_count Preserved per DD7=c "keep the observation chain": 380..383 SP13 P0a hold-pricing observer (per-step writes still active) 460 MIN_HOLD_TEMPERATURE_ADAPTIVE (separate producer chain) 468..474 MHT_* Welford accumulators (MIN_HOLD_TEMPERATURE) Pattern: SP14-C.1 RESERVED-gap precedent — retired slots stay in the layout_fingerprint_seed + state_reset_registry (sentinel 0.0 rewritten each fold) so checkpoints with the old slot names continue to load. Legacy consumers fall back to scale=1.0 (bit-identical pre-Phase-2 behavior). New `sp16_t2_t3_slots_marked_retired` lock test asserts the "SP18 v2: RETIRED" marker is on each retired slot's description, AND that the preserved MHT_* slots are NOT marked retired. Audit doc updated per Invariant 7. Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
2b91f587b7 |
plan(sp18 v2): 22 fold-reset registry entries + dispatch arms (PP.3)
Adds FoldReset registry coverage for all 22 SP18 ISV slots [483..505)
landed in PP.2 (10 D-leg + 12 B-leg), in lockstep with matching match
arms in `reset_named_state` per the existing
`every_fold_and_soft_reset_entry_has_dispatch_arm` contract test (which
catches the "add registry entry, forget dispatch arm → runtime panic
at fold boundary" pattern surfaced twice historically — SP5 #281, SP7
T7 commit
|
||
|
jgrusewski
|
16100912c3 |
plan(sp18 v2): allocate 22 ISV slots [483..505) for D-leg + B-leg
Adds combined SP18 D-leg Hold-reward adaptive-cap slots [483..493) + B-leg TD(λ) Q(s') bootstrap diagnostics + PopArt reset flag at [493..505) per spec DD6 (D-leg 10 slots) + B-DD13 (B-leg 12 slots). D-leg [483..493) — 10 slots: 483 HOLD_REWARD_POS_CAP (sentinel 5.0, bounds [0.5, 50]) 484 HOLD_REWARD_NEG_CAP (sentinel -10.0) 485 HOLD_REWARD_DECOMP_DIAG (sentinel 0.0) 486 HOLD_OPP_COST_FIRE_RATE_EMA (sentinel 0.0) 487..493 HRC_* Welford accumulators (mirrors SP16 T3 HCS_* pattern) B-leg [493..505) — 12 slots: 493 TD_ERROR_MAG_EMA (HEALTH_DIAG B-DD9 ratio gate input) 494 Q_NEXT_TARGET_P99 (target-Q bootstrap bound check) 495 Q_NEXT_MINUS_REWARD_P99 (sanity: should be O(1) post-fix) 496 V_SHARE_TREND_DIAG (B-leg synergy probe) 497 POPART_RESET_FLAG (sentinel 1.0 — one-shot, B-DD11) 498..504 TDB_* Welford accumulators (mirrors HRC_* pattern) 504 RESERVED (B-leg follow-up) Pearl-A first-observation bootstrap sentinels match position-side SP14 P0-A REWARD_POS_CAP_ADAPTIVE pattern (POS=5.0, NEG=-10.0). Slot 497 POPART_RESET_FLAG sentinel = 1.0 per B-DD11 — host writes 1.0 once at first SP18 epoch, kernel zeroes after consuming, gating the per-fold PopArt slot 63 EMA reset at the SP18 deployment boundary. ISV_TOTAL_DIM bumped 483 → 505; layout_fingerprint_seed updated with all 22 new slot names; state_layout.cuh C-side mirror in lockstep (continues SP14-P0A/P1/audit-fix-4A/4B mirror precedent — SP16/SP17 slots intentionally not mirrored per existing pattern, only SP18 gets fresh mirror entries). `feedback_no_partial_refactor`: both legs share an ISV section + a single fingerprint bump; SP13 [380..383) and SP16 [461..474) slots remain ALLOCATED but RETIRED in PP.4 (sentinel 0.0, no producer launch — RESERVED-gap pattern from SP14-C.1 preserves checkpoint compatibility). Audit doc updated per Invariant 7 with Pre-Phase PP.2 entry. Spec: docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
44e1f58b69 |
docs(sp18): copy spec + plan to feat/sp18-combined branch
Mirrors the SP17 PP.1 plan-copy pattern. Plan and spec source-of-truth live on the sister `feat/sp17-dueling-q-network` worktree; this commit copies them onto the SP18 branch so subsequent commits reference local paths. Spec: docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
0c57a5a31f |
docs(sp17-3.3): Phase 3 close-out — canonical HEALTH_DIAG line + pearl
Phase 3.3 finalises the SP17 dueling-Q diagnostic chain. Canonical HEALTH_DIAG line (already emitted by Phase 3.2 commit |
||
|
jgrusewski
|
b6b17d46bb |
feat(sp17-3.2): V_share + advantage_clip_bound producers + extended emit
Phase 3.2 lands the remaining two SP17 dueling-Q diagnostic producers
atomically with kernel + launcher + Rust wrapper + extended HEALTH_DIAG
emit + GPU oracle tests per `feedback_wire_everything_up`.
V_share[d] = |E[V]| / (|E[V]| + |E[A_centered, picked]|)
where picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic — tractable
per-batch without depending on actions_history_buf which is collector-
time state stale relative to the cuBLAS forward at HEALTH_DIAG cadence).
Pearl-A bootstrap (sentinel 0.5) + α=WELFORD_ALPHA_MIN=0.4 + bilateral
[0, 1] clamp per `pearl_symmetric_clamp_audit`. 4 blocks × 256 threads.
advantage_clip_bound = p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5
via sp4_histogram_p99 (block tree-reduce + per-warp tile binning, NO
atomicAdd per `pearl_fused_per_group_statistics_oracle`). EMA α=0.01
slow per-fold + bilateral clamp [0.1, 100.0] per
`pearl_symmetric_clamp_audit`. Pearl-A bootstrap (sentinel 1.0).
Single block × 256 threads + flat |A_centered| scratch buffer
(mapped-pinned, sized to B × Σ_d b_d × NA).
Observability-only — the actual clipping wire-up is Phase 5 follow-up.
The Phase 1 mean-zero contract (commits eabcf8d52..6f53d676f) makes
A_centered a meaningful signal; this commit observes it.
Extended HEALTH_DIAG line:
HEALTH_DIAG[N]: dueling [v_share=(d=X m=Y o=Z u=W)]
[a_var=(d=A m=B o=C u=D)] [clip=K]
GPU oracle tests on RTX 3050 Ti (all pass, 13/13 SP17 tests):
- v_share_per_branch_matches_closed_form: synthetic V=2.0 + linear A
per branch; closed-form V_share = 2/(2 + |K_d × (n_d-1)/2|);
ε=1e-4. Pearl-A bootstrap REPLACES on first launch.
- advantage_clip_bound_tracks_p99_safety: synthetic A with action-
dominant + per-(i,z) jitter (the jitter is REQUIRED — pathologically
lockstep values undercount in sp4_histogram_p99's non-atomic warp
tile binning per the kernel's documented "1/(256×32) loss for
uniformly distributed signals" qualifier; concentrated values violate
the assumption. Real |A_centered| in production is continuous, so
this is a test-data-only effect.) ε=0.20 (jitter + linear histogram
quantization).
Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
Phase 3.2.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
1e70cd5e59 |
feat(sp17-3.1): A_var_ema per-branch producer + HEALTH_DIAG emit
Phase 3 of SP17 dueling-Q identifiability — first of three diagnostic
producers landed atomically with kernel + launcher + Rust wrapper +
HEALTH_DIAG emit + GPU oracle test per `feedback_wire_everything_up`.
Per branch d ∈ {dir, mag, ord, urg}:
Var_d = (1/(B × n_d × NA)) Σ_{i, a, z} (A[i, a, z] − mean_a A[*, z])²
Block tree-reduce (no atomicAdd, `feedback_no_atomicadd`); 4 blocks ×
256 threads. Pearl-A first-observation bootstrap (sentinel 0.0 →
REPLACE on first launch); steady-state α = WELFORD_ALPHA_MIN=0.4 per
`pearl_wiener_alpha_floor_for_nonstationary` — the structural-control
floor preserves catch-up bandwidth without storing 24 Welford
accumulator slots for a cold-path-cadence diagnostic.
Cold-path emit: single launch per HEALTH_DIAG cadence (epoch boundary)
right after `v_a_means`. New line:
HEALTH_DIAG[N]: dueling [a_var=(d=X m=Y o=Z u=W)]
The line will be extended with V_share + advantage_clip_bound in
Phase 3.2, then finalised in Phase 3.3.
GPU oracle test on RTX 3050 Ti: synthetic A constructed so each branch
d has a closed-form Var(A_centered); kernel readback matches expected
value within ε=1e-4 (f32 rounding budget for ~8×n×51 accumulator
length).
Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
Phase 3.1.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
079a06e485 |
test(sp17): asymmetric A → deterministic argmax under MIN_TEMP=0.5
Verifies that when one direction's centered advantage dominates
strongly enough, action selection picks it deterministically across
all N=1000 Philox-keyed runs.
Plan specifies "Thompson temp = 0.0 → pure argmax E[Q]", but the
kernel floors thompson_temp at MIN_TEMP=0.5 per
pearl_blend_formulas_must_have_permanent_floor — pure τ=0 isn't
ISV-accessible. With τ=0.5 the blend is q_eff = 0.5·E[Q] +
0.5·q_sample; deterministic argmax across all τ=0.5 draws requires
the E[Q] gap to exceed atom_span/2.
Construction: NA=3 atoms [-0.1, 0, +0.1] (atom_span=0.2). A_dir
puts +300 at z=2 for Long, -100 at z=2 for others (per-atom mean
zero ⇒ centering preserves shape). Long centered E[Q] ≈ +0.1, others
≈ -0.05 (gap 0.15 > 0.1). q_sample[Long] = +0.1 with prob ≈ 1
(softmax(300) ≈ delta at z=2); q_sample[other] ∈ {-0.1, 0} (zero
prob for +0.1). q_eff[Long] = 0.1; q_eff[other] ≤ -0.025. Long
strictly wins all draws.
If the centering breaks (Long no longer dominant under centered E[Q])
or τ blends a non-Long sample over Long, this test fires.
Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
f31a6b7ff0 |
test(sp17): symmetric A ⇒ identical per-direction E[Q]
Verifies that under uniform A (== 0 across all action × atom slots),
every direction has identical centered E[Q] regardless of V.
The plan's original wording probes this through Thompson selector
("uniform action distribution"), but the kernel's first-wins-strict-
`>` argmax over four i.i.d. Thompson samples produces a structurally
non-uniform distribution under symmetric A even with correct centering
(closed-form earlier-bias predicts ≈[44%, 26%, 18%, 11%] across
Short/Hold/Long/Flat from the tie statistics). The Thompson distribution
is V-dependent through tie statistics — NOT a centering regression.
Restated as the structural pre-Thompson property: with A=0 and
V arbitrary, centered logit = V + 0 is identical across all directions
⇒ per-direction E[Q] identical to ε=1e-5. The Thompson selector reads
these centered logits; if A=0 produced non-zero per-direction E[Q]
spread, *that* would be the centering regression — exactly what this
test catches.
Probed via compute_expected_q (reads back per-action E[Q] directly,
no Thompson noise as red herring). V-non-uniform sanity check confirms
the kernel reads V (non-zero E[Q] when V ≠ 0).
Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
c2dd6917e0 |
test(sp17): A-centered invariance under V shift
Verifies the architectural contract: a uniform additive shift to V across atoms cannot leak into the post-centering distribution. The plan specifies reading A_centered directly, but A_centered is a register-local quantity inside compute_expected_q; this test restates the property as the equivalent behavioral assertion that adding a uniform constant to V leaves every per-action E[Q] identical (softmax translation invariance). A regression that accidentally reduced over (V + A) instead of A alone would shift the per-atom mean by V_SHIFT and corrupt the centered logits; the per-action E[Q] would diverge by O(1), failing the ε=1e-4 assertion. Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
10fafe8e60 |
test(sp17): V-invariance behavioral test
Verifies the architectural contract: V depends only on state, not on the specific advantage tensor. Two A tensors with same per-atom mean produce different post-centering shapes ⇒ different per-action E[Q] (centering is shape-sensitive), but the V-only diagnostic readout is identical across the two runs (V cannot leak in via the per-atom mean reduction). This is the structural property that makes dueling work — without it, the model conflates "state value" with "action value" and gradient updates corrupt V via raw-A noise. Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
6f53d676fb |
docs(sp17): annotate c51_loss/c51_grad SP17 compliance (Commit E)
Audit-only commit per user design call DD9. The plan's Task 1.4
incorrectly claimed c51_loss_batched + c51_grad_kernel needed migration;
both kernels have been mean-zero centered since commit
|