Two `gpu_backtest_validation` tests were failing with bit-identical
deterministic values for 2+ months: `test_always_long_on_downtrend`
(expected negative PnL, got +0.00023627281) and
`test_multiple_windows_produce_results` (uptrend < downtrend instead
of > ).
Root cause: SP21 Phase 8.5 (2026-05-12) wired the factored 4-3-3-3
action decoder into the eval (`backtest_state_gather` + env_step),
but the test's hardcoded `constant_action_model(4, ...)` integer
literal wasn't migrated. Pre-Phase-8.5 the eval used a flat
4-action enum where `4` reportedly meant Long100; the factored
decoder now interprets `4` as:
decode_direction_4b(4, b1=3, b2=3, b3=3) = 4 / 27 = 0 = DIR_SHORT
decode_magnitude_4b(4, b1=3, b2=3, b3=3) = (4/9) % 3 = 0 = MAG_QUARTER
So the test was running Short-Quarter (-0.25 position) on the trend
fixtures. On random-walk synthetic prices with drift ±0.001 vs σ=0.01
noise per bar, the 24-step eval window has S/N ≈ 0.49 — specific
seeds can produce net-against-drift trajectories, making the actual
short-quarter PnL small but deterministic, with sign flipped relative
to test intent.
Fix: change `constant_action_model(4, ...)` → `constant_action_model(72, ...)`
in the 2 failing tests. Action 72 = dir=LONG (2) * 27 + mag=FULL (2)
* 9 + 0 + 0 — the actual "Long100" under 4-3-3-3 factoring. Both
tests now pass; no regressions on the 4 previously-passing tests.
Verification
────────────
- gpu_backtest_validation pre-fix: 4 passed, 2 failed
- gpu_backtest_validation post-fix: 6 passed, 0 failed
Out of scope for this commit (follow-up audit needed)
─────────────────────────────────────────────────────
Three other tests in the same file have the same stale `4` constant
with misleading "Always Long100" comments, but currently pass
incidentally:
- `test_always_long_on_uptrend` (line 218): asserts `total_pnl > 0`.
Currently passes BY ACCIDENT — action=4 (Short-Quarter) on seed-42's
net-down 24-bar trajectory produces +PnL, satisfying the assertion
for the wrong reason. Fixing to action=72 alone would break this
test (true Long100 on seed-42's net-down trajectory is negative);
the test needs BOTH the action fix AND a seed/window change so the
"uptrend" trajectory actually trends up over the eval window
(e.g., 250-bar window or drift=0.01).
- `test_extended_metrics_populated` (line 465) and
`test_active_model_records_trades` (line 524): assertions are
direction-agnostic (VaR/CVaR/Calmar/Omega NaN-check + CVaR≤VaR;
total_trades > 0 + win_rate range), so they pass legitimately
under whatever-direction action=4 produces. The "Long100" comments
are misleading but the tests are correctly covering their stated
behavior.
A separate audit-and-fix pass should address all three at once:
either correct the action constants + adjust seed/window to ensure
each test's named trajectory direction is statistically reliable,
or introduce a named constant (e.g., LONG100_ACTION) and helper to
prevent the same drift recurring.
Refs
────
- SP21 T2.2 Phase 8.5 commit 5694eb4df: "wire factored-action branch
sizes into closure-based eval (atomic)"
- crates/ml/src/cuda_pipeline/trade_physics.cuh: `decode_direction_4b`,
`decode_magnitude_4b` — canonical factored action decoders
- crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs:189:
`DqnBacktestConfig::from_network_dims` — sets branch_sizes (4,3,3,3)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
v8 smoke (train-96wfk, commit 5694eb4df) eval pod crashed at the
same point as v7 with CUDA_ERROR_ILLEGAL_ADDRESS despite Phase 8.5's
set_branch_sizes fix. 15-minute runtime confirmed env_step decoder
no longer crashes (8.5 worked); a later kernel still OOBed.
Localization via local compute-sanitizer (RTX 3050 Ti):
Updated gpu_backtest_validation.rs to mirror production eval-baseline
(FEATURE_DIM=42, portfolio_dim=24, set_branch_sizes), reproducing the
v8 crash in 1.66s locally.
compute-sanitizer --tool=memcheck pinpointed:
Invalid __global__ read of size 4 bytes
at cost_net_sharpe_kernel+0x90
by thread (32,0,0) in block (0,0,0)
Access at 0x65c is out of bounds
0x65c = 1628 bytes = float index 407 = OFI_IMPACT_LAMBDA_INDEX.
Kernel reads isv[407]; isv pointer was 0 (null) because
isv_signals_ptr defaults to 0 in constructor and set_isv_signals_ptr
is only called by production training. Closure-path callers
(eval-baseline) dereferenced null + slot×4 bytes.
Fix:
Allocate zero-filled default_isv_buf of size ISV_TOTAL_DIM=536 f32
in constructor. Wire isv_signals_ptr to its dev_ptr by default.
Production training still overrides via set_isv_signals_ptr.
Zero-init semantics:
- isv[407]=0 → ofi_lambda=0 → c_ofi=0 in cost-net sharpe
(degraded but valid; matches LobBar.ofi=0.0 placeholder)
- Other slots default to 0 — Kelly health, controller anchors,
etc. all see degraded-but-valid defaults
Test updates (consumer migration):
- FEATURE_DIM 10 → 42 (production value; FEATURE_DIM < 32 makes
gather kernel's `market_dim = feat_dim - SL_OFI_DIM (32)` negative
→ OOB; previous tests were already broken even before our changes)
- portfolio_dim 3 → 24 (Phase 8.3+9 contract)
- set_branch_sizes call added (Phase 8.5 contract)
Verification:
cargo test -p ml --test gpu_backtest_validation \
gpu_tests::test_always_long_on_uptrend --features cuda --release
# PASSES
compute-sanitizer --tool=memcheck <test_bin>
# 0 CUDA errors across all 6 tests
# (2 PnL-assertion test failures are pre-existing data-expectation
# issues with new FEATURE_DIM=42, not OOB bugs)
Pearls honoured:
- feedback_no_hiding: null pointer surfaced via compute-sanitizer
instead of silently crashing in production
- feedback_no_partial_refactor: closure-path callers now have
self-contained ISV setup matching training path; consumer
migration (test FEATURE_DIM/portfolio_dim/set_branch_sizes)
atomic with the producer-side default ISV alloc
- pearl_no_deferrals_for_complementary_fixes: same SP cycle as
8.3+9, 8.5 (the layer-by-layer eval rot peeling)
Co-Authored-By: Claude Opus 4.7 (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>
8 test files had stale types from the bf16→f32 conversion:
- gpu_smoketest: missing adam_epsilon in DQNConfig
- gpu_backtest_validation: closure params bf16→f32
- gpu_kernel_parity_test: market data, weight readback bf16→f32
- gpu_per_integration_test: weights readback bf16→f32
- target_update_tests: varstore register bf16→f32
- smoke_test_real_data: market buffers bf16→f32
- activation_tests, dropout_scheduler_tests: forward() signature change
These tests only compile with --features cuda (CI path), which is why
they passed locally with cargo test --lib.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Final cleanup:
- 61 test files + 5 example files: candle imports replaced
- 8 testing/integration files: migrated to cudarc/ml-core types
- 3 services/trading_service test files: migrated
- Root Cargo.toml: candle-core, candle-nn removed from [workspace.dependencies]
- crates/ml/Cargo.toml: candle-nn dependency removed
- testing/e2e/Cargo.toml: candle-core dependency removed
Zero active candle_core/candle_nn/candle_optimisers code references remain.
Zero candle dependency declarations in any Cargo.toml.
Remaining "candle" strings are exclusively in doc comments.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
- Remove ALL #[cfg(feature = "cuda")] guards (~400+ occurrences)
- Remove ALL #[cfg_attr(not(feature = "cuda"), ignore)] test annotations (~250)
- Make cuda default feature in 9 ML crates (ml, ml-core, ml-dqn, ml-ppo, etc.)
- Convert nvrtc JIT compilation to precompiled nvcc (searchsorted, prefix_sum)
- Move compile_ptx_for_device() to ml-core for shared access
- Delete dead CPU code: multi_step.rs, self_supervised_pretraining.rs,
training_guard_gpu_tests.rs, CPU PER buffer paths, CPU Q-diagnostics
- Replace unwrap_or(Device::Cpu) with hard errors everywhere
- Remove dead is_cuda() else branches in DQN/PPO/hyperopt trainers
- Change config defaults from "cpu" to "cuda" (rainbow, tlob, pipeline)
- Port IQL value network to GPU kernel (5 CUDA entry points)
- Port HER goal relabeling to GPU kernel (warp-per-sample)
- Wire DSR GPU-to-CPU sync in training loop
- cfg!(feature = "cuda") → true in inference_validator
Zero warnings, zero errors across entire workspace.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Adds crates/ml/tests/gpu_backtest_validation.rs with 6 deterministic
tests that validate GpuBacktestEvaluator produces reasonable metrics:
always-long on uptrend (positive PnL), always-long on downtrend
(negative PnL), always-flat (~zero PnL), multi-window ordering,
extended metrics finiteness/self-consistency, and trade count.
All tests are #[ignore] gated and skip gracefully on CPU-only machines
(CI passes with 6 ignored; intended for GPU development machines).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
