ShortSmall=0, ShortHalf=1, ShortFull=2, Flat=3, LongSmall=4, LongHalf=5,
LongFull=6. Updates all array sizes, bounds checks, names, and distribution
logic across gpu_monitoring, monitoring, training_loop, metrics, and financials.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Task 3: Enable PopArt reward normalization (default true, all 3 TOML configs,
wiring confirmed in fused_training.rs submit_forward_ops_main()).
Task 4: Couple PopArt variance with tau reset — add prev_popart_var to
FusedTrainingCtx, read_popart_var() to GpuDqnTrainer, read_popart_variance()
+ should_reset_tau() to FusedTrainingCtx, tau reset injected at epoch boundary
in training_loop.rs after log_phase_timing().
Task 5: Add best_sharpe/best_epoch/best_val_loss reset to reset_for_fold() so
each walk-forward fold competes independently. Also wire v8 reward fields
(popart_enabled, micro_reward_scale, td_lambda, max_trace_length,
hindsight_fraction, hindsight_lookahead) through training_profile.rs apply_to().
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The 4-branch DQN (direction x magnitude) had 3 degenerate variants
(Short25, Flat, Long25) that all mapped to 0.0 target exposure when
direction=Flat, causing 82% Flat collapse. Collapse these into a
single Flat variant, giving 7 levels (ShortSmall/Half/Full, Flat,
LongSmall/Half/Full) and 63 total factored actions (7x3x3).
- ExposureLevel enum: 9 variants -> 7 (add direction/magnitude/from_dir_mag)
- FactoredAction: 81 -> 63 total actions, from_index/to_index updated
- DQN epsilon-greedy: use from_dir_mag() instead of dir*3+mag indexing
- DQN config: num_actions default 9 -> 7
- PPO action space: 45 -> 63 actions, action masking updated
- Signal adapter CUDA kernel: 5-bin -> 7-bin exposure aggregation
- All tests updated for new variant names and index ranges
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Fuses GEMM + bias-add + ReLU into a single cublasLtMatmul kernel via
CUBLASLT_EPILOGUE_RELU_BIAS. Eliminates 7 separate add_bias_relu
kernel launches per forward pass (3 trunk + 4 branch hidden layers).
Creates separate cached descriptors with epilogue enabled at init time.
Falls back to separate kernels if the epilogue heuristic isn't available.
Bias pointer set dynamically per-call via set_matmul_desc_attribute.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Performance:
- Cache cuBLAS descriptors: pre-create matmul_desc + layouts + algo at init.
requestedAlgoCount=3 for better algorithm selection. Zero per-GEMM overhead.
- Per-branch workspace: 4 × 32MB separate workspace buffers for multi-stream
branch dispatch. Eliminates workspace contention on parallel execution.
Training stability:
- Remove hardcoded shrink-perturb that fired every epoch on short runs (3-5 epochs).
With epochs=5, interval = epochs/4 = 1 → fired every epoch, destroying epoch 1
learned weights. This caused Sharpe to collapse from +0.60 to -0.29 after epoch 1.
- Phase 3 shrink-perturb now uses config values (was hardcoded alpha=0.9, sigma=0.01).
- The config-defined shrink_perturb_interval=20 now controls all shrink-perturb timing.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The backward GEMM heuristic on H100 with CUDA 13.0 requires 32MB
workspace for TF32 algorithm selection. 16MB causes
CUBLAS_STATUS_NOT_SUPPORTED on dW_only (m=256,n=128,k=16384).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
IQN was recomputing target h_s2 via iqn_trunk_forward_kernel using the
same target weights on the same next_states that DQN Pass 2 already
computed into tg_h_s2_buf. Now IQN reads from the pre-computed buffer
via a DtoD copy (negligible) instead of re-running 2 cuBLAS GEMMs.
Expected: -15ms GPU time on H100 (batch=16384, shared_h2=256).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Events recorded during CUDA graph capture cannot be queried with
cuEventElapsedTime after graph replay — the graph creates internal
copies. Reverted to original mega-graph capture structure.
Sub-graph timing fields kept in PhaseEvents for future use with
split-graph diagnostic mode or nsys profiling.
Also kept submit_loss_and_grad_ops() extraction and pub(crate) visibility
on launch_cublas_forward/backward for future per-phase graph splitting.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Split fwd_bwd timing into 5 sub-phases: spectral, forward, loss,
backward, aux. Events recorded during graph capture replay with the
graph on every step.
Also extract submit_loss_and_grad_ops() from submit_forward_ops_main()
and make launch_cublas_forward/backward pub(crate) for sub-graph timing.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
CI builder uses CUDA 13.0 cudarc which exports _v2 suffix.
Local CUDA 12.x has both names but CI only resolves _v2.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: cuStreamSynchronize blocked the thread, preventing tokio
single-threaded runtime from progressing RwLock .write().await calls.
- Remove cuStreamSynchronize from training loop (stalls pipeline)
- Move log_phase_timing() after process_epoch_boundary (DtoH syncs stream)
- compute_epoch_q_diagnostics takes &mut DQNAgentType param (no re-lock)
- Disable Q-value gap diagnostics (deadlock in single-threaded runtime)
- smoke_trainer helpers: always set fxcache dir
- test_no_hang_single_epoch: use fxcache direct loading
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Delete 7 dead files: prioritized_replay.rs, prioritized_replay_staleness.rs,
replay_buffer.rs, hindsight_replay.rs, rainbow_agent.rs, checkpoint.rs,
strategy_dqn_bridge.rs (-4,760 lines)
- Rewrite replay_buffer_type.rs: ReplayBufferType enum → StagedGpuBuffer struct
(GPU PER is the only replay path, no CPU fallbacks)
- Strip agent.rs to data types only (TradingState + AgentMetrics)
- Convert DQNAgentType from enum to struct wrapping RegimeConditionalDQN
(Standard variant was never constructed, 43 dead match arms removed)
- Remove dead CPU methods: store_experience, fused_post_step,
fused_post_step_no_ema, adaptive_buffer_resize, refresh_stale_per_priorities
- Add always-on CUDA event per-phase profiling (8 events, 4 phases:
upload/fwd_bwd/adam/per_update) with Drop cleanup and error checking
to identify 329ms/batch H100 bottleneck
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: weight tensors packed sequentially in the flat params buffer
had non-aligned start offsets when preceding tensors had odd element counts
(e.g. bias of 51 atoms = 204 bytes, 204 % 16 = 12). cublasLtMatmul with
CUBLAS_COMPUTE_32F_FAST_TF32 requires 16-byte aligned buffer pointers.
Fix: pad each tensor to 4-element boundary (16 bytes) in both
f32_weight_ptrs_from_base (pointer computation) and compute_total_params
(buffer allocation). Added align4() and padded_byte_offset() helpers,
fixed shrink_perturb skip range and bottleneck gradient offset.
Switched compute type: CUBLAS_COMPUTE_32F → CUBLAS_COMPUTE_32F_FAST_TF32
(forward + backward). Explicit TF32 tensor core path, required by cuBLAS
13.0 on H100 SM90.
Deleted dead bf16_weight_ptrs function.
19/19 smoke tests pass on RTX 3050.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Add last_epoch_loss metric. All assertions now compare the run's own metrics
against each other: convergence = loss stability (not divergence), overfitting =
best_oos_sharpe >= first_sharpe, OOS collapse = val_loss swing/magnitude ratio.
Walk-forward fold resets make epoch-1 loss unreliable for convergence comparison
(pre-trained weights), so we check stability instead.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- 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>
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>
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>
