User directive: "we don't remove functionality at all, we fix what's
broken". Standing rule captured in
~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/feedback_no_functionality_removal.md.
Phase 2 plan reframed:
* Task 2.3 (R5 micro-reward): was DELETE; now FIX — either (a) set a
non-zero scale that delivers real signal, or (b) document why it's
intentionally disabled while preserving the code path.
* Task 2.4 (R6 loss-aversion): unchanged — the relocation to C51
Bellman target is consolidation, not removal. Invariant preserved.
* Task 2.6 (E4 entropy-reg): was DELETE-CANDIDATE; now TUNE — if the
signal doesn't reach magnitude, re-route it rather than remove.
* Task 2.7 (C4 grad-clip): was DELETE-CANDIDATE; now
KEEP-AND-DOCUMENT — gradient clipping is a safety feature; run the
ablation for diagnostic, not for removal.
* H9 delete-magnitude-branch fallback: rejected permanently. On
magnitude convergence failure the fallback is per-magnitude reward
shaping / per-bin advantage weighting / curriculum / state enrichment.
Task 2.5 wiring-bug sweep unchanged (correctness fixes, not removal).
This commit updates only the planning narrative. Task execution hasn't
started on 2.3/2.6/2.7 yet, so no code changes required here. When
those tasks run, they will implement the fix-paths, not the
delete-paths.
Task 2.0 instrumentation (commits d60e5375a / 980f3b07f / 41b0c559c)
revealed two silent bugs in Task 0.4's grad_ratio_mag_dir accessor:
Bug A — readback size mismatch: pinned buffer allocated at
total_params, but grad_buf length is total_params+cutlass_tile_pad,
so size check always failed → Err silently coerced to 0.0 by proxy.
Bug B — readback timing wipe: estimate_avg_q_value_with_early_stopping
in process_epoch_boundary replays the forward graph, which zeros
grad_buf. Any subsequent readback sees all zeros.
Both fixed in 41b0c559c; snapshot hoisted to top of process_epoch_boundary.
With those bugs fixed, the measured gradient ratio is NOT 0.0000 — it is
50–400× mag/dir across 60 epochs. Magnitude branch is over-fed, not
starved. Direction gradient is small but non-zero (~2e-2 to 7e0).
Direction policy is observably healthy (Short/Hold/Long/Flat 38/12/42/14%).
Track 1 triage's H4 CONFIRMED verdict was a measurement artefact
produced by Bugs A+B. H4 as originally defined is now REJECTED.
Plan changes:
- Add new "Task 2.0 findings" section after cross-cutting concerns,
documenting bug A, bug B, observed data, and revised strategy.
- Task 2.1 marked DEFERRED (not deleted — kept for reference).
- Task 2.2 promoted to PRIMARY fix.
- New fallback: H9 delete-magnitude-branch as replacement for Task 2.1
if Task 2.2 alone is insufficient.
- Task 2.0 inventory row updated with LANDED status and commit chain.
Net effect: Phase 2 simplifies. The hardest task (2.1 three-branch
architectural fix) is skipped; primary path is Task 2.2 (~25 LOC).
First Task 2.0 dispatch escalated BLOCKED: the four loss-component
backward kernels are captured inside the fused training graph, so
host-side snapshot-between-components isn't possible mid-graph without
a force-ungraphed diagnostic step (~210 LOC + cross-stream sync risk).
Revised approach (chosen after cost analysis):
cudaMemcpyAsync(device → pinned host) IS captureable in a CUDA graph.
Even better: DtoD into per-component scratch buffers, then an in-graph
reduction kernel computes per-component (mag_norm, dir_norm) and writes
8 floats to a pinned result slot. Only the 8-float result crosses
PCIe (at epoch boundary), keeping per-step PCIe traffic to zero.
Changes to the plan's Step 2 + Step 3 + Step 4:
- Step 2: added 4 device-side scratch buffers (one per component,
~10 MB each = 40 MB device) + 8-float pinned result slot + new
reduction kernel grad_decomp_kernel.cu spec'd out.
- Step 3: clarified that DtoD snapshot + backward + reduction kernel
are ALL captured in the graph; graph replays them every step;
no force-ungraphed dance needed.
- Step 4: added refresh_grad_component_norms() accessor that reads
the 8-float pinned slot at epoch boundary (zero-copy) and populates
the host-side cache.
Approach matches Task 0.4 pattern (commit bb42c9963) extended four-fold.
No atomicAdd (reduction uses shared-mem tree), no non-captured replay,
no cross-stream sync risk.
LOC estimate: ~90 (was ~210 for the rejected option).
Three polish items from the plan review:
1. Task 2.4 Step 5 — replaced vague "must not regress" with a numeric
tolerance-band table: multi_fold best_val_metric ±15% per fold,
Best Sharpe ≥ 20 floor, and hard F_Half/F_Full ≥ 0.05 + cf_flip ≥ 0.1
BLOCKERS. Anchors to the baseline metrics doc values captured at
policy-quality-baseline.
2. Task 2.6 — converted the ad-hoc text decision at Step 1 into the same
five-row decision table format Task 2.1 uses (DELETE / KEEP / KEEP-as-
safety-net / INCONCLUSIVE / DELETE-because-inseparable). Step 2
ablation got its own 4-row outcome table keyed to ent_mag delta,
multi-fold Sharpe regression, and NaN appearance.
3. Task 2.7 Step 1 — added "Repeat 3× with different seeds" clause and
explicit total wall-clock note: ~75 min (3 × 25 min) for the ablation
pass before the Step 2 decision. Aligns with the Cross-cutting concern
#6 about sample-noise rejection.
No task count change (still 11: 2.0–2.10). Net code-delta estimate
unchanged. Standing-rule compliance unchanged (no stubs / no atomics
/ no quickfixes / no hiding / no feature flags / no push-per-task).
User preference: everything lands on main, ~500 LOC is small enough
that feature-branch isolation isn't needed.
Spec changes:
* §2.3 outcome paths: drop "unmerged branch" vocabulary
* §3 architecture: main-branch timeline, no feat/, no wip/
* §4 Phase 0: commits on main, tag policy-quality-baseline at end
* §5 Phase 1: experiments are worktree edits reverted after, only
triage docs commit (to main)
* §6 Phase 2: commits land on main, author chooses granularity,
each leaves smokes green
* §7.3 outcome handling: iterate on main, baseline tag for rollback
* §7.4 closing: tag policy-quality-v1, update status
Plan changes:
* Task 0.1: verify-starting-state (not branch creation)
* Task 1.x: temp-edit-then-revert (not wip branches)
* Task 1.5: triage docs commit directly to main
* Phase 4: tag + close (not merge + cleanup)
* Rollback: one git reset --hard command
target_dim expansion: adds raw_open (OHLCV) and mid_price_open
(MBP-10 midpoint at bar formation) to fxcache targets. FXCACHE_VERSION
2→3 for auto-rebuild. Legacy v2 files handled with close-price fallback.
Spec v5 adds 3 pearls:
- Bar duration encoding in Mamba2 (continuous-time SSM awareness)
- Order book center of mass from all 10 MBP-10 levels (aggression signal)
- Retrospective hold quality bonus (teaches exit timing)
Plus: Hold action (4th direction), DSR Sharpe EMA fix, counterfactual
magnitude/order sign fix, MFT mid-price mark-to-market.
OFI embed MLP now 18→10 (was 16→8). Mamba2 width SH2+10 (was SH2+8).
Attention width D+10 (was D+8).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Task 9d: Q-mean EMA drift penalty (lambda=0.01) in C51 grad kernel.
Prevents +0.47 Q-mean drift observed in train-vpb4w.
Task 9e: 500-step LR warmup for new tensor groups (regime gate,
adaptive atoms, Mamba2, multi-horizon). Prevents gradient shock
on new components introduced to trained network.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Task 9a: Adaptive atom position training via atom entropy gradient.
Numerical finite-difference gradient on spacing_raw, SGD at LR=1e-3.
Task 9b: Mamba2 BPTT backward through K=8 scan steps.
d_W_A/B/C accumulated, separate Adam at LR=1e-4.
Task 9c: Multi-horizon C51 loss — 3 n-step passes (1/5/20 bar),
3 C51 loss launches, regime-weighted gradient blend (ADX-based).
Fixes the "untrained component" gap identified in plan review.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Task 10: q_mean_reduce + q_mean_subtract two-phase kernel. Zero params.
Fixes Q-mean drift 0→+0.47 from bootstrapping bias.
Task 11: Adaptive entropy_coeff based on utilization_ema. Zero params.
Fixes atom utilization collapse 100%→20%.
Execution order: 10, 11, then 1-9.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Raw E[Q] ~ 0.1 while h_s2 ~ 1.0 post-ReLU. Without normalization,
Q_dir columns learn 10x slower. Dividing by delta_z scales to ~1-10.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
6 new CUDA kernels, C51/grad/epsilon kernel mods, dual-tau IQL,
v_range dead code deletion, per-sample everything.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Fundamental fix: C51 atom spacing (dz=0.6) was larger than the Q-value
range (0.02), making the distributional loss unable to resolve rewards.
The network was blind to its own learning signal.
Adaptive v_range via device buffer (same pattern as tau_buf):
- v_range_buf[2] on GPU, all C51/MSE/CQL/expected_q kernels read via
pointer (graph captures stable address, reads value at replay time)
- Cold start: v_range=±1.0 → dz=2/51=0.039 → 25× atom resolution
- Monotonic expansion based on Q-stats every 50 training steps
- Never shrinks (avoids invalidating Bellman targets in replay buffer)
Kernel changes (5 files):
- c51_loss_kernel.cu, mse_loss_kernel.cu, mse_grad_kernel.cu,
cql_grad_kernel.cu, compute_expected_q: scalar v_min/v_max → pointer
Also fixed:
- PopArt warmup 100→1 (start normalizing from step 1)
- PopArt variance floor 0.0001→1e-8 (allow sigma < 0.01)
- Backtest evaluator: allocate own v_range_buf for compute_expected_q
- num_atoms 51→52 in all TOMLs (TF32 4-element alignment)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
DuelingWeightSet becomes raw u64 pointers into params_buf, not separate
CudaSlice allocations. Eliminates ALL D2D copies between weight
representations. flatten/unflatten cycle deleted entirely.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Complete bf16 elimination across all crates (ml, ml-core, ml-dqn, ml-ppo,
ml-supervised). Zero half::bf16, __nv_bfloat16, or CudaSlice<half::bf16>
references remain (verified by grep).
CUDA: All 60+ .cu kernels and .cuh headers converted to native float.
- Half-precision intrinsics (__hmul, __hadd, __hdiv) → float operators
- atomicAddBF16 → native atomicAdd(float)
- bf16 wrapper functions → f32 identity passthroughs
Rust: All CudaSlice<half::bf16> → CudaSlice<f32> across 90+ files.
- htod_f32_to_bf16/dtoh_bf16_to_f32 → htod_f32/dtoh_f32 (direct, no conversion)
- Deleted bf16 mirror infrastructure (DuelingWeightSetBf16, alloc_bf16_mirror, etc.)
- Renamed params_bf16→params_flat, d_value_logits_bf16→d_value_logits, etc.
- Fixed .to_f32() sed damage on Decimal::to_f32() and rng.f32()
FxCache: Single f32 disk format (was bf16/f64 dual-version).
- Deleted --bf16 CLI flag from precompute_features
- PVC cache files need regeneration via precompute_features
TF32 tensor cores activated via cublasLtMatmul CUBLAS_COMPUTE_32F — no
explicit TF32 types needed. Storage is pure f32 everywhere.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
CUBLAS_COMPUTE_32F_FAST_TF32 returns NOT_SUPPORTED on RTX 3050 (SM86)
for certain (m,n,k) via cublasLtMatmul (e.g., 128×512×64). Plain
CUBLAS_COMPUTE_32F works for all dimensions. TF32 tensor cores are
still enabled via CUBLAS_TF32_TENSOR_OP_MATH math mode on the cuBLAS
handle — this activates TF32 on Ampere+ automatically when the
hardware supports the specific matrix dimensions.
Also: use sys::cublasLtMatmul directly (bypass result:: wrapper)
for clearer argument ordering.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause of H100 gradient collapse: c51_grad, mse_grad, and cql_grad
kernels wrote d_adv_logits in sample-major interleaved layout [B, TBA]
but cuBLAS backward read it as branch-major [B*B0*NA | B*B1*NA | ...].
At batch=16384, the GEMM read cross-contaminated data from wrong branches,
producing zero weight gradients by symmetry cancellation at scale.
Fix: changed indexing in all 3 grad kernels to branch-major.
Reverted diagnostic hacks (ungraphed forward, forced diagnostics).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- iqn_lambda restored to 0.25 (was 0.0 for NaN isolation)
- STEP_DIAG per-step logging removed (was temporary)
- FOXHUNT_GRAD_DIAG env var removed from Argo template
IQN NaN root cause fixed in 8cbabcef5 (f32-as-bf16 type mismatch).
Ready for H100 baseline deployment.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
C51 cross-entropy structurally prefers low-variance actions (Small positions
have tighter return distributions). MSE is variance-neutral. By reducing
C51's gradient to 20% and amplifying MSE to 4× for the magnitude branch,
MSE becomes the dominant loss signal for position sizing.
Direction/order/urgency keep full C51 gradient for distributional risk awareness.
Replaces the previous mag_amp=1.5 which addressed a symptom (Flat gradient
starvation) rather than the root cause (C51 variance bias).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The mirror universe feature negated only features[0..3] (4 return-like
features) but inverted the direction action (Short↔Long). The remaining
38 directional features (momentum, trend, patterns) were NOT negated.
This taught the model the WRONG direction on 50% of epochs, cancelling
out the directional signal and producing worse-than-random performance.
Disabled until proper full-feature mirroring is implemented.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
