- Add final_sharpe and final_val_loss to TrainingMetrics (last-epoch values)
- Rewrite test_walk_forward_no_overfitting_50_epochs: all checks are relative
to the run's own metrics — no hardcoded thresholds
- Checks: finiteness, gradient health, OOS collapse ratio, IS→OOS gap ratio
- 19/19 smoke tests pass
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Extract shared resolve_data_dir() function — resolves relative data paths
against workspace root via CARGO_MANIFEST_DIR. Used for mbp10_data_dir
and trades_data_dir (was only mbp10 before, causing smoke test failures).
All 19/19 smoke tests pass with pure f32 pipeline.
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>
Root cause: cuMemAllocAsync (stream-ordered memory pools) produces
memory that cublasLtMatmul rejects with CUBLAS_STATUS_NOT_SUPPORTED
on H100 (CUDA 13.0 driver) when the matmul runs on a different
stream than the allocation. Despite CUDA documentation stating
same-device async allocations should be accessible from all streams
with proper event synchronization, cublasLtMatmul disagrees.
Proof chain:
- C++ test with cudaMalloc: ALL dimensions pass on H100 ✓
- Rust test with cuMemAlloc: ALL dimensions pass on H100 ✓
- Rust test with cudarc alloc_zeros (cuMemAllocAsync): FAILS ✗
- Same dimensions, same handle, same workspace — only alloc differs
Fix: set has_async_alloc=false unconditionally in CudaContext::new().
This forces cuMemAlloc_v2 for ALL allocations. Zero performance impact
(allocations at construction time, not in hot loops).
Also: converted has_async_alloc from bool to AtomicBool for safety.
Removed diagnostic code (CUBLASLT_DIAG, test_lt_matmul_raw).
Reverted cuMemAlloc workspace hack (now uses normal alloc_zeros).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: cudarc uses cuMemAllocAsync (stream-ordered memory pools)
on H100 where CU_DEVICE_ATTRIBUTE_MEMORY_POOLS_SUPPORTED=1. These
allocations are only accessible from the allocating stream.
cublasLtMatmul dispatches GEMM work on branch streams (multi-stream
fork-join for parallel advantage heads), which can't access
stream-local memory from the main stream — returns NOT_SUPPORTED.
Proof: C++ test with cudaMalloc passes ALL dims on H100 ✓
Rust test_lt_matmul_raw with cuMemAlloc passes on H100 ✓
Same Rust code with cudarc alloc_zeros (cuMemAllocAsync) FAILS ✗
Fix: add CudaContext::disable_async_alloc() to vendored cudarc.
Called in DQN::new_with_stream() to force cuMemAlloc_v2 for ALL
allocations. No performance impact (allocations at construction,
not in hot loops).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cudarc's stream.alloc_zeros uses cuMemAllocAsync (stream-ordered).
These allocations are only accessible from the allocating stream.
When cublasLtMatmul runs on a forked BRANCH stream (multi-stream
branch dispatch), the stream-ordered workspace is inaccessible,
causing CUBLAS_STATUS_NOT_SUPPORTED on H100 at batch=4096.
Fix: allocate workspace via cuMemAlloc_v2 (synchronous, globally
accessible from all streams). This matches the C++ debug test
which uses cudaMalloc and works on H100 for all dimensions.
Also: add d_layout destroy (leaked handle cleanup).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The forward lt_matmul destroyed a_layout, b_layout, c_layout, matmul_desc,
and matmul_pref — but NOT d_layout. The leaked handle caused cublasLt to
return stale descriptors on subsequent create calls. On H100 at batch=4096,
the dimension mismatch between the stale first-call layout and the actual
GEMM dimensions triggered CUBLAS_STATUS_NOT_SUPPORTED.
On RTX 3050 with small batches, the stale layout happened to be compatible
(oversized) for all subsequent GEMMs, masking the bug.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Experience collector: exp_h_s1, exp_h_s2, exp_h_v changed from
CudaSlice<half::bf16> to CudaSlice<f32>. These were unused legacy
fields but their bf16 type was inconsistent with the f32 pipeline.
Backtest evaluator: states_buf and chunked_states_buf padded to
state_dim_padded (pad128). gather_states kernel updated with padded_sd
parameter. DtoD copy uses padded row bytes.
C++ debug test confirms cublasLtMatmul works on H100 for all
dimensions including (128,4096,256) and (128,16384,256) with
both COMPUTE_32F and FAST_TF32.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause of cublasLtMatmul CUBLAS_STATUS_NOT_SUPPORTED:
- states_buf allocated [batch * state_dim] (unpadded)
- cuBLAS forward GEMM reads with stride state_dim_padded (pad128)
- Buffer overflow: GEMM reads past buffer end
cublasLtMatmul validates buffer sizes against layout descriptors and
returns NOT_SUPPORTED for undersized buffers. cublasSgemm silently
read garbage — this was the source of the "parallel test congestion
errors" seen previously.
Fix:
- Allocate states_buf with state_dim_padded stride (3 allocation sites)
- gather_states kernel: add padded_sd parameter, write with padded stride
- DtoD copy: use padded row bytes
- Both gather_states call sites updated with padded_sd arg
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>
cublasLtMatmul passes workspace+stream per-call (no handle state),
eliminating the cublasSetStream workspace conflict that hung CUDA
Graph mega-capture on H100 with cublasGemmEx.
Forward (22 call sites via sgemm_f32/sgemm_f32_ldb):
- cublasLtMatmul with CUBLAS_COMPUTE_32F + per-call stream
- No more cublasSetStream for branch dispatch (stream passed directly)
- 32MB workspace unconditionally (TF32 HMMA needs it even on Ampere)
Backward (7 call sites via sgemm_f32):
- Same cublasLtMatmul pattern with flexible transa/transb
- Stream parameter threaded through backward_fc_layer methods
TF32 tensor cores enabled via CUBLAS_TF32_TENSOR_OP_MATH math mode
on the cuBLAS handle. CUBLAS_COMPUTE_32F (not FAST_TF32) used in
matmul descriptors for SM86 compatibility.
19/19 smoke tests pass (sequential). Ready for H100 graph_mega test.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cublasSetStream() resets workspace to the default pool, which uses
stream-ordered allocation — incompatible with CUDA Graph capture.
This caused cublasGemmEx TF32 to hang inside graph_mega on H100.
Fix:
- Store workspace raw pointer + size in CublasForward struct
- Call cublasSetWorkspace_v2 after every cublasSetStream (5 sites)
- Add 32MB workspace to CublasBackward (had none — used default pool)
- Increase forward workspace 4MB → 32MB (TF32 HMMA needs more)
Root cause from NVIDIA docs: "cublasSetStream() unconditionally
resets the cuBLAS library workspace back to the default workspace pool"
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cublasSgemm was running on pure f32 CUDA cores (~60 TFLOPS on H100),
causing 15× slowdown vs the old bf16 tensor core path (77s vs 5s/epoch).
cublasGemmEx with CUBLAS_COMPUTE_32F_FAST_TF32 + CUBLAS_GEMM_DEFAULT_TENSOR_OP
uses TF32 tensor cores (~500 TFLOPS on H100) — 8× faster than pure f32.
The earlier NOT_SUPPORTED error was caused by bf16 weight pointers feeding
f32 GEMMs (stride mismatch), not API incompatibility. Now that all pointers
are f32, TF32 GemmEx works on both RTX 3050 (SM86) and H100 (SM90).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The evaluate closures now receive &CudaSlice<f32> from the backtest
evaluator. Updated all 3 closures (DQN, PPO, supervised) to match.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Logs L2 norms of critical buffers after forward+backward+aux, before Adam.
Fires on first 3 steps of each epoch only. Identifies exactly which buffer
goes zero on H100 — is it the loss gradient or the cuBLAS backward output?
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Blend used b*na but cast used b*pad32(na). Now both use the padded
size so scale/saxpy process the same element count as the cast.
Padding elements are zero — no behavioral change, just consistency.
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>
The old sample-major interleaved layout did not match cuBLAS backward
branch-major pointer arithmetic. At batch=16384, backward GEMM read
cross-contaminated data producing zero weight gradients.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The grad_norm_finalize kernel writes both f32 (sum-of-squares) and bf16
(L2 norm). The IQN and ensemble trunk gradient computations were passing
the MAIN grad_norm_buf for the bf16 output, overwriting it mid-pipeline.
While compute_grad_norm_outside_graph() overwrites it again before Adam,
this was still a correctness hazard. Now uses a dedicated 1-element bf16
scratch buffer (aux_norm_bf16_scratch) for IQN/ensemble finalize calls.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
compute_grad_norm_outside_graph() called launch_grad_norm() which uses
the graph-captured CUfunction handle. After graph_mega capture, this
handle cannot be launched outside the graph — CUDA silently corrupts
kernel state on H100 (SM_90). The corrupted grad_norm_f32_buf caused
Adam to read bogus norms, zeroing all gradient updates.
Fix: use launch_grad_norm_standalone() which has its own separately
loaded kernel handle, safe for post-graph launches.
This was the ROOT CAUSE of the H100 gradient collapse at epoch 3-4.
The collapse happened after mega-graph capture (step 2 of each epoch),
and the corrupted norm persisted for all subsequent steps.
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>
The IQN backward kernel writes f32 via atomicAdd into d_h_s2_buf
(allocated as CudaSlice<f32>). But apply_iqn_trunk_gradient called
bf16_to_f32_kernel on this buffer, reinterpreting f32 bits as bf16.
This produced garbage/NaN that poisoned the trunk gradient on H100.
Replace with direct f32 DtoD copy — no cast needed.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Graph recapture ruled out as NaN cause (NaN persists without graph
invalidation). Now testing: does NaN disappear when IQN is disabled?
If yes → IQN backward/Adam produces NaN on H100.
If no → NaN is in the main C51/MSE pipeline.
Also reverts graph invalidation hack from previous diagnostic commit.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
TEMPORARY: graphs keep stale alpha (wrong training, but isolates NaN source).
If NaN disappears → recapture path is the bug.
If NaN persists → something else causes it.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The f32→bf16 cast in cast_d_logits_to_bf16() reads
b*(b0+b1+b2+b3)*na + 32*4 elements, but d_adv_logits_buf and
d_adv_logits_bf16 were allocated with only +32*3 padding (96 vs 128).
The 32 extra f32 reads went past the allocation into adjacent GPU
memory. On H100 with 80GB and different memory layout, this grabbed
NaN/garbage values. The NaN propagated through the cuBLAS backward
pass into grad_buf, killing training at step ~720 (non-deterministic
depending on what lives in adjacent memory).
Fix: allocate +32*4 padding in d_adv_logits_buf, d_adv_logits_bf16,
and d_adv_logits_mse to match the cast kernel's access pattern.
Also reverts c51_alpha_max default to 1.0 — the gradient collapse
was from this buffer bug, not C51 premature convergence. The
c51_alpha_max parameter is kept for hyperopt exploration.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Logs grad_norm, loss, c51_alpha every 50 steps via existing async readback
scalars. No stream sync or extra GPU work. Works with mega-graph path.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Temporary diagnostic: logs per-stage gradient norms (pre-clip, post-clip,
post-IQN, final) to identify where H100 gradient collapse occurs.
Remove after root cause is found.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
H100 baseline (20 epochs) showed gradient collapse at epoch 10: C51
cross-entropy converges its distributional fit before the policy converges,
leaving zero gradient signal. The collapse happened 5 epochs after C51
reached alpha=1.0 (pure C51, zero MSE).
Fix: cap the C51 alpha ramp at c51_alpha_max (default 0.5) so MSE always
contributes (1 - alpha_max) of the primary gradient. MSE loss measures
Q-error directly and only goes to zero when Q-values are correct, not
just when the distribution shape is right.
- c51_alpha_max added to DQNHyperparameters (default 0.5)
- Added to PSO search space as 15th dimension (range [0.3, 0.9])
- Added to TOML profiles: smoketest, production, hyperopt
- Training loop caps alpha ramp at alpha_max instead of 1.0
- All 907 ml tests + 300 ml-core tests + 6 smoke tests pass
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three root causes found and fixed:
1. SUM-reduced gradients without 1/N: all CUDA loss gradient kernels
(C51, MSE, IQN backward, CQL) accumulated per-sample gradients as
raw SUM. At batch=16384 (H100) the raw norm was 282x larger than
batch=58, causing gradient clipping to destroy signal-to-noise ratio
and collapse training at epoch 2-3. Now all kernels multiply by
1/batch_size, making gradient scale batch-invariant.
2. ExposureLevel::target_exposure() used a flat 9-level scale that did
not match the 4-branch dir*mag encoding. The Rust backtest evaluator
computed wrong position sizes (e.g. 4x oversize for Short+Small).
Now uses dir x mag formula. Also fixed is_buy/is_sell/is_hold and
from_trading_action for 4-branch semantics.
3. Rust epsilon-greedy only explored 5/9 exposure combos (0..5 instead
of dir*3+mag), ignored the magnitude branch on greedy, and used wrong
indices for order/urgency (get(1)/get(2) instead of get(2)/get(3)).
LR recalibrated: old gradient_clip_norm was accidentally a batch-size-
dependent LR reducer (~60x at batch=58, ~16000x at batch=16384). With
mean-reduced gradients the clip rarely fires, so LR is now the sole
training speed control. Smoketest 1e-4 -> 2e-6, production 1e-4 -> 1e-5,
hyperopt range [1e-5,3e-4] -> [1e-7,1e-4].
Diagnostics: FOXHUNT_GRAD_DIAG=1 enables per-stage gradient norm logging.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
RUN 4 (1540c0287) survived 9 epochs with stable grad_norm=0.44.
That commit INCLUDED reward normalization (÷pos_frac). The revert
in 09fb80baa removed it, and the resulting H100 run collapsed at
epoch 7 (grad_norm=0.000).
The reward normalization is part of the stable configuration.
experience_kernels.cu restored to exact RUN 4 state.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
