disable_event_tracking() was called globally in GpuDqnTrainer::new(),
polluting the CUDA context for ALL subsequent operations including
other models (TFT, PPO). This caused test_tft_adapter_deterministic
to fail when run after DQN tests.
Fix: disable/enable only inside capture_training_graph(), not at
construction. The capture region is the only place where CudaEvents
are disallowed.
Result: 855/855 ml lib tests pass (was 854/855 with 1 flaky failure).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cuCtxSetLimit(STACK_SIZE) reserves stack_bytes × max_threads of VRAM
upfront. On 4GB GPUs, multiple cuCtxSetLimit calls (experience: 8KB,
curiosity: 16KB, training: 64KB) caused CUDA_ERROR_OUT_OF_MEMORY.
Fix: set stack ONCE to 64KB in DQNTrainer::new_internal() — before
any kernel initialization. All subsequent kernels (experience,
curiosity, training) share this limit.
Removed cuCtxSetLimit from:
- GpuExperienceCollector::new() (was 8KB)
- GpuCuriosityTrainer::new() (was 16KB)
- GpuDqnTrainer::new() (was 64KB, moved to constructor)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The fused curiosity kernel (curiosity_fused_zero_fwd_bwd_adam) used
__threadfence() + __syncthreads() + atomicAdd block counter for
inter-block synchronization. This crashed on RTX 3050 (and potentially
other consumer GPUs) — the synchronize() after launch deadlocked.
Fix: use separate kernels instead:
1. curiosity_shift_states (shift next_states)
2. curiosity_forward_backward (accumulate gradients via atomicAdd)
3. curiosity_adam_step (4 launches, one per param group)
This adds 4 kernel launches (~5μs overhead) but eliminates the
fragile inter-block sync pattern. Gradients are zeroed via
memset_zeros before forward_backward (GPU-side, stream-ordered).
Curiosity re-enabled in smoke test. Full e2e DQN training passes
on RTX 3050 in 0.67s with branching + C51(11) + curiosity + NoisyNets.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Full branching DQN training pipeline validated on consumer GPU:
- 2 epochs × 8 training steps × 32 batch
- Branching DQN with C51 (11 atoms, dynamic scaling)
- GPU experience collection (2 episodes × 10 timesteps)
- Fused CUDA Graph training (forward+loss+backward+Adam)
- GPU PER insert + monitoring reduce
- All GPU, zero CPU roundtrips in hot path
Curiosity disabled for smoke test — the curiosity CUDA kernel
still crashes on RTX 3050 (needs separate investigation).
Removed all debug eprintln! traces from training_loop.rs.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
ROOT CAUSE FOUND AND FIXED: The fused training kernel allocates
DIST_SIZE(HIDDEN_DIM) * NUM_ATOMS per-thread local arrays:
- With NUM_ATOMS=51: ~46KB/thread → stack overflow on ALL GPUs
- With NUM_ATOMS=11: ~12KB/thread → fits with 64KB stack limit
Three critical fixes:
1. cuCtxSetLimit(STACK_SIZE, 65536) in GpuDqnTrainer::new() — the
forward+loss kernel needs up to 46KB/thread for distributional
arrays. Default 1KB causes ILLEGAL_ADDRESS from __local__ overflow.
2. disable_event_tracking() in GpuDqnTrainer::new() — CudaEvents are
disallowed inside CUDA Graph capture. Without this, cudarc's
device_ptr/device_ptr_mut record events that cause
STREAM_CAPTURE_INVALIDATED or INVALID_VALUE on stale event refs.
3. Dynamic C51 atom scaling in smoke test: <8GB VRAM → 11 atoms,
<16GB → 21 atoms, >=16GB → 51 atoms. Prevents stack overflow
on consumer GPUs (RTX 3050: 4GB) while keeping full C51 on H100.
Verified: GPU at 100% utilization, training loop running successfully
on RTX 3050 with NUM_ATOMS=11 and 64KB stack. No more ILLEGAL_ADDRESS.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three fixes:
1. disable_event_tracking() before CUDA Graph capture, enable after.
cudarc's device_ptr/device_ptr_mut record CudaEvents which are
disallowed inside graph capture (STREAM_CAPTURE_INVALIDATED).
2. Replaced raw_device_ptr + dtod_copy scalar gather with direct
stream.memcpy_dtoh from each 1-element GPU buffer. Eliminates
the ManuallyDrop guard leak anti-pattern from the readback path.
3. curiosity kernel: cuCtxSetLimit(STACK_SIZE, 8192) + removed
in-kernel gradient zeroing (inter-block race), replaced with
host-side memset_zeros (GPU cuMemsetD8Async, stream-ordered).
Remaining: CUDA_ERROR_ILLEGAL_ADDRESS in fused forward/loss kernel
during graph replay — the captured training kernel has a buffer
overrun. Needs investigation of submit_training_ops() kernel args.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
1. cuCtxSetLimit(STACK_SIZE, 8192) in curiosity trainer — prevents
stack overflow in fused kernel (6 arrays of 42-128 floats/thread)
2. Removed in-kernel gradient zeroing (Phase 1) — had inter-block race
where fast blocks atomicAdd while slow blocks still zero. Now uses
host-side memset_zeros (GPU cuMemsetD8Async, stream-ordered)
3. cuda_slice_to_tensor_f32 now uses safe stream.memcpy_dtod() instead
of raw device_ptr + memcpy_dtod_async (cudarc event tracking fix)
4. Debug traces in smoke test and training loop for deadlock diagnosis
Investigation ongoing: deadlock in init_gpu_experience_collector —
the collector constructor hangs during initialization, not during
experience collection or training.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Zero candle_core, candle_nn, or candle_optimisers references remain
in the entire ml crate source code. Workspace compiles clean with
0 errors and 0 clippy warnings.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three fixes:
1. count_cuda_devices() uses CudaContext::device_count() instead of
probing with CudaContext::new(i) — failed probes on non-existent
devices pollute cudarc's error_state and invalidate CUDA Graph
capture for the fused training path.
2. Curiosity sync guard: stream.synchronize() after curiosity
training with auto-disable if async error detected. Prevents
poisoned stream from deadlocking the PER insert path.
3. Smoke test sets curiosity_weight=0.0 temporarily — the curiosity
CUDA kernel has a stack/memory issue that needs separate
investigation. Disabling it unblocks the full training pipeline.
4. Removed unused DevicePtr/DevicePtrMut imports after memcpy fix.
Remaining: CUDA_ERROR_STREAM_CAPTURE_INVALIDATED during fused
training graph capture — a prior error invalidates the capture.
This is NOT a deadlock (test completes in 8s), just needs the
capture error source identified and fixed.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Three fixes:
1. FusedTrainingCtx::new now accepts RegimeConditionalDQN by using
primary_head() — the old code rejected it, causing silent fallback
to non-fused train_step() which has action-space mismatch with
branching DQN (5-action Q-table vs 45-action factored indices →
CUDA_ERROR_ILLEGAL_ADDRESS buffer overrun)
2. ensure_fused_ctx() returns Result instead of () — fused init
failure is now a hard error (no silent CPU fallback)
3. GpuTensor::cat now supports dim>0 concatenation (was unimplemented,
caused "dim=1 > 0 not yet implemented" error in validation path)
Root cause chain:
RegimeConditional rejected by fused init
→ silent fallback to DQN::train_step()
→ compute_loss_internal() uses num_actions=5 but actions are 0-44
→ gather with out-of-bounds offsets → CUDA_ERROR_ILLEGAL_ADDRESS
→ async error poisons CUDA context
→ next stream.synchronize() deadlocks forever
Remaining: curiosity kernel crash also poisons the stream. The
curiosity training runs inside collect_gpu_experiences() and its
async error blocks the PER insert. This needs separate investigation
of curiosity_training_kernel.cu.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Fixes cudarc 0.19 event tracking corruption caused by mixing safe
device_ptr/device_ptr_mut wrappers with raw memcpy_dtod_async.
The anti-pattern: manually calling device_ptr() to get raw pointers,
then using low-level memcpy_dtod_async, then dropping SyncOnDrop
guards — this bypasses cudarc's event management and corrupts the
synchronization state of CudaSlice objects.
Fixed in 7 call sites across 3 files:
- gpu_experience_collector.rs: dtod_clone_f32, dtod_clone_i32,
build_next_states_dtod (replaced pointer arithmetic with slice views)
- gpu_weights.rs: extract_one, sync_one
- training_loop.rs: cuda_slice_to_tensor_f32
Investigation ongoing: deadlock persists in PER insert path
(cuda_slice_to_tensor_f32 -> stream.synchronize()). The memcpy fix
is correct but there's an additional issue in the CudaSlice->GpuTensor
conversion that needs further debugging.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Reverts 4 commits (8139911c, 282f3aff, b3aca96d, dad61e4e) that
tried to workaround slow CI by limiting data subsets and cleaning
caches. The real issue is debug-mode .dbn.zst parsing taking minutes.
Kept: --test-threads=1 fix (root cause), hard CUDA error, action
range fixes, #[ignore] for heavy tests.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
DQNHyperparameters.max_bars caps total bars loaded from .dbn files.
CI smoke test now loads 2000 bars (not 600K) — validates the full
pipeline in seconds instead of minutes of I/O.
Default: 0 (unlimited, for production training).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
cuda_if_available() silently fell back to CPU when CUDA init failed,
causing DQNTrainer to crash later with "cuda_stream() called on CPU
device". This was the root cause of the CI "hang" — the test failed
immediately but the error was invisible due to buffered output.
- DQNTrainer::new() now uses MlDevice::cuda(0)? with explicit error
- cuda_if_available() now logs the CUDA error before falling back
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Tests that load full 163K-bar training datasets take 30+ min in CI.
Mark them #[ignore] — run via nightly cron or manual trigger.
Affected: gpu_residency(2), training_stability(2), performance(2),
feature_coverage(1), ppo_benchmark(1), cache(1)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Final 31 ml crate fixes: unsafe_code allows, unused vars prefixed,
boolean simplification, dead code removal, integer suffix, drop cleanup.
cargo fix auto-removed ~30 unused imports from ml crate.
Total clippy cleanup: 278 errors → 0 across all ML crates.
Full workspace: `cargo clippy --workspace --lib -- -D warnings` = 0 errors.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Removed dead code after early return in liquid round-trip test.
Checkpoint is implemented — removed "Skip weight equality check
until checkpoint is implemented" comment.
12/12 lightweight smoke tests pass. 6 heavy tests need >4GB VRAM (H100).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
gpu_portfolio.rs: deleted CPU simulate_batch() path (GPU path exists),
get_portfolio_state() returns &CudaSlice (was: download to [f32;8])
branching.rs: GPU-native argmax per branch, GPU gather+mean for Q
aggregation, GPU affine+floor for action decomposition, GPU sub→abs→
max for weight comparison in tests
gpu_tensor.rs: added pub cuda_data() accessor
stream_ops.rs: fixed CudaView type mismatch in dtod_copy
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
apply_static_context() tried gpu_add([1,9,32], [1,1,32]) — shape
mismatch. Static context is time-invariant, must broadcast across
sequence dimension before adding to temporal features.
Both TFT adapter tests now pass. Zero ignored in ensemble adapters.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
test_model_registry_new and test_register_and_retrieve_model no longer
ignored. foxhunt-postgres docker container available locally.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
test_optimization_rosenbrock and test_optimization_deterministic were
marked #[ignore] with no valid reason. They complete in <50ms.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Removed #[ignore] from tests that have local infrastructure:
- 3 data_loader tests: auto-detect test_data/real/databento/ via workspace
- 3 memory_profiler tests: nvidia-smi at /usr/bin/nvidia-smi
- 4 benchmark tests (TFT, Mamba2, DQN, PPO): GPU + DBN data available
- 1 inference test: model loading (slow but should run)
- 3 DQN performance smoke tests: GPU available
PPO benchmark: fixed data_path to test_data/real/databento/6E.FUT
Sequential: added vars_mut() accessor
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
PpoExperienceBatch: all 6 fields Vec<T>→CudaSlice<T>. Data stays on GPU
from kernel collection through training. Zero DtoH for batch data.
- collect_experiences(): returns CudaSlice via D2D copy, no memcpy_dtoh
- gpu_batch_to_trajectory_batch(): DELETED (was CPU conversion)
- PPO::update_gpu(): reads states via D2D mini-batch, forward on GPU
- compute_metrics_from_gpu(): downloads only scalars (returns/advantages/
actions), NOT the 123 MB state tensor
- download_*() methods for debug/checkpoint only
States (123 MB/epoch) never leave GPU. Only 2 scalar losses come to CPU.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Deleted ExperienceBatch (CPU Vec<f32>) struct and collect_experiences()
method (120 lines of memcpy_dtoh + CPU next_state computation).
Production DQN training already uses collect_experiences_gpu() which
returns GpuExperienceBatch (CudaSlice, zero CPU download).
Tests rewritten as pure index-math validation without GPU types.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
New ml-core/cuda_autograd/reductions.rs:
- fused_stats_reduce: min/max/sum/sum_sq/count in single pass with
warp-level __shfl_down_sync + atomicCAS (20 bytes DtoH)
- argmax_flat: single u32 result (4 bytes DtoH)
- argmax_rows: per-row argmax for 2D data
- sum_reduce: standard tree reduction
- col_sum_reduce: per-column sum along rows
DQN trainer integration:
- collect_qvalue_statistics: 4 DtoH → 1 (single stats() call)
- compute_q_diagnostics_fused: replaces Candle sort/narrow pipeline
with argmax_rows + stats + col_sums (3 kernels, 3 readbacks)
- ReductionKernels lazy-initialized in training loop
All kernels compiled via compile_ptx_for_device (native cubin + disk cache).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Experience collector: PinnedHostBuf<T> via cuMemHostAlloc(PORTABLE) for
6 DtoH buffers — ~25-30% faster transfers vs pageable memory.
NVRTC caching: 4 OnceLock additions in ml-ppo cuda_nn:
- softmax.rs: was re-compiling NVRTC on EVERY forward pass batch (!)
- linear.rs, lstm.rs: cached for multi-layer construction
Device attribute: OnceLock<u32> for max_threads in gpu_replay_buffer
pfx_sum — eliminates ~5µs driver query per call.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
reset_episodes() was allocating 5 CPU Vecs per epoch for HtoD zero-fill.
Replaced with:
- memset_zeros() for barrier/diversity buffers (async GPU-side, no HtoD)
- Pre-allocated host staging vectors for portfolio_init and RNG seeds
gpu_curiosity_trainer already optimized (fused kernel zeros internally).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
