After sp5 merge with -X theirs, two compile errors needed manual fix:
- gpu_her.rs lost the DevicePtrMut trait import
- gpu_tlob.rs Fix 20 regression test used the 4-arg adam_step signature
before sp5's Pearl 4 added β1/β2/ε params.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Conflicts auto-resolved with -X theirs strategy. sp5 has more
evolved MappedF32Buffer API + more recent work; main's via-pinned
cleanup landed in parallel as a lateral migration. Taking sp5's
side preserves the cleaner consolidated upload pattern.
The SP7 controller's CQL reference signal was reading cql_sx (post-SAXPY,
budget-scaled) which created a self-perpetuating deadlock at small
budget values: cql_sx_norm = budget × raw_grad → small budget → small
cql_sx → controller can't update → budget stays small.
The earlier offset 6 → 3 attempt failed because grad_decomp_launch_cql
was never effectively populating slot 3 — the snapshot pattern measures
‖grad_buf − snapshot‖, but apply_cql_gradient writes to cql_grad_scratch
(separate buffer), not grad_buf, so the snapshot delta is always 0.
This commit adds a real producer kernel cql_raw_norm_compute that reads
cql_grad_scratch directly and computes ‖raw_cql‖ over mag/dir/trunk
slices. Wired to fire AFTER apply_cql_gradient and BEFORE
apply_cql_saxpy, populating grad_decomp_result_pinned[3..6] with the
raw norm independent of cql_budget.
SP7 launcher updated to read from offset 3 (now: real raw CQL norm,
not the never-populated cql snapshot delta). HEALTH_DIAG label renamed
cql_sx → cql_raw to reflect the new contract; component index in the
cached grad_component_norms_* arrays switched 2 → 1.
The historical grad_decomp_launch_cql() call is removed — keeping it
would overwrite the slot with 0 after cql_raw_norm_compute fires. The
paired grad_decomp_snapshot_cql snapshot is left in place to scope the
diff to Path A; buffer cleanup (grad_snapshot_cql allocation +
grad_decomp_launch_cql definition) belongs in a follow-up commit per
feedback_no_partial_refactor.
Files: cql_raw_norm_kernel.cu (new, 97 LOC), build.rs,
gpu_dqn_trainer.rs (struct field + cubin static + load + launcher +
SP7 read offset 6→3), loss_balance_controller_kernel.cu (docstring +
arg comment), fused_training.rs (4 launch_cql_raw_norm call sites,
1 dead grad_decomp_launch_cql call removed), training_loop.rs
(HEALTH_DIAG label + index), audit doc Fix 31 SP7 Path A entry.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The SP7 smoke at 237b3dbfb showed all 8 (head, branch) combinations
stuck at bootstrap across 5 epochs, despite Q-variance and gradient
norms being non-zero. The activation flag mechanism is monotonic per
fold (once active, stays active), so 8/8 inactive across all observed
epochs implies the kernel never hits the active path — but we can't
confirm whether that's because (a) grad_decomp_result_pinned reads
zero at SP7's epoch-boundary call site, or (b) something else.
Added two HEALTH_DIAG lines:
- grad_decomp_pinned — reads bytes [0..12,24..36,36..48] of the
pinned buffer the SP7 kernel reads; surfaces iqn/cql_sx/c51 mag/
dir/trunk norms exactly as the kernel sees them.
- lb_active_per_branch — reads LB_{CQL,C51}_ACTIVE_BASE+0..4 from
ISV; surfaces the activation flag state per (head, branch).
Together these disambiguate "grad_decomp not populated at SP7 read
time" from "values populated but kernel cold-start branch fires for
other reasons".
Purely additive — no behavior change. Audit doc Fix 31 sub-bullet.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
SP7 T1 added 24 ISV slots (LB_DIFF_VAR_CQL_BASE=297 ... LB_C51_ACTIVE
through 321) without bumping ISV_TOTAL_DIM, leaving SP7's wiener-state
and activation slots out-of-bounds of the allocated pinned buffer
(294 * 4 = 1176 bytes; SP7 slots write to bytes 1188..1284). GPU
direct-pointer writes silently corrupted memory in the next page;
CPU read_isv_signal_at returned garbage in release builds.
This explains the SP7 smoke's confusing zero-budget-everywhere
observation: the activation flag was reading OOB memory, the consumer
saw inconsistent values, and the controller's actual ISV state was
never visible.
Bumped ISV_TOTAL_DIM to 321 (max valid index 320, covering
SP5_SLOT_END-1). Made pub(crate) so the new contract test can reference
it from sp5_isv_slots.rs. Updated layout_fingerprint_seed()'s slot
entries to include the previously missing D3 and SP7 slots
(TRAINING_SHARPE_EMA=294 through LB_C51_ACTIVE_BASE=317) and updated
ISV_TOTAL_DIM= literal to 321 in lockstep per feedback_no_partial_refactor.
Added contract test all_sp5_slots_fit_within_isv_total_dim to
permanently gate this class of bug at cargo test time. The test would
have caught SP7 T1 instantly; future slot allocations cannot regress
this.
Files: gpu_dqn_trainer.rs (ISV_TOTAL_DIM + comment + fingerprint),
sp5_isv_slots.rs (test), docs/dqn-wire-up-audit.md (Fix 31 sub-bullet).
Cargo check workspace clean. Cargo test ml --lib 935 passed (934
baseline + 1 new contract test); 16 pre-existing GPU-hardware failures
unchanged.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The SP7 controller's flatness gate reads per-branch Q-variance from
ISV[Q_VAR_PER_BRANCH_BASE=222..226) but that signal was not visible in
HEALTH_DIAG. The existing "var_q" in the main HEALTH_DIAG line is
realized step-return variance per magnitude bin from
gpu_experience_collector — a trade-outcome metric, semantically
distinct from per-branch Q-output variance.
Added one emit line immediately after cql_budget_per_branch:
HEALTH_DIAG[E]: q_var_per_branch [dir=X.XXXX mag=X.XXXX ord=X.XXXX urg=X.XXXX]
Reads ISV[222..226) via read_isv_signal_at — the same slots written by
q_branch_stats_kernel.cu (scratch slot 2 per branch) and routed into ISV
by apply_pearls_ad_kernel in launch_sp5_pearl_1_atom. No new ISV slots,
no kernel change, no StateResetRegistry entry.
Class 2 signal (mag_concat_scale / q_rms): Option A infeasible — q_rms
is a per-sample register variable in mag_concat_qdir with no existing ISV
slot; h_s2_rms_ema at ISV[96] is the only available proxy. Option B
(new ISV slot) blocked pending explicit controller OK. See audit doc for
full stop-and-report rationale.
Files touched:
crates/ml/src/trainers/dqn/trainer/training_loop.rs (+28 LOC)
docs/dqn-wire-up-audit.md (+13 LOC)
[ISV slot decision: Option A reused existing Q_VAR_PER_BRANCH_BASE=222..226]
Cargo check workspace clean. State-reset contract test passes.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Forward W_Q SGEMM stored col-major [K, M] (lda=K) while backward dW_Q
wrote col-major [M, K] (ldc=M). When M ≠ K (TLOB: M=16, K=32), Adam's
element-wise update applied gradients computed at position (m, k) to
weights stored at position (k, m) — silent learning corruption at
every flat index ≠ 0 (511 of 512 W_Q slots updated using wrong-position
gradients, matched in W_K/V; W_O is square so unaffected).
Standardised backward dW_Q/K/V SGEMM output to col-major [K, M] (ldc=K)
matching the forward layout (Strategy A from the audit brainstorm — the
forward layout is the definitive weight storage; Adam's flat layout
follows forward's allocation). The fix flips the cuBLAS strided-batched
operands: backward now computes `dW^T = ofi @ d_proj^T` instead of
`dW = d_proj @ ofi^T`. Same gradient values, just re-laid-out so flat
indexing matches `params`. No new kernel; no kernel-internal layout
change (the SDP forward/backward kernels still read `proj_qkv_buf` /
`d_proj_qkv_buf` as [M, B] col-major — those buffers are untouched).
The QKV-fusion `cublasSgemmStridedBatched(batch=3)` semantics are
preserved: ofi is the new shared operand (strideA=0), d_proj is the
per-batch operand (strideB=M·B), strideC=M·K=512 unchanged.
Phase-1 reproduction (`tlob_dw_layout_alignment_repro`,
#[ignore = "requires GPU"]) ran the broken and fixed cuBLAS dispatches
side-by-side on identical sentinel inputs (`d_proj[m=0,b=0]=1`,
`ofi[k=1,b=0]=1`, all else 0); broken `[M, K]` placed the `1.0`
gradient at flat 16, fixed `[K, M]` placed it at flat 1 — O(1)
cross-layout delta exactly matching the audit prediction. Pre-fix Adam
would have updated `W_Q[m=0, k=16]` (the forward layout's flat-16 slot)
using the gradient computed for `W_Q[m=0, k=1]` — the silent
corruption.
Phase-3 regression (`tlob_dw_layout_alignment_regression_full_chain`,
#[ignore = "requires GPU"]) exercises the full forward → backward →
Adam → forward chain with random Xavier-init weights (W_O seeded to
break the production-zero-init that would collapse the gradient chain
to all-zero in a synthetic test). Asserts (1) GPU dW_Q matches a CPU
reference computed in the post-fix [K, M] layout within TF32 tolerance,
and (2) the second forward Q matches the analytical [K, M]
interpretation of the post-Adam W_Q — locks in cross-step layout
agreement and would fail if any future refactor accidentally
re-permutes `params` between Adam and the next forward.
Existing inline `tlob_sgemm_parity_with_cpu_reference` still passes
(its CPU dW_Q/K/V reference was updated in lockstep to the [K, M]
layout per `feedback_no_partial_refactor`; pre-fix the GPU produced
[M, K] and the new CPU reference would diverge element-wise — a clean
no-skip parity check that locks the layout convention end-to-end).
`tlob_qkv_fusion_equivalence` unchanged (the fix only touches the
backward call, forward QKV fusion is bit-identical pre/post).
Local verification (RTX 3050 Ti, batch=256 for fusion test):
tlob_dw_layout_alignment_repro: PASS
tlob_dw_layout_alignment_regression_full_chain: PASS
tlob_qkv_fusion_equivalence: PASS (3.79× speedup retained)
tlob_sgemm_parity_with_cpu_reference: PASS
Fix 20 in docs/dqn-gpu-hot-path-audit.md updated FIXED with verdict
+ strategy + test list. Forward SGEMM call site got an inline comment
block documenting the [K, M] convention and pointing at the
`tlob_dw_layout_alignment_*` regression coverage.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ZN.FUT tests in crates/ml/src/data_loader.rs were failing because no
valid ZN.FUT DBN data is available locally; gated with #[ignore].
ml-asset-selection's universe definition and backtesting's
zn_futures() slippage profile remain untouched — those are production
references to ZN as a candidate symbol, distinct from data availability.
Migrated 4 deprecated cudarc memcpy_stod/memcpy_dtov sites in the
test function test_eval_action_select_eval_argmax_picks_best in
crates/ml/src/cuda_pipeline/mod.rs to mapped-pinned per
feedback_no_htod_htoh_only_mapped_pinned:
- 3x memcpy_stod (f32 input uploads) → MappedF32Buffer::new +
write_from_slice + dev_ptr as raw u64 kernel arg; kernel reads
directly from mapped-pinned pages, no DtoD copy needed
- 1x memcpy_dtov (i32 output readback) → MappedI32Buffer::new +
dev_ptr as kernel arg + read_all() after stream sync
The cudarc deprecation suggested clone_htod/clone_dtoh as replacements
but those still perform HtoD/DtoH copies — violating the strict rule.
Mapped-pinned with direct dev_ptr kernel args is the correct pattern
(matches distributional_q_tests.rs).
Note: DqnGpuData/PpoGpuData upload paths also in mod.rs still use
clone_to_device_f32_via_pinned; migrating those requires changing
CudaSlice<f32> struct fields to MappedF32Buffer which is blocked until
gpu_dqn_trainer.rs consumers are also updated (separate scope).
Workspace cargo check warnings: 15 → 15 (test-only deprecated calls
not visible to cargo check; ZN gate adds 3 to ignored count).
cargo test -p ml --lib failures: 16 → 13 (3 ZN tests now ignored).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Fixes the SP7 controller dormancy discovered in smoke-test-8556k:
per-branch budgets stuck at exactly bootstrap constants because the
kernel's cold_start_basis numerically equaled the consumer's bootstrap
fallback, AND the Wiener-α was clamped at ALPHA_FLOOR=1e-4 from step 1.
Architectural change:
- 8 new ISV slots LB_{CQL,C51}_ACTIVE_BASE per (head × branch).
Activation flag is monotonic per fold, FoldReset on boundary.
- Kernel only sets active=1 when grads are populated AND on subsequent
active-state computation; cold-start branch leaves active=0 (fresh
branch) or holds prior budget steady (transient grad gate).
- Consumer dispatches on activation: bootstrap when active<0.5,
controller verbatim when active>=0.5. No more spurious bootstrap
when controller writes legitimate small values.
- Welford-α hybrid (max of 1/max(1,epoch_idx_in_fold) and Wiener-α)
gives full update on first active step, falls off as 1/N until
Wiener takes over with meaningful variance estimates. EPOCH_IDX_INDEX
is the existing per-fold-reset counter (no new tuned constants).
State reset registry: 2 new sp7_lb_*_active FoldReset entries +
matching dispatch arms in reset_named_state. Contract test
(every_fold_and_soft_reset_entry_has_dispatch_arm) gates compile.
GPU unit test sp7_loss_balance_controller_activation_flag_transitions
exercises 3 transitions (cold start → both flags 0; active → both
flags 1 with controller-computed budget != bootstrap; transient grad-
gate → flags hold at 1, prior budget held verbatim). Passes on local
RTX 3050 Ti.
Audit doc: Fix 31 sub-bullet describing the activation-flag fix.
Memory pearl out-of-tree (controller will dispatch separately).
Touched:
crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
crates/ml/src/trainers/dqn/fused_training.rs
crates/ml/src/trainers/dqn/state_reset_registry.rs
crates/ml/src/trainers/dqn/trainer/training_loop.rs
crates/ml/tests/sp5_producer_unit_tests.rs
docs/dqn-wire-up-audit.md
Cargo check workspace clean. Cargo test ml --lib clean (incl. contract
test + 6 sp5_isv_slots tests). 16 pre-existing failures unchanged.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces three back-to-back `cublasSgemm_v2` calls (one per Q/K/V
projection, M=16 K=32 N=B at TF32) with a single
`cublasSgemmStridedBatched(batch=3)` launch in both the forward and the
dW_Q/K/V backward paths. Cuts cuBLAS heuristic-lookup + kernel-launch
overhead 3× on the TLOB hotspot identified by task #218 nsys profiling.
Strategy chosen: strided batched (Strategy 2 from the worktree brief),
NOT the originally-recommended concatenated-W approach. Reason: the
concat-W path requires the SDP kernel to read with stride-3M
(col-major [3M, B], ldc=3M), forcing a kernel signature change and
breaking bit-equivalence with the prior 3-SGEMM path. Strided batched
keeps the per-projection [M, B] memory layout intact, so the SDP
forward + backward kernels are byte-identical pre/post fusion (only the
buffer layout is fused: 3 contiguous M·B-float chunks at offsets
0, M·B, 2·M·B inside `proj_qkv_buf` and `d_proj_qkv_buf`).
Param flat layout `[W_Q | W_K | W_V | W_O]` is unchanged
(strideA=M·K reads the existing weights in order), so the
checkpoint/save/load contract is unaffected (TLOB has no on-disk
checkpoint; weights are Xavier-init random).
Numerical equivalence + microbenchmark (RTX 3050 Ti, batch=256, TF32):
- max abs diff Q=3.77e-4, K=3.41e-4, V=4.29e-4
→ within 2e-3 TF32 tolerance (matches inline parity test's TOL_GEMM)
- per-call latency over 200 iters:
fused 1× SgemmStridedBatched batch=3: 5–6 µs
ref 3× cublasSgemm_v2 back-to-back: 19–22 µs
→ ~3.5–3.8× speedup on the QKV-projection portion alone (forward;
backward dW_Q/K/V fusion has the same shape and the same gain).
Tolerance rationale documented inline (`TOL_FUSION = 2e-3`): the shared
classic-cuBLAS handle is bound to `CUBLAS_TF32_TENSOR_OP_MATH`
(`shared_cublas_handle::create_handles_and_workspace`); the
strided-batched dispatch can pick a different internal algo than
back-to-back single calls and the K=32 reduction amplifies TF32 rounding
to a few × 1e-4. Both paths are mathematically equivalent within TF32
precision; sub-1e-5 bit-equivalence is not achievable on a TF32 handle
and is not what the fusion is supposed to provide. Layout/stride/offset
bugs would show up as O(1) deltas, which the 2e-3 threshold catches
trivially.
Tests:
- `cuda_pipeline::gpu_tlob::tests::tlob_sgemm_parity_with_cpu_reference`
(existing inline parity vs CPU SGEMM reference): still PASSES — the
fused path produces the same Q/output/dW values to within 2e-3 of the
hand-rolled CPU reference.
- `cuda_pipeline::gpu_tlob::tests::tlob_qkv_fusion_equivalence`
(NEW, `#[ignore = "requires GPU"]`): runs both the new fused path and
a private 3-SGEMM reference helper on identical inputs, asserts max
abs diff ≤ TOL_FUSION, and prints a fused-vs-3-call latency
microbenchmark over 200 iters. Reverts the fusion if it ever stops
helping.
Audit doc updated: `docs/dqn-gpu-hot-path-audit.md` Fix 20 records the
strategy, bench numbers, and a pre-existing forward/backward
W_Q-vs-dW_Q lda/ldc transposition observation surfaced during
analysis (orthogonal to QKV fusion; flagged for a separate audit
pass — the fusion preserves the existing per-projection layouts
byte-for-byte).
via_pinned migration (overlap with `wt/via-pinned-cleanup`):
The repo's pre-commit `check_no_dtod_via_pinned` guard rejects ANY
staged .rs file containing `upload_f32_via_pinned` or
`clone_to_device_*_via_pinned`. Three pre-existing call sites in
gpu_tlob.rs (line ~235 production param upload + 2 inline-test
uploads) plus one new site I added in the equivalence test would have
blocked this commit. Per the worktree brief I was instructed to leave
the existing line ~235 alone for the parallel `wt/via-pinned-cleanup`
worktree (commit 072c1d3f9), but the hook applies to the whole file
content not the diff, so a partial migration is not viable: I migrated
all 4 call sites in gpu_tlob.rs to the canonical
`MappedF32Buffer + memcpy_dtod_async + sync` pattern that
072c1d3f9 already applies to every other crate-ml caller.
The shape of the migration is identical to 072c1d3f9, so when the
controller merges both worktrees back to main the gpu_tlob.rs hunks
should resolve to the same final content (or a trivial whitespace
merge); no additional functional reconciliation is needed.
Constraints respected:
- `feedback_no_partial_refactor`: kernel sig preserved (offset device
pointers); param + grad buffer layouts unchanged on disk and in
memory; no stale call sites left behind.
- `feedback_no_cpu_compute_strict`: fused dispatch is GPU-only
(cublasSgemmStridedBatched).
- `feedback_isv_for_adaptive_bounds`: no new tunable constants —
QKV_BATCH=3 and W_QKV_STRIDE_FLOATS=M·K are structural.
- `feedback_trust_code_not_docs`: docstrings (`Architecture`,
`Backward`, `cuBLAS API choice`, forward/backward step comments,
buffer field docs) all updated.
- `feedback_no_htod_htoh_only_mapped_pinned`: all CPU↔GPU uploads in
the file now go through `MappedF32Buffer` direct staging (host_ptr
writes, kernel/cublas reads dev_ptr) — zero `via_pinned` calls in
the file after this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The *_via_pinned helpers (clone_to_device_f32_via_pinned,
upload_f32_via_pinned, clone_to_device_i32_via_pinned,
upload_i32_via_pinned, upload_u32_via_pinned) were temporary scaffolding
during the SP4 mapped-pinned migration. The canonical replacement is
MappedF32Buffer / MappedI32Buffer / MappedU32Buffer in the same module —
direct allocation + write_from_slice + device_ptr with no additional
indirection.
Each caller now inlines the staging + DtoD pattern directly:
- Allocate MappedXxxBuffer (cuMemHostAlloc DEVICEMAP)
- Write data via write_from_slice (no memcpy, mapped pinned coherence)
- alloc_zeros::<T> CudaSlice for device-resident destination
- memcpy_dtod_async from staging.dev_ptr to dst
- stream.synchronize() before staging drops
This commit migrates all callers atomically (per
feedback_no_partial_refactor) and removes the now-unused helpers in the
same commit. grep -rn 'via_pinned' returns 0 matches after this lands.
The two local file-private helpers in gpu_experience_collector.rs
(upload_host_to_cuda_f32_via_pinned / upload_host_to_cuda_i32_via_pinned)
were already correctly using MappedF32Buffer directly; they were renamed
to drop the _via_pinned suffix.
Touched: gpu_attention.rs, gpu_backtest_evaluator.rs, gpu_dqn_trainer.rs,
gpu_experience_collector.rs, gpu_her.rs, gpu_iql_trainer.rs,
gpu_iqn_head.rs, gpu_ppo_collector.rs, gpu_tlob.rs, gpu_walk_forward.rs,
gpu_weights.rs, mapped_pinned.rs, mod.rs, hyperopt/adapters/mamba2.rs,
hyperopt/adapters/ppo.rs, trainers/dqn/trainer/training_loop.rs,
trainers/ppo.rs, docs/dqn-wire-up-audit.md (18 files, +930/-363 LOC).
Cargo check workspace clean. Cargo test ml --lib: 925 passed, 17 failed
(all 17 are pre-existing GPU/data infrastructure failures unrelated to
this change).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Surfaced the SP7 T7 dispatch-arm bug at runtime (unknown-name panic at
fold boundary). This test enumerates RegistryEntry names from
StateResetRegistry::new() and asserts each has a match arm in
reset_named_state, catching the bug at `cargo test -p ml --lib`
instead of mid-training.
Source-introspection design — no production-code change. Test fails
fast if a future contributor adds a registry entry without the
corresponding dispatch arm.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The pearl_loss_balance_controller memory pearl is at
~/.claude/projects/.../memory/pearl_loss_balance_controller.md (not in
git — user-memory subsystem). MEMORY.md index entry updated.
This audit-doc commit records the T8 landing for the Fix 31 chain.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Pre-existing test bug: line 164 asserted `ef >= 0.05` (post-Kelly-cap
eval_dist), which gates on Kelly cold-start warmup completing within
the smoke's training horizon — not on whether the Q-head learned to
prefer Full magnitude. On the local-laptop smoke (1 quarter MBP-10),
Kelly warmup never completes, pinning eval_dist[Full] = 0 even when
Q(Full) > Q(Half) clearly.
Per `project_magnitude_eval_collapse_kelly_capped`, the diagnostic
split landed in #212: intent_dist measures policy learning, eval_dist
measures policy + Kelly-cap. Tests asserting on Q-learning success
must use intent_dist; the test was never updated.
Smoke now PASSES with intent_full=0.057 (intent_full=0.866 in the
prior re-run — both above the 0.05 threshold; Q(Full) is clearly
preferred when Kelly cap doesn't suppress it).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
T2 (commit aa2854017) registered 4 SP7 ResetEntries:
sp7_lb_diff_var_cql → ISV[297..301)
sp7_lb_sample_var_cql → ISV[301..305)
sp7_lb_diff_var_c51 → ISV[305..309)
sp7_lb_sample_var_c51 → ISV[309..313)
But never added their dispatch arms in `reset_named_state`. At fold
boundary, the dispatcher panics with "unknown name 'sp7_lb_diff_var_cql'"
because every FoldReset entry must have a matching match arm.
Same shape as bug #281 (SP5 Layer A bug-fix). Surfaced by the SP7 T7
local sanity smoke (test_magnitude_distribution).
The 4 new arms mirror the existing `sp5_budget_cql` / `sp5_budget_c51`
template — `for b in 0..4 { write_isv_signal_at(BASE + b, 0.0); }`.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three compile errors in crates/ml/tests/smoke_test_real_data.rs at lines 568,
644, and 715 — all three were calls to `collect_experiences_gpu(&market_buf,
&target_buf, ...)` where `market_buf` and `target_buf` were `CudaSlice<f32>`
allocated via `stream.alloc_zeros + memcpy_htod`.
The production API changed in cba9f25ed (Bug-1 close-out): both parameters now
require `&MappedF32Buffer` (mapped pinned DEVICEMAP, no HtoD copy). The test
helper `real_market_data` was never updated.
Migration applied to smoke_test_real_data.rs:
- Added `use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer` import
- Removed now-unused `type CudaSlice<T>` alias
- Changed `real_market_data` return type from `(CudaSlice<f32>, CudaSlice<f32>, usize)`
to `(MappedF32Buffer, MappedF32Buffer, usize)`
- Replaced `stream.alloc_zeros + memcpy_htod` with
`unsafe { MappedF32Buffer::new(len) } + write_from_slice` at lines 511-517
- Parameter renamed `stream` → `_stream` since it is no longer used by the buffer
construction (stream is still used by callers via `GpuExperienceCollector::new`)
Pattern mirrors production caller in
crates/ml/src/trainers/dqn/trainer/training_loop.rs:1322-1337.
The two `unsafe` block warnings in the test are expected (same `warn(unsafe_code)`
lint applied project-wide; production code carries the same warnings).
mamba2_hyperopt_p0_p1_fixes.rs was also reported as failed but its failure was
purely a cascading linker OOM kill from the smoke_test_real_data compile error,
not an independent type error — confirmed by `cargo build --test
mamba2_hyperopt_p0_p1_fixes` succeeding independently.
NOTE: The SP7 T7 magnitude_distribution smoke does NOT pass — see report in
commit message body. Root cause: four sp7_lb_* entries were registered in
state_reset_registry.rs (commit aa2854017) but their dispatch arms in
DQNTrainer::reset_named_state (training_loop.rs) were never wired. Error at
runtime: "unknown name 'sp7_lb_diff_var_cql'". This is a T7 wire-up regression
separate from the test-compile fixes in this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two clarity-only edits surfaced by code review:
- M1: compute_adaptive_budgets header docstring claimed floors of
"0.05 C51, 0.05 IQN, 0.02 CQL, 0.02 Ens" — stale post-T7 (IQN is now
0.11, others are sentinel-aware bootstrap). Rewritten to reflect the
SP7 contract: IQN keeps BASE_IQN=0.11 structural floor; C51/CQL/ENS
use sentinel-aware bootstrap.
- M2: last_ens_budget_eff docstring misattributed ENS to "SP5 Pearl 2
(flatness-modulated)" — Pearl 2 stopped writing BUDGET_ENS_BASE in
T6. Rewritten to state ENS has no active controller driver in SP7;
bootstrap-only.
Comment-only — zero runtime effect. cargo check clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Atomic per feedback_no_partial_refactor — launch site + consumer floor
change + stale doc are three halves of the same contract:
(a) launch_loss_balance_controller wired into the per-step pipeline
after launch_sp5_pearl_2_budget (FLATNESS_BASE populated) and after
backward (grad_decomp pinned slots populated).
(b) compute_adaptive_budgets replaces hard floors (0.02/0.05) with
sentinel-aware bootstrap. The controller may now drive budgets near
zero when the structural target says so. IQN keeps BASE_IQN=0.11
structural floor (reference, can't be 0).
(c) All 4 stale "B4/G5" budget-eff docstrings rewritten to reflect
actual ownership: last_iqn_budget_eff (Pearl 2 flatness),
last_cql_budget_eff (SP7 controller), last_c51_budget_eff (SP7
controller), last_ens_budget_eff (Pearl 2 flatness +
sentinel-aware bootstrap). The previous claims (e.g.,
"0.10×(1−regime)×health") were never implemented; per
feedback_trust_code_not_docs, code wins over docs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Four comment/docstring sites still described the pre-SP7 Pearl 2 as
writing "20 floats / 5 outputs / c51/cql/ens" after the T6 kernel-
signature shrink. Updated to reflect the current 6-arg kernel writing
8 floats (IQN + flatness only) into scratch[115..119, 127..131).
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Kernel signature reduced from 9 to 6 args; pearl_2_budget_update now
writes only budget_iqn[4] and flatness[4]. Launcher migrated atomically:
arg list shrinks, apply_pearls smoothing loop shrinks from 5 slot-blocks
to 2.
SCRATCH_PEARL_2_C51, SCRATCH_PEARL_2_CQL, SCRATCH_PEARL_2_ENS deleted as
orphan constants per feedback_wire_everything_up. The corresponding
producer scratch slots (111..115, 119..127) become reserved-for-future
inside the buffer.
CQL/C51/ENS budget ownership now lives in the SP7 loss-balance
controller (loaded in T5; wired in T7).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two clarity-only edits surfaced by code review:
- Rename `cql_dev` → `cql_sx_dev` in `launch_loss_balance_controller`.
The launcher uses `grad_decomp_result_dev_ptr + 6 * f32_size` which
is the cql_sx (post-budget) slot, not cql (pre-budget at offset 3).
The variable name now reflects that.
- Add `// 213` inline annotation on `base_wiener_offset` to match the
pattern used by all four sibling launchers (lines 11062, 11194, 11303).
No semantic impact — pointer arithmetic and parameter order unchanged.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
LOSS_BALANCE_CONTROLLER_CUBIN static, kernel slot in GpuDqnTrainer,
6 SCRATCH_LB_* index constants (218..242), pub(crate) launch_loss_balance_controller
that runs the producer + 24 apply_pearls_ad smoothing launches.
SP5_SCRATCH_TOTAL bumped 218→242 with updated docblock and allocation
comment to match the 24 new slots.
Component pointers into grad_decomp_result_pinned: IQN at offset 0,
CQL_SX at offset 6, C51 at offset 9 (3-float layout per
launch_grad_decomp docstring).
Producer-only — call site in training_loop.rs lands at T7 atomically
with the consumer floor change and stale-doc deletion per
feedback_no_partial_refactor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds the kernel to the build.rs manifest after pearl_1_ext_num_atoms_kernel
and before the SP5 Layer D producers (matching the chronological SP5/SP7
producer ordering). nvcc-compiled cubin lands under
$OUT_DIR/loss_balance_controller_kernel.cubin and is loaded by
gpu_dqn_trainer.rs in T5.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three comment-only edits surfaced by code review:
- M1: Annotate `tid >= 8` guard with `// 8 = 2 heads × 4 branches`.
- M2: Header pseudocode `actual_ratio = h_norm / iqn_norm` clarified —
the implementation correctly omits the inline `fmaxf` because the
outer guard already enforces `iqn_n >= EPS_DIV`. Comment now states
the invariant explicitly.
- M3: Header "No atomicAdd" line now cites `feedback_no_atomicadd` for
consistency with sibling pearl kernels.
Comment-only — zero runtime effect. cargo check clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Single-block 8-thread kernel: 2 heads (CQL, C51) × 4 branches. Reads
3-float views into grad_decomp_result_pinned for IQN/CQL_SX/C51, plus
FLATNESS_BASE for the per-branch target modulator. Writes new budgets +
raw Wiener observations (diff² and sample²) to producer scratch.
Per-branch flatness-modulated target ratio (CQL → ANCHOR×(1-flatness),
C51 → ANCHOR×flatness). Wiener-optimal α from per-branch EMA state
slots, clamped to [1e-4, 0.5]. Cold-start seeds at consumer-side
bootstrap (0.02 CQL, 0.05 C51) on first observation; sentinel-aware.
Producer-only commit; build.rs entry, kernel slot, launch site land in
subsequent tasks.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Rewrite the three sp5_budget_c51/cql/ens registry descriptions to be
accurate at HEAD aa2854017: replace present-tense "driven by"/"no longer
writes" with future-tense "will be driven by"/"will stop writing", replace
non-existent C51_BOOTSTRAP_BUDGET/CQL_BOOTSTRAP_BUDGET/ENS_BOOTSTRAP_BUDGET
named constants with their actual anonymous .max() literal line references,
and apply the "(Pearl 2 will stop writing...)" parenthetical uniformly to all
three slots (cql was missing it). Audit doc T2 bullet updated accordingly.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
4 new ResetEntries × 4 slots each cover ISV[297..313). Sentinel 0
triggers Pearl A first-observation bootstrap on fold boundary, matching
the existing Pearl 2/3/4/5/6/8 pattern.
Also: corrected the stale sp5_budget_cql description claiming a
non-existent "regime_stability allocator" (per
feedback_trust_code_not_docs); clarified sp5_budget_c51 / sp5_budget_ens
to reflect SP7 ownership.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Minor item 1: Remove extra column-alignment spaces from the 4 SP7
constants (LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE,
LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE) so they match the
no-alignment style of surrounding SP5 constants (BUDGET_C51_BASE etc.).
Constant values (297, 301, 305, 309) and inline comments are unchanged.
Minor item 2: Soften the Fix 31 T7 stub in dqn-wire-up-audit.md from
"sentinel-aware consumer" to "sentinel-aware bootstrap with bootstrap
constants matching the kernel's cold-start basis (defined in T7)"
so the audit entry does not forward-reference specific constant names
before they land.
Cargo check: clean (zero warnings, zero errors).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Update stale (118) → 137 unique-slot count and [174..294) = 120 → [174..313) = 139
range in the SP5_PRODUCER_COUNT docstring. Update inline comment unique-slot count
121 → 137 with SP7 T1 breakdown entry. Append SP7 Task 1 history paragraph matching
the established Layer D D3 paragraph style.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
LB_DIFF_VAR_CQL_BASE=297, LB_SAMPLE_VAR_CQL_BASE=301,
LB_DIFF_VAR_C51_BASE=305, LB_SAMPLE_VAR_C51_BASE=309. Bumps SP5_SLOT_END
297 → 313 and SP5_PRODUCER_COUNT 123 → 139. Helper fns mirror the
existing budget_*/flatness/q_var_per_branch pattern. Layout fingerprint
extended. slot_layout_no_overlaps_and_total_correct asserts the new
total.
Audit doc Fix 31 stub added; will be extended commit-by-commit through
Tasks 2-9 of the SP7 plan.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Critical re-review of v1 caught:
1. Task 6 retained Pearl 2 kernel signature with (void) no-op args
"to save 3-site cascade" — that's the partial-refactor anti-pattern
feedback_no_partial_refactor explicitly forbids. v2 does the full
contract change: signature shrinks 9→6 args, launcher migrates
atomically.
2. SCRATCH_PEARL_2_C51/_CQL/_ENS would become orphan constants —
feedback_wire_everything_up says delete or wire. v2 deletes them in
T6.
3. Audit-doc updates were separated into a single late commit — would
fail check_audit_doc_updates pre-commit gate on T1-T6. v2 folds Fix
31 stub into T1 and extends it commit-by-commit.
4. T5 referenced grad_decomp_*_result_dev_ptr field names that don't
exist — actual layout is one shared 27-float pinned buffer with
per-component byte offsets (iqn=0, cql_sx=24, c51=36). v2 uses
pointer arithmetic from grad_decomp_result_dev_ptr.
5. Tasks 1, 2, 4 had "if file shape X then Y, otherwise Z" placeholder
branches — writing-plans skill forbids these. v2 has concrete patches
based on reading the actual files.
Also corrected: T2 fixes the stale "regime_stability allocator"
description in state_reset_registry.rs alongside the new entries.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Outcome-driven controller replacing Pearl 2's hardcoded-0 CQL budget and
floor-pinned C51 budget with multiplicative ratio adaptation. Targets
`grad_split_bwd cql/iqn = 2.0` and `c51/iqn = 1.0` per slice (trunk, dir,
mag). Slow EMA (α=0.01) consistent with existing controller pearls; Pearl
A sentinel-bootstrap on fold boundary; Pearl D Wiener-optimal smoothing.
Replaces ghost docstring at fused_training.rs:3409 ("0.10×(1−regime)×health"
formula was never implemented). Pearl 2 keeps owning IQN budget (reference
denominator) and FLATNESS_BASE (consumed by NoisyNet σ pearl); stops
writing CQL/C51/ENS slots so the new controller is the sole driver.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The smoke template already compiled `train_baseline_rl`. Now also build
`evaluate_baseline` + `precompute_features` and copy + strip all 3 to
`/data/bin/$SHORT_SHA/` after PASS — exactly the layout that
train-multi-seed-template.yaml's `ensure-binary` cache check
(lines 235-246) looks up by SHA. A smoke-then-train sequence at the same
SHA now hits the cache and skips the ~6-min compile, exiting in the
~1-min pod-startup baseline.
- Drop readOnly:true on training-data PVC mount (cargo writes binaries
into /data/bin/$SHORT_SHA; read-side fxcache + market data unchanged)
- Build all 3 example binaries in a single cargo invocation (sccache-warm)
- Idempotent SHA-keyed dest dir; PASS-only so broken bins never cache
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes Fix 29 audit row #13 — the last ⚠ Stale row from the Bug-1
contract drift triage. Pre-fix, `backtest_plan_state_isv` extracted
raw_close as `features[bar*feat_dim + 0]`. Post Bug-1 (commit
`5a5dd0fed`) `features[..+0]` is z-normed log-return, NOT raw_close.
The resulting `equity = cash + position*raw_close` and
`unrealized = position*(raw_close - entry_price)` formulas mixed
z-normed-log-return as a dollar price, corrupting val plan_isv
slots [PNL_VS_TARGET] (slot 1) and [PNL_VS_STOP] (slot 2). Other
plan_isv slots (progress, conviction, drift, regime, remaining)
don't depend on raw_close and were correct pre-fix.
Resolution: route raw_close from the upload-once `prices` buffer
(layout `[n*max_len*4]` raw OHLC). Close is at column index 3 — the
same index `backtest_env_kernel.cu` already reads from for portfolio
mark-to-market (line 64 of that kernel; OHLC layout is canonical
across the val backtest path). Reading from `prices` aligns the val
plan_isv path with env_step's source-of-truth, eliminating the
mixed-units pathology end-to-end.
Sites fixed:
- crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu:74-100
Kernel signature: `const float* features` and `int feat_dim`
parameters dropped, replaced by `const float* prices` (the
[n*max_len*4] raw OHLC buffer). Raw_close read becomes
`prices[(w*max_len + current_step)*4 + 3]`. Multi-line comment
block documents the Bug-1 origin and the env_step parity
reference. The `bool have_close` / `prices != nullptr` guard
semantics preserved — callers without OHLC data fall back to
`unrealized = 0` exactly as pre-fix.
- crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs:~1956
Single launcher (`evaluate_dqn_graphed` chunk loop, the only
invocation site of `plan_state_isv_kernel`) updated to pass
`&self.prices_buf.dev_ptr` instead of `&self.features_buf.dev_ptr`
and to drop the now-unused `feat_dim_i32` local + `.arg()` call.
Inline comment explains the Bug-1 origin.
- docs/dqn-wire-up-audit.md
Stale-B row appended to Fix 30's table. The standalone Stale-B
DEFERRED paragraph at the bottom replaced by the commit summary
+ the Fix 30 closure note (all 4 ⚠ Stale and 1 ❓ Ambiguous rows
from Fix 29's deferred follow-ups now resolved).
Migration scope per `feedback_no_partial_refactor`: kernel signature
changed → every consumer migrates in the same commit. Single launcher;
verified via `grep -rn backtest_plan_state_isv` (only the kernel
definition + the gpu_backtest_evaluator launcher + load site appear).
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline (43.68s) clean.
- SQLX_OFFLINE=true cargo build -p ml --release --offline
--features cuda (1m 30s) clean; cubin recompiled via nvcc.
- Pre-commit DtoD-via-pinned guard passes (the prereq commit
`4d966e62f` migrated `gpu_backtest_evaluator.rs`'s buffers to
MappedF32Buffer, eliminating the 5 `_via_pinned` callers that
had blocked any prior staging of this file).
Refs Fix 29 row #13. `feedback_no_partial_refactor` (single launcher
migrated alongside kernel signature change in one commit),
`feedback_no_functionality_removal` (PNL_VS_TARGET / PNL_VS_STOP
slots preserved — only their data source corrected; the audit's
"drop the slots" alternative explicitly rejected),
`feedback_no_hiding` (no fallback to z-normed reads remaining;
kernel either gets real raw_close or falls through to
`have_close=false` with `unrealized=0`, identical to pre-fix
smoke-test semantics where prices==NULL),
`feedback_no_cpu_compute_strict` n/a (zero new host-side compute),
`feedback_no_htod_htoh_only_mapped_pinned` already satisfied
(`prices_buf` is `MappedF32Buffer` post the prereq migration),
`feedback_trust_code_not_docs` (the kernel comment said
"raw_close from features buffer" for months — accurate-when-written
pre-Bug-1, stale-after-Bug-1; verify-against-code disambiguates).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Prerequisite for the deferred Fix 30 Stale-B kernel-side fix
(`backtest_plan_kernel.cu` raw_close source). The DtoD-via-pinned
pre-commit guard (`scripts/pre-commit-hook.sh::check_no_dtod_via_pinned`,
commit `5275932f4`) blocks any commit that stages `gpu_backtest_evaluator.rs`
against the file's 5 pre-existing `clone_to_device_*_via_pinned` callers
that landed before the guard. Stale-B has to stage this file (to thread
`prices_buf` into the `backtest_plan_state_isv` launcher), so the
migration must land first as its own atomic commit per
`feedback_no_partial_refactor`.
Migration scope. Single file, four buffer fields:
- prices_buf CudaSlice<f32> → MappedF32Buffer [n*max_len*4]
- features_buf CudaSlice<f32> → MappedF32Buffer [n*max_len*feat_dim]
- portfolio_buf CudaSlice<f32> → MappedF32Buffer [n*8] (kernel-mutated)
- window_lens_buf CudaSlice<i32> → MappedI32Buffer [n]
Init sites (lines ~651-664 post-edit). 4 `clone_to_device_*_via_pinned`
calls replaced with `MappedF32Buffer::new(host.len())` +
`write_from_slice(host)`. No memcpy_dtod_async + stream.synchronize()
pair at construction — mapped-pinned coherence makes the kernel see
host writes after the next stream-sync barrier.
Reset site (`reset_evaluation_state`, line ~1485 post-edit). In-place
`self.portfolio_buf.write_from_slice(&portfolio_init)` replaces the
prior buffer-replacement
`self.portfolio_buf = clone_to_device_f32_via_pinned(...)`. No alloc
churn per epoch; the device pointer stays stable across resets which
matches MappedF32Buffer's intended use.
Consumer sites (17 kernel arg passes). `arg(&self.X_buf)` →
`arg(&self.X_buf.dev_ptr)` so the launcher receives the device pointer
the kernel expects. Sites: launch_gather (×2), launch_gather_chunk,
launch_env_step, env_batch_kernel chunked path, plan_state_isv kernel,
metrics_kernel. CUdeviceptr (u64) is passed by reference exactly as
metrics_dev_ptr already does at the metrics launch site (~line 2497).
Kernel-mutated MappedF32Buffer precedent. portfolio_buf is mutated by
the env_step kernel every step. The same file's `plan_diag_buf` (a
MappedF32Buffer) is also kernel-written via dev_ptr (lines ~1230-1256)
and host-read via host_ptr — direct precedent. The IQN τ migration
(commit `facbf76eb`) confirms cuMemHostAlloc DEVICEMAP for kernel
reads. The mapped_pinned.rs module docstring states explicitly:
"Kernels write through dev_ptr (with __threadfence_system())".
What this change touches:
- crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
Field types (lines 299-321), init block (lines 641-665), reset
(lines 1462-1471), 17 kernel arg passes across 6 launchers.
- docs/dqn-wire-up-audit.md
Fix 30 Stale-B paragraph extended with the prereq commit summary.
Stale-B itself remains DEFERRED (kernel-side fix lands next).
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline (44.38s) clean.
- SQLX_OFFLINE=true cargo build -p ml --release --offline
--features cuda (53.78s) clean.
- grep -n "clone_to_device_.*_via_pinned\|upload_.*_via_pinned"
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs returns
zero hits.
- grep -n "self\.\(prices\|features\|window_lens\|portfolio\)_buf"
shows every kernel-arg site followed by `.dev_ptr`. The only
non-`.dev_ptr` references are the field declarations, the
constructor moves into `Self { ... }`, and the
`write_from_slice` calls in `reset_evaluation_state`.
- Pre-commit DtoD-via-pinned guard now passes on staging this file.
Eliminates the last 5 `_via_pinned` callers in
`gpu_backtest_evaluator.rs`. Per
`feedback_no_htod_htoh_only_mapped_pinned`, `feedback_no_partial_refactor`
(every consumer of a field-type change migrates in one commit),
`feedback_no_hiding` (no `--no-verify`; the migration IS the fix the
guard is asking for).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes Fix 29 audit row #18. The synthetic-data overlay kernel
`phantom_liquidity_gbm` in `dqn_utility_kernels.cu:1050-1096` was
flagged ❓ Ambiguous because its writer wrote `log_ret` to
`market_features[bar*market_dim + 0]` and `+3` — the WRITER's
downstream consumer contract was unverifiable post Bug 1 (commit
`5a5dd0fed`) which redefined feature index 0 from raw log-return to
z-normed log-return.
Investigation result: zero callers in the entire codebase. The audit
grep `grep -rn "phantom_liquidity_gbm" crates/ml/src/ services/
crates/ bin/` returns ONLY the kernel definition itself — no
`load_function`, no `launch_builder`, no Rust-side launcher. The
kernel has been compiled but never invoked since at least early 2026.
Resolution: delete the kernel definition and replace with an
explanatory comment block. Per `feedback_no_hiding` an orphan kernel
with an ambiguous post-Bug-1 contract is exactly the failure mode the
rule warns against — a future re-wirer would have written
synthetic z-normed log-returns into a slot the production pipeline
expects raw log-returns at, re-introducing the very Bug-1 drift the
audit was triaging. Per `feedback_wire_everything_up` a kernel that
compiles but is unconsumed is either wired in the same commit or
deleted; this kernel was never wired since its writer existed.
The replacement comment block (23 lines) documents the design intent
(GBM-based synthetic-data augmentation) so any future re-introduction
has the spec available, plus the re-introduction conditions:
caller wiring must land in the same commit, and the writer/consumer
contract for `market_features[..+0]` (z-normed vs raw log-return)
must be explicit at the kernel boundary.
What this change touches:
- crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu:1037-1096
47-line kernel definition + ══ comment header deleted; 23-line
DELETED notice replaces it. Net -24 lines.
- docs/dqn-wire-up-audit.md
Fix 30 Ambiguous-A row appended with the dead-code-deleted
resolution. Stale-B deferral note also added (separate paragraph)
documenting the pre-existing `_via_pinned` guard block requires a
dedicated `gpu_backtest_evaluator.rs` migration commit before the
Stale-B kernel-side fix can land.
Verification:
- SQLX_OFFLINE=true cargo build -p ml --release --offline
--features cuda (1m 30s) clean; cubin recompiled via nvcc.
- Pre-commit guards pass.
- Final grep for `phantom_liquidity_gbm` returns only the comment
block in the kernel file.
Refs Fix 29 row #18 deferred follow-up. `feedback_no_hiding`
(delete orphans, do not leave them as ambiguous landmines),
`feedback_wire_everything_up` (every module/feature/kernel built
must be wired to a production path or removed),
`feedback_no_functionality_removal` does NOT apply: a kernel with
zero callers is not a functional feature, and the audit doc retains
the design intent so any future re-introduction can rebuild the
mechanism with the correct post-Bug-1 contract.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes Fix 29 audit row #12. Pre-fix, `scripted_policy_select` read
`state[MARKET_START]` as `close_now` and computed
`recent_ret = (close_now - prev_close) / prev_close`. Post Bug-1
(commit `5a5dd0fed`) `state[MARKET_START]` is z-normed log-return at
the WRITER, not raw_close, so the formula mixed two unrelated signals
(z-normed return vs dollar price). The seed-phase MOMENTUM /
MEAN_REV / VWAP_DEV branches degenerated to noise-floor below the
±0.0001f cutoffs, leaving 60% of seed-phase episodes (the non-UNIFORM
20%+20%+20%) effectively in DIR_HOLD instead of expressing the
intended scripted-policy diversity.
Resolution: route raw_close from the same fxcache target buffer the
env_step kernel reads from. Kernel signature gains
`const float* targets`, `const int* episode_starts`, `int t`,
`int total_bars` parameters. Per-thread `bar_idx = episode_starts[i]
+ t` (matching `experience_kernels.cu:1769`'s indexing convention),
then `close_now = targets[bar_idx*6 + TARGET_RAW_CLOSE]`. Bounds-clamp
to `[0, total_bars-1]` mirrors env_step's out-of-bounds early-return.
Sites fixed:
- crates/ml/src/cuda_pipeline/scripted_policy_kernel.cu
Kernel signature extended; close_now read source switched;
`FXCACHE_TARGET_STRIDE` / `FXCACHE_TARGET_RAW_CLOSE` mirrored as
#defines (kernels can't import Rust constants; co-locating the
literals in a comment-block keeps the cross-language contract
visible at the call site for any future TARGET_DIM bump). The
previous `bar_idx` parameter renamed `t`; the LCG seed now mixes
per-thread `bar_idx = episode_starts[i] + t` instead of the
per-step `t` — stronger entropy across episodes, no behavioural
regression (UNIFORM policy still produces 4-direction uniform).
- crates/ml/src/cuda_pipeline/gpu_experience_collector.rs:~3735
Single launcher updated to pass `&targets_buf.dev_ptr`,
`&self.episode_starts_buf`, `t as i32`, and `total_bars` (already
in scope from line 3301). Inline comment explains Bug-1 origin.
Migration scope per `feedback_no_partial_refactor`: single launcher;
no other call sites. Verified via
`grep -rn scripted_policy_select crates/ml/src/`. NO new `PS_*` slot
or `PORTFOLIO_STRIDE` bump required — the targets buffer is the
canonical raw_close source and routing it directly avoids cascading
through the ~12 PORTFOLIO_STRIDE consumers (kelly_cap_update_kernel,
trade_stats_kernel, gpu_experience_collector, ml-core mirror, etc.).
This is the "route raw_close directly" branch of the audit's contract
decision; the alternative "add an extra state slot" branch was
explicitly rejected as cascade-heavy for a seed-phase-only signal.
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline (43.87s) clean.
- SQLX_OFFLINE=true cargo build -p ml --release --offline
--features cuda (1m 30s) clean; cubin recompiled via nvcc.
- No host-side compute added; all reads on GPU.
- `targets_buf` is mapped pinned by upstream caller (per
`feedback_no_htod_htoh_only_mapped_pinned`).
Refs Fix 29 row #12. `feedback_no_partial_refactor` (single launcher
migrated in same commit), `feedback_no_functionality_removal`
(`recent_ret` signal preserved — only its data source is fixed; the
CUSUM-substitution alternative path explicitly rejected because the
recent_ret signal IS the scripted-policy contract, not an
implementation accident), `feedback_no_hiding` (no fallback to
z-normed-log-return reads; kernel either gets real raw_close or
clamps to total_bars-1 boundary), `feedback_no_cpu_compute_strict`
n/a (zero new host-side compute), `feedback_trust_code_not_docs` (the
kernel comment said `state[MARKET_START + 0] = close_now` for months
— accurate-when-written, stale-after-Bug-1; verify-against-code
disambiguates).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes Fix 29 audit rows #14 and #15 (Bug-1 contract drift in val/HPO
close-price extraction). Both sites read `target[0]` thinking it is
raw_close, but post Bug-1 (commit `5a5dd0fed`) `target[0]` is
preproc_close (z-normed log-return). The `fv[3]` fallback path is
unreachable on every production code path because
`set_val_data_from_slices` (in `dqn/trainer/mod.rs:1704`) always yields
`Vec<f64>` of length 6 from `[f64; 6]` slices post `TARGET_DIM=6` bump
(commit `063fd2716`), so `target.len() >= 2` is an always-true guard.
Per `feedback_no_hiding` the dead fallback is removed in the same edit
rather than left as a silent wrong-units path.
Sites fixed:
- crates/ml/src/trainers/dqn/trainer/metrics.rs:576
(val window_prices for GpuBacktestEvaluator)
- crates/ml/src/hyperopt/adapters/dqn.rs:2492
(HPO adapter val_close_prices for window-aggregated backtest)
Both now read `target[TARGET_RAW_CLOSE]` (col 2). Both import the named
constant from `crate::fxcache` so a future column rename moves the call
site with the writer (Fix 27 prevention pattern).
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline (8.03s) clean.
- No GPU code changed; cubin not affected.
Affects val Sharpe annualization + window equity curves on training
metrics; affects HPO val score on every trial. Pre-fix would yield
"prices" of magnitude ~stddev(log_return) (~7e-5 for ES 1-min) feeding
into PnL math that expects dollar prices, producing degenerate
backtest output. Post-fix prices are real raw_close values.
References Fix 27 Bug B (host-side Welford) — same kind of bug surfaced
at val/HPO consumer rather than training Welford. References
feedback_no_hiding (delete unreachable fallbacks rather than leaving
silent wrong-units fall-through), feedback_no_partial_refactor (both
consumers of the same (fv, target) tuple convention migrate together),
feedback_trust_code_not_docs (the `target.len() >= 2` guard read as
defensive but was masking a contract drift).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The kernel block at experience_kernels.cu:698-704 divided
market_features[0..3] by vol_normalizer at runtime — designed for the
pre-Bug-1 pipeline where features arrived as RAW log returns (~±0.001).
After Bug 1 fix (commit 5a5dd0fed) moved z-normalization to the WRITER
(precompute_features.rs::NormStats::normalize_batch before fxcache write,
plus data_loading.rs DBN-fallback path applying the same op), features
arrive already-z-normalized. The runtime division then created a
1000-13000× DOUBLE NORMALIZATION inflating column 0 of next_states to
raw-price magnitude.
DIAG_AUX_LABEL diagnostic ground truth (production train-multi-seed-bn42w):
- features_raw_cuda col 0 mean_abs = 0.443 (clean z-norm at SOURCE)
- aux_nb_label_buf mean_abs = 5398 (= 0.443 × inv_vol ~ 13245)
- ratio matches 1 / vol_normalizer for ES 1-min realised vol ~7.5e-5
This bug caused label_scale=5481 in 50-epoch validation (cancelled
train-multi-seed-bn42w) while smoke ran with smaller window producing
smaller inv_vol → label_scale ~25-432. Both wrong, just different
inflation factors.
What changed:
- experience_kernels.cu:698-704 block deleted; replaced with header
comment explaining why. Kernel parameter `vol_normalizer` retained
in signature with `(void)vol_normalizer;` to silence the
unused-warning — removing the param would cascade through
gpu_experience_collector.rs config struct + launcher + per-epoch
Welford in training_loop.rs (60 lines). Bounded scope: leave the
pipe wired, gut the consumer.
- DIAG_AUX_LABEL diagnostic removed per its Fix 28 removal gate
(gpu_dqn_trainer.rs ~12981 diagnostic block ~165 lines + the
DIAG_AUX_LABEL_SOURCE_PTRS OnceLock static ~22 lines +
training_loop.rs populate site ~22 lines).
- Audit doc Fix 29 entry with audit results + open follow-ups.
Bug-1 contract audit (this commit's exhaustive re-check, 18 sites
classified):
- 1 ⚠→✅: experience_kernels.cu:698-704 (this fix)
- 13 ✅: env-step kernel body, mirror universe, feature mask/noise,
target reads at col 2 raw_close (Fix-27 already correct), DT
rewards kernel, curriculum/hindsight/portfolio_sim kernel target
reads, kernel header docs, PS_PREV_CLOSE state-layout slot
- 4 ⚠ Stale (deferred — separate triage commits):
* scripted_policy_kernel.cu:59-65 — seed-phase momentum reads
state[MARKET_START] as raw_close, but post-Bug-1 it is z-normed
log-return. Affects seed-phase scripted policy quality only.
* backtest_plan_kernel.cu:77-100 — val plan_isv reads
features[bar*feat_dim+0] as raw_close. Corrupts val plan_isv
slots [PNL_VS_TARGET]/[PNL_VS_STOP].
* metrics.rs:576 — val_data → window_prices reads target[0] as
close (clones Fix-27 Bug B at host-side; affects val Sharpe).
* hyperopt/adapters/dqn.rs:2492 — same target[0] pattern in HPO
val_close_prices.
- 1 ❓ Ambiguous: dqn_utility_kernels.cu:1089-1093 synthetic feature
overlay; needs downstream-consumer contract verification.
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline (47.87s) clean.
- SQLX_OFFLINE=true cargo build -p ml --release --offline
--features cuda (1m 30s) clean; cubin recompiled via nvcc.
- Next L40S production run should show label_scale ~0.8 (matching
kernel docstring expectation).
Refs: DIAG_AUX_LABEL diagnostic from Fix 28 in dqn-gpu-hot-path-audit.md
(commit 2683d4637) which captured the ground truth that pinned this
bug. Cancelled validation runs train-multi-seed-p5qzw (label_scale=808)
and train-multi-seed-bn42w (label_scale=5481) both blocked on this —
Fix 27 cleared the host-side variant (Welford reading target[0] as
raw_close); this Fix 29 closes the kernel-side variant. Bug 1 chain
(label_scale=5443 from 4-month-old #193) closes here.
feedback_trust_code_not_docs (the kernel comment said `#13 Vol
normalization` for months — accurate-when-written, stale-after-Bug-1).
feedback_no_partial_refactor does not apply because the kernel
parameter is retained as a no-op, deliberately leaving the
launcher/config-struct contract intact while the consumer is gutted.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Production training run train-multi-seed-bn42w showed label_scale=5481
(raw_close magnitude) at epoch 4 vs smoke's ~25 at the same commit
29b1d34c6. The 5-layer data-loading defense + DBN spread filter + target
stride/column fix all in place — yet column 0 of next_states_buf still
sees raw-price-magnitude values in production but not in the
multi_fold_convergence test path.
This one-shot diagnostic fires once at the first training batch and
prints:
- aux_nb_label_buf stats (the kernel input; should be ~0.8 z-norm)
- next_states_buf col 0 sample (what strided_gather reads)
- features_raw_cuda col 0 sample (the source feature buffer)
- targets_raw_cuda col 2 sample (raw_close — the suspected leak)
- Verdict line interpreting the 4 stats
Three numbers will pin the source. If aux_buf ~5500: raw_close leaked
into next_states. If features_raw_cuda col 0 ~5500: the writer didn't
normalize. If targets_raw_cuda col 2 ~5500 but aux_buf ~0.8: only
production reads raw_close, smoke doesn't.
One-shot via static AtomicBool. Pre-graph-capture, no perf impact.
Removal gate: delete once the leak source is identified and fixed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two latent bugs converged in the cancelled 50-epoch run
(train-multi-seed-p5qzw at 96769d171, label_scale=808 vs smoke
baseline 22-28):
Bug A (latent since 063fd2716, 2026-04-19): TARGET_DIM was bumped 4→6
in fxcache (added raw_open at col 4 + mid_price_open at col 5), but
N consumer kernels in experience_kernels.cu + dt_kernels.cu hardcoded
stride 4 when indexing targets[bar*4+col]. Reading at the old stride
against the stride-6 buffer slid every lookup into the wrong bar's
data. Smoke harness paths (multi_fold_convergence) didn't exercise
expert_action_override / compute_difficulty_scores / hindsight_relabel
heavily; production 368k-bar run surfaced the drift via the
compounded label_scale = 30× expected magnitude.
Bug B (introduced today at 5a5dd0fed, 2026-05-02): the Bug 1 fix
changed targets[0] from raw_close to log-return-normalized
preproc_close. training_loop.rs::epoch_vol_normalizer's Welford pass
still read targets[0] expecting raw_close → output was
ln(log_return/log_return) ~ stddev 0.6 → triggered sanity-band warning
+ default fallback (5e-4); but downstream label_scale consumers fed
the corrupt value through.
Sites fixed:
HOST (1 site):
- training_loop.rs:570-573, :603 — w[i].1[0] → w[i].1[TARGET_RAW_CLOSE]
+ warning message updated
GPU (5 sites in experience_kernels.cu):
- line 3768 — kernel param doc OHLCV → fxcache TARGET_DIM=6
- lines 3806-3808 — bar*4+3 → bar*6+2 (raw_close column)
- lines 3974-3975 — i*4+2 → i*6+2 (stride-only)
- lines 4015, 4020 — bar*4+2 → bar*6+2 (stride-only)
- lines 3400-3404 — t_offset = global_idx*4 → *6 (stride-only)
- lines 1553-1561 — kernel header doc updated to 6-column layout
GPU (1 site in dt_kernels.cu):
- line 793 — kernel param doc OHLC → fxcache TARGET_DIM=6
- lines 803, 807 — i*4+3 → i*6+2 (raw_close col 2)
HOST docstring (decision_transformer.rs:1049) — [num_bars, 4]
→ TARGET_DIM=6 with reference to fxcache constant module.
Structural prevention: 6 named column constants
(TARGET_PREPROC_CLOSE / TARGET_PREPROC_NEXT / TARGET_RAW_CLOSE /
TARGET_RAW_NEXT / TARGET_RAW_OPEN / TARGET_MID_OPEN) added to
crates/ml/src/fxcache.rs co-located with the existing TARGET_DIM
(promoted from private to pub). Co-locating in the contract-owner
module means future renames or column adds force every consumer to
update at the same call site. Compile-time density test asserts
columns are dense and exhaustive. Host-side consumer training_loop.rs
imports TARGET_RAW_CLOSE; GPU kernels reference the constant module
in comments + use literal stride 6 with a header block documenting
the contract (kernels can't import Rust constants).
PPO consumer NOT in scope: ppo_experience_kernel.cu reads at stride 4
but PPO has its own set_raw_market_data path uploading 4 columns from
Vec<f64>. PPO write/read pair is internally consistent at stride 4,
unaffected by fxcache stride bump. Production train_baseline_rl.rs
flow doesn't invoke set_raw_market_data, so PPO GPU collector is
currently orphaned. Consolidating PPO onto fxcache target buffer is a
separate refactor.
Verification:
- SQLX_OFFLINE=true cargo check -p ml --offline clean
- cargo build -p ml --release --offline --features cuda clean
(cubin compiles via nvcc; release build under 2 min)
- cargo test -p ml --lib -- target_layout 1/1 pass
(target_columns_dense_and_exhaustive)
- audit grep "targets[var * 4 +]" returns ZERO hits in production
- re-run of train-multi-seed-p5qzw will show label_scale ~22-30
(not 808) — gating production validation
Refs: cancelled train-multi-seed-p5qzw at 96769d171, 50-epoch
validation blocker. Bug origins: 063fd2716 (target_dim 4→6 bump) and
5a5dd0fed (Bug 1 column-0 semantic fix). feedback_no_partial_refactor
(every consumer of shared buffer contract migrates in same commit),
feedback_trust_code_not_docs (kernel doc OHLCV/OHLC strings stale
post-bump), feedback_no_quickfixes (proper structural fix not
one-line patch), feedback_wire_everything_up (column constants land
co-located with fxcache::TARGET_DIM so writer + caller share single
contract module).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
ml::validation is gated behind validation-mod, but dqn/config.rs:177
references crate::validation::RegimeMetrics unconditionally. Both
services pull in dqn::config via the trainer pipeline, so they need
the gate opened to compile. Cheaper than refactoring 700+ refs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan task #296. Pearl 7 was an INVESTIGATION task in the SP5 brainstorm:
the pre-SP5 50-epoch baseline (train-multi-seed-cv2mw, F0 epochs 4-9)
showed intent_dist freezing at exact Bin(2, 0.5) ratios (0.25/0.50/0.25),
suggesting a hidden binary action decomposition somewhere downstream.
The plan §C4 closure rule: if post-SP5 smokes show intent_dist drifting
normally (no freeze), Pearl 7 closes with no code changes.
Verdict from 3 retained SP5-era smokes: intent_dist drifts smoothly each
epoch. No Bin(2, 0.5) freeze observed at:
- smoke-test-ks2wf (post-spread-filter, 5845e4403)
- smoke-test-7pv9v (Layer D additive, f42b5fff8)
- smoke-test-w9nsw (D4 atomic, 2e9e276a0)
(smoke-test-cnlrw (sanitize-only, 8434737a6) cached log was pruned
before evidence-collection; 3 retained smokes sample post-spread-filter
and full Layer D production paths and decisively satisfy the closure
rule on their own.)
Likely cause: SP5 Layer A's per-branch parameter lifting (C51 atom span,
NoisyNet σ, IQN τ schedule, loss budgets, Adam β/ε, Kelly floors) added
enough independent variability at every shared site that no single
2-state decomposition can dominate intent_dist in steady state.
What lands:
- docs/dqn-wire-up-audit.md Pearl 7 closure entry
- memory/pearl_intent_dist_freeze_resolved.md (new)
- memory/MEMORY.md Topic Files Index entry
No code changes. No follow-up spec opened. The escalation trigger
(re-freeze for ≥3 consecutive epochs in a future run) is documented
in the audit entry.
Closes plan task #296.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
