docs: CUDA backtest & GPU-resident training design spec
Three-phase plan to eliminate all GPU→CPU roundtrips from training: - Phase 1: Seal training loop (persistent GPU epoch state, async monitoring) - Phase 2: Vectorized CUDA backtest kernel for hyperopt evaluation - Phase 3: General-purpose GPU backtester replacing CPU SIMD path Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
434
docs/plans/2026-03-11-cuda-backtest-gpu-residency.md
Normal file
434
docs/plans/2026-03-11-cuda-backtest-gpu-residency.md
Normal file
@@ -0,0 +1,434 @@
|
||||
# CUDA Backtest & GPU-Resident Training — Design Spec
|
||||
|
||||
> **For agentic workers:** REQUIRED: Use superpowers:subagent-driven-development (if subagents available) or superpowers:executing-plans to implement this plan. Steps use checkbox (`- [ ]`) syntax for tracking.
|
||||
|
||||
**Goal:** Eliminate all GPU→CPU roundtrips from the DQN/PPO training loop and port the walk-forward backtesting engine to CUDA, enabling fully GPU-resident training and evaluation.
|
||||
|
||||
**Architecture:** Three-phase approach — (1) seal remaining CPU roundtrips in the training loop, (2) build a vectorized GPU backtest environment for hyperopt evaluation, (3) extend to a general-purpose GPU backtester replacing the CPU SIMD path.
|
||||
|
||||
**Tech Stack:** Rust, cudarc 0.17, NVRTC, Candle (tensor ops + model forward), CUDA C (custom kernels)
|
||||
|
||||
---
|
||||
|
||||
## Context & Motivation
|
||||
|
||||
The DQN training loop currently has ~80% GPU residency:
|
||||
|
||||
- `GpuExperienceCollector` — full episode kernel (action selection, portfolio sim, barriers, fill sim, reward, TD error)
|
||||
- `collect_experiences_gpu()` → `insert_batch_tensors()` — zero-roundtrip DtoD path to GPU replay buffer
|
||||
- `GpuReplayBuffer` — GPU-resident PER sampling (proportional + rank-based)
|
||||
- `GpuTrainingGuard` — on-device NaN/loss-clip/grad-collapse checks
|
||||
|
||||
**Remaining CPU roundtrips to eliminate:**
|
||||
|
||||
1. **Training step readbacks**: `to_scalar` for loss, `to_vec1` for Q-values, gradient norms
|
||||
2. **Monitoring downloads**: rewards_cpu/actions_cpu from experience kernel (per kernel launch)
|
||||
3. **Epoch-boundary state**: vol EMA, portfolio reset/compound, DSR normalizer reset — force `cudaStreamSync`
|
||||
4. **Hyperopt evaluation**: trained model weights downloaded → CPU walk-forward backtest → metric uploaded
|
||||
5. **Standalone evaluation**: `evaluate_baseline` binary runs CPU backtesting with SIMD AVX2
|
||||
|
||||
---
|
||||
|
||||
## Phase 1: Seal the Training Loop
|
||||
|
||||
Close remaining GPU→CPU gaps in `crates/ml/src/trainers/dqn/trainer.rs` and CUDA pipeline.
|
||||
|
||||
### 1.1 Eliminate train_step readbacks
|
||||
|
||||
**Loss scalar** (`trainer.rs:4797`):
|
||||
- `GpuTrainingGuard` already checks loss on-device via pinned host memory
|
||||
- Defer `to_scalar` readback: accumulate losses on GPU, single async readback at epoch end
|
||||
- Replace per-step loss with GPU-side running mean (EMA in training guard kernel)
|
||||
|
||||
**Q-value monitoring** (`trainer.rs:4834,5495`):
|
||||
- `gpu_qvalue_stats` kernel already computes mean/std/min/max on device
|
||||
- Drop the CPU `to_vec1` → `to_vec2` fallback paths entirely when `GpuTrainingGuard` is active
|
||||
- Q-value divergence check: already has GPU kernel path — make it exclusive on CUDA
|
||||
|
||||
**Gradient norm** (`trainer.rs:5419`):
|
||||
- Candle computes grad norm during `optimizer.step()` — stays on GPU
|
||||
- Current `to_scalar` readback is for logging only
|
||||
- Fix: write to pinned host memory via `cuMemHostAlloc`, read asynchronously at epoch end
|
||||
|
||||
### 1.2 Batch monitoring downloads
|
||||
|
||||
**Current**: `collect_experiences_gpu()` downloads `rewards_cpu` and `actions_cpu` after every kernel launch for monitoring (mean reward, Sharpe, action diversity).
|
||||
|
||||
**Fix**: Accumulate monitoring stats on GPU across all kernel launches within an epoch. Single `memcpy_dtoh` of summary struct at epoch end:
|
||||
```c
|
||||
struct EpochMonitoringSummary {
|
||||
float mean_reward;
|
||||
float reward_std;
|
||||
float sharpe_estimate; // mean/std * sqrt(steps)
|
||||
int action_counts[5]; // per-exposure-level counts
|
||||
float max_reward;
|
||||
float min_reward;
|
||||
int total_experiences;
|
||||
};
|
||||
```
|
||||
|
||||
Add a small `monitoring_reduction_kernel` that reduces per-experience rewards/actions into this summary. Launch once at epoch end. Download 48 bytes instead of `2 * N_episodes * timesteps * 4` bytes per kernel launch.
|
||||
|
||||
### 1.3 GPU-persistent epoch-boundary state
|
||||
|
||||
Move epoch-boundary state management from CPU to persistent GPU buffers.
|
||||
|
||||
**Persistent `CudaSlice<f32>` buffers on `GpuExperienceCollector`:**
|
||||
|
||||
| Buffer | Size | Purpose |
|
||||
|--------|------|---------|
|
||||
| `epoch_state` | 8 × f32 | vol_ema, median_vol, portfolio_value, portfolio_position, portfolio_cash, dsr_mean, dsr_var, step_count |
|
||||
|
||||
**Kernel changes:**
|
||||
- Last step of experience kernel writes final state to `epoch_state` buffer
|
||||
- Next epoch's kernel launch reads `epoch_state` as initial state
|
||||
- New kernel arg `reset_flags: u32` (bitfield): bit 0 = reset portfolio (DSR mode), bit 1 = reset DSR normalizer, bit 2 = reset vol EMA
|
||||
|
||||
**Result**: Zero `cudaStreamSynchronize` between epochs. The only CPU→GPU communication is updating kernel config args (epsilon decay, learning rate) which are cheap scalar copies.
|
||||
|
||||
### 1.4 Remove CPU fallback codepaths (CUDA builds)
|
||||
|
||||
When compiled with `feature = "cuda"` and running on a CUDA device:
|
||||
- Remove `#[cfg(not(feature = "cuda"))]` branches from the hot path in `train_epoch()`
|
||||
- The CPU fallback in the experience collection loop (`if !gpu_experiences_collected`) should be unreachable when GPU collector is initialized
|
||||
- Add `debug_assert!` guards confirming GPU path was taken
|
||||
- Keep CPU paths for `feature = "cpu-only"` builds and test harness
|
||||
|
||||
### 1.5 Expected outcome
|
||||
|
||||
| Metric | Before | After |
|
||||
|--------|--------|-------|
|
||||
| `cudaStreamSync` per epoch | ~8-16 (per batch readback) | 1 (epoch-end monitoring) |
|
||||
| CPU roundtrips per experience batch | 2 (rewards + actions download) | 0 |
|
||||
| CPU roundtrips per train step | 3 (loss + Q-stats + grad norm) | 0 |
|
||||
| Epoch boundary sync | 1 (portfolio/vol state) | 0 |
|
||||
|
||||
---
|
||||
|
||||
## Phase 2: CUDA Backtest Kernel for Hyperopt Evaluation
|
||||
|
||||
### 2.1 Problem statement
|
||||
|
||||
Hyperopt evaluation loop:
|
||||
```
|
||||
for trial in 0..N_trials:
|
||||
train model on GPU (epochs)
|
||||
for fold in walk_forward_windows:
|
||||
download model weights to CPU # SYNC
|
||||
run CPU backtest on test window # SLOW
|
||||
compute Sharpe/PnL on CPU # SLOW
|
||||
upload aggregated metric to optimizer # SYNC
|
||||
```
|
||||
|
||||
On H100 with 20 hyperopt trials × 8 walk-forward folds × ~100K bars per window:
|
||||
- CPU backtest: ~2-5s per fold × 8 folds = 16-40s per trial
|
||||
- GPU training: ~30s per trial
|
||||
- **Evaluation is 30-60% of total hyperopt time**
|
||||
|
||||
### 2.2 Architecture: Vectorized GPU Environment
|
||||
|
||||
Pattern: NVIDIA Isaac Gym / Google Brax style — parallelize across environments (walk-forward windows), sequential within each.
|
||||
|
||||
```
|
||||
┌─────────────────────────────────────────────────────┐
|
||||
│ GPU Memory │
|
||||
│ │
|
||||
│ ┌──────────────────────────────────────────────┐ │
|
||||
│ │ Market Data [N_windows × max_len × feat_dim] │ │
|
||||
│ │ (uploaded once, read-only, SoA layout) │ │
|
||||
│ └──────────────────────────────────────────────┘ │
|
||||
│ │
|
||||
│ for step in 0..max_steps: │
|
||||
│ ┌────────────────────────────────────────────┐ │
|
||||
│ │ 1. Gather states [cudarc kernel] │ │
|
||||
│ │ states[w] = features[w][step] ++ port[w]│ │
|
||||
│ └──────────────┬─────────────────────────────┘ │
|
||||
│ ▼ │
|
||||
│ ┌────────────────────────────────────────────┐ │
|
||||
│ │ 2. Batch forward [Candle, on-device] │ │
|
||||
│ │ q_values = model.forward(states_batch) │ │
|
||||
│ └──────────────┬─────────────────────────────┘ │
|
||||
│ ▼ │
|
||||
│ ┌────────────────────────────────────────────┐ │
|
||||
│ │ 3. Action select [cudarc kernel] │ │
|
||||
│ │ actions = argmax(q_values, dim=-1) │ │
|
||||
│ └──────────────┬─────────────────────────────┘ │
|
||||
│ ▼ │
|
||||
│ ┌────────────────────────────────────────────┐ │
|
||||
│ │ 4. Env step [cudarc kernel] │ │
|
||||
│ │ For each window w (parallel): │ │
|
||||
│ │ execute_trade(actions[w], port[w]) │ │
|
||||
│ │ compute_reward(port[w], prices[w]) │ │
|
||||
│ │ update portfolio state │ │
|
||||
│ │ check done (window exhausted) │ │
|
||||
│ └──────────────┬─────────────────────────────┘ │
|
||||
│ ▼ │
|
||||
│ end for │
|
||||
│ │
|
||||
│ ┌──────────────────────────────────────────────┐ │
|
||||
│ │ 5. Metrics reduction [cudarc kernel] │ │
|
||||
│ │ Per window: Sharpe, total_pnl, max_dd │ │
|
||||
│ └──────────────┬───────────────────────────────┘ │
|
||||
│ ▼ │
|
||||
│ Single memcpy_dtoh: [N_windows × 3] floats │
|
||||
└─────────────────────────────────────────────────────┘
|
||||
```
|
||||
|
||||
### 2.3 Kernel designs
|
||||
|
||||
#### `backtest_env_kernel.cu` — Vectorized environment step
|
||||
|
||||
```c
|
||||
// One thread per walk-forward window
|
||||
// Each thread steps sequentially through its window
|
||||
__global__ void backtest_env_step(
|
||||
// Market data (read-only, SoA)
|
||||
const float* __restrict__ features, // [N_windows, max_len, feat_dim]
|
||||
const float* __restrict__ prices, // [N_windows, max_len, 4] (open, high, low, close)
|
||||
const int* __restrict__ window_lens, // [N_windows] actual length per window
|
||||
|
||||
// Actions from model (read-only)
|
||||
const int* __restrict__ actions, // [N_windows] current step action
|
||||
|
||||
// Portfolio state (read-write, persistent across steps)
|
||||
float* portfolio_state, // [N_windows, PORTFOLIO_STATE_SIZE]
|
||||
// Layout per window: [value, position, cash, entry_price, max_equity, step_pnl, cum_return, step_count]
|
||||
|
||||
// Step rewards (write)
|
||||
float* step_rewards, // [N_windows]
|
||||
int* done_flags, // [N_windows]
|
||||
|
||||
// Config
|
||||
float max_position,
|
||||
float spread_cost,
|
||||
float tx_cost_bps,
|
||||
int current_step
|
||||
) {
|
||||
int w = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (w >= N_WINDOWS) return;
|
||||
if (done_flags[w]) return; // already done
|
||||
if (current_step >= window_lens[w]) { done_flags[w] = 1; return; }
|
||||
|
||||
// Read current price
|
||||
int price_idx = w * max_len * 4 + current_step * 4;
|
||||
float close = prices[price_idx + 3];
|
||||
|
||||
// Read portfolio state
|
||||
int ps_idx = w * PORTFOLIO_STATE_SIZE;
|
||||
float value = portfolio_state[ps_idx + 0];
|
||||
float position = portfolio_state[ps_idx + 1];
|
||||
float cash = portfolio_state[ps_idx + 2];
|
||||
// ... (execute trade, compute reward, update state)
|
||||
|
||||
// Write back
|
||||
portfolio_state[ps_idx + 0] = new_value;
|
||||
step_rewards[w] = step_return;
|
||||
}
|
||||
```
|
||||
|
||||
#### `backtest_metrics_kernel.cu` — Per-window reduction
|
||||
|
||||
```c
|
||||
// Reduction kernel: compute Sharpe, total PnL, max drawdown per window
|
||||
// Uses warp-level primitives for efficient reduction
|
||||
__global__ void compute_backtest_metrics(
|
||||
const float* step_returns, // [N_windows, max_len]
|
||||
const int* window_lens, // [N_windows]
|
||||
float* metrics_out, // [N_windows, 3]: (sharpe, total_pnl, max_drawdown)
|
||||
int max_len
|
||||
) {
|
||||
int w = blockIdx.x;
|
||||
// Thread-parallel reduction over step_returns[w][0..window_lens[w]]
|
||||
// Compute: mean, variance (Welford's online), cumsum for drawdown
|
||||
// Write: sharpe = mean/sqrt(var) * sqrt(252), total_pnl, max_dd
|
||||
}
|
||||
```
|
||||
|
||||
### 2.4 Rust orchestrator: `GpuBacktestEvaluator`
|
||||
|
||||
```rust
|
||||
pub struct GpuBacktestEvaluator {
|
||||
stream: Arc<CudaStream>,
|
||||
env_kernel: CudaFunction,
|
||||
metrics_kernel: CudaFunction,
|
||||
// Persistent GPU buffers
|
||||
features_buf: CudaSlice<f32>, // [N_windows, max_len, feat_dim]
|
||||
prices_buf: CudaSlice<f32>, // [N_windows, max_len, 4]
|
||||
window_lens_buf: CudaSlice<i32>, // [N_windows]
|
||||
portfolio_buf: CudaSlice<f32>, // [N_windows, PORTFOLIO_STATE_SIZE]
|
||||
rewards_buf: CudaSlice<f32>, // [N_windows, max_len]
|
||||
done_buf: CudaSlice<i32>, // [N_windows]
|
||||
actions_buf: CudaSlice<i32>, // [N_windows]
|
||||
metrics_buf: CudaSlice<f32>, // [N_windows, 3]
|
||||
// Config
|
||||
n_windows: usize,
|
||||
max_window_len: usize,
|
||||
feature_dim: usize,
|
||||
}
|
||||
|
||||
impl GpuBacktestEvaluator {
|
||||
/// Upload walk-forward test windows to GPU (called once per hyperopt trial).
|
||||
pub fn upload_windows(
|
||||
windows: &[WalkForwardWindow],
|
||||
feature_extractor: &impl FeatureExtractor,
|
||||
device: &Device,
|
||||
) -> Result<Self, MLError>;
|
||||
|
||||
/// Run full evaluation: step loop with model forward + env kernel.
|
||||
/// Returns per-window metrics without leaving GPU until final readback.
|
||||
pub fn evaluate(
|
||||
&mut self,
|
||||
model: &dyn ModelForward, // Candle model with .forward(Tensor) -> Tensor
|
||||
device: &Device,
|
||||
) -> Result<Vec<WindowMetrics>, MLError>;
|
||||
}
|
||||
|
||||
pub struct WindowMetrics {
|
||||
pub sharpe: f32,
|
||||
pub total_pnl: f32,
|
||||
pub max_drawdown: f32,
|
||||
}
|
||||
```
|
||||
|
||||
### 2.5 Integration with hyperopt adapters
|
||||
|
||||
Each hyperopt adapter's `evaluate()` method gains a GPU path:
|
||||
|
||||
```rust
|
||||
// In hyperopt/adapters/dqn.rs (and ppo.rs, tft.rs, etc.)
|
||||
fn evaluate_walk_forward(&self, model: &TrainedModel, windows: &[WalkForwardWindow]) -> f64 {
|
||||
#[cfg(feature = "cuda")]
|
||||
if self.device.is_cuda() {
|
||||
// GPU path: zero roundtrips
|
||||
let mut evaluator = GpuBacktestEvaluator::upload_windows(windows, &self.feature_extractor, &self.device)?;
|
||||
let metrics = evaluator.evaluate(model, &self.device)?;
|
||||
return aggregate_sharpe(&metrics);
|
||||
}
|
||||
// CPU fallback: existing walk_forward_backtest()
|
||||
self.cpu_evaluate_walk_forward(model, windows)
|
||||
}
|
||||
```
|
||||
|
||||
### 2.6 Memory budget
|
||||
|
||||
On H100 80GB with 8 walk-forward windows × 100K bars × 48-dim features:
|
||||
|
||||
| Buffer | Size |
|
||||
|--------|------|
|
||||
| Features | 8 × 100K × 48 × 4B = 147 MB |
|
||||
| Prices | 8 × 100K × 4 × 4B = 12 MB |
|
||||
| Portfolio state | 8 × 8 × 4B = 256 B |
|
||||
| Step rewards | 8 × 100K × 4B = 3 MB |
|
||||
| Actions/done/lens | < 1 MB |
|
||||
| Metrics output | 96 B |
|
||||
| **Total** | **~163 MB** (~0.2% of H100 VRAM) |
|
||||
|
||||
Trivial. Could run 256+ parallel windows if needed for ensemble cross-validation.
|
||||
|
||||
### 2.7 Expected speedup
|
||||
|
||||
| Component | CPU (current) | GPU (target) | Speedup |
|
||||
|-----------|--------------|-------------|---------|
|
||||
| Feature gather per step | 1-5 μs | <0.1 μs (coalesced read) | 10-50× |
|
||||
| Model forward (batch of 8) | 100-500 μs | 20-50 μs (already on GPU) | 5-10× |
|
||||
| Env step (8 windows) | 8-40 μs | 0.5-2 μs (parallel threads) | 10-20× |
|
||||
| Metrics reduction | 5-20 ms | 0.1-0.5 ms | 20-40× |
|
||||
| Data transfer overhead | 2-5 ms (weight download) | 0 (stays on GPU) | ∞ |
|
||||
| **Per-trial eval (8 folds × 100K bars)** | **16-40s** | **1-3s** | **8-15×** |
|
||||
|
||||
---
|
||||
|
||||
## Phase 3: General-Purpose GPU Backtester
|
||||
|
||||
Extends Phase 2 to serve as a drop-in replacement for `crates/backtesting/strategy_runner.rs`.
|
||||
|
||||
### 3.1 Additional capabilities
|
||||
|
||||
- **Ensemble inference**: Batch forward through K models, GPU-side confidence aggregation (weighted vote)
|
||||
- **Full metrics suite**: Sharpe, Sortino, Calmar, VaR (percentile), CVaR (tail mean), win rate, avg trade return, profit factor — all via GPU reduction kernels
|
||||
- **Slippage modeling**: `VolumeImpactSlippage` (Almgren-Chriss `sqrtf()`) computed in env kernel
|
||||
- **Triple barrier episodes**: Port existing logic from `dqn_experience_kernel.cu` — already GPU-proven
|
||||
- **Position sizing**: Kelly criterion scaling (already in `dqn_experience_kernel.cu`)
|
||||
|
||||
### 3.2 New kernel: `backtest_full_metrics_kernel.cu`
|
||||
|
||||
Extends `backtest_metrics_kernel.cu` with:
|
||||
- Sortino ratio (downside deviation only)
|
||||
- Rolling max drawdown with recovery time tracking
|
||||
- VaR/CVaR via parallel sort + percentile extraction
|
||||
- Win rate / profit factor from per-trade P&L buffer
|
||||
- Annualization with configurable trading days (252 default)
|
||||
|
||||
### 3.3 Integration with evaluate_baseline binary
|
||||
|
||||
```rust
|
||||
// In bin targets: evaluate_baseline, evaluate_supervised
|
||||
fn run_evaluation(model_path: &Path, test_data: &[OHLCVBar], config: &EvalConfig) -> EvalReport {
|
||||
#[cfg(feature = "cuda")]
|
||||
if let Ok(device) = Device::new_cuda(0) {
|
||||
let evaluator = GpuBacktestEvaluator::new_full(test_data, config, &device)?;
|
||||
return evaluator.evaluate_full(model, &device)?;
|
||||
}
|
||||
// CPU fallback: existing AdaptiveStrategyRunner with SIMD
|
||||
run_cpu_evaluation(model_path, test_data, config)
|
||||
}
|
||||
```
|
||||
|
||||
### 3.4 Compatibility
|
||||
|
||||
- CPU path preserved for non-CUDA builds and CI testing
|
||||
- f32 GPU results validated against f64 CPU results (max 0.1% relative error for Sharpe, 0.01% for PnL)
|
||||
- `Decimal` precision maintained in CPU path; GPU uses f32 (sufficient for relative model comparison)
|
||||
|
||||
---
|
||||
|
||||
## File Map
|
||||
|
||||
### Phase 1 (modified files)
|
||||
- `crates/ml/src/trainers/dqn/trainer.rs` — remove CPU readback paths, add GPU-persistent epoch state
|
||||
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — persistent epoch_state buffer, monitoring reduction
|
||||
- `crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu` — epoch state read/write, monitoring accumulation
|
||||
- `crates/ml/src/cuda_pipeline/gpu_training_guard.rs` — async pinned-memory readback for grad norm
|
||||
|
||||
### Phase 2 (new files)
|
||||
- `crates/ml/src/cuda_pipeline/backtest_env_kernel.cu` — vectorized environment step kernel
|
||||
- `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` — per-window Sharpe/PnL/drawdown reduction
|
||||
- `crates/ml/src/cuda_pipeline/gpu_backtest_env.rs` — Rust wrapper for env kernel
|
||||
- `crates/ml/src/cuda_pipeline/gpu_backtest_metrics.rs` — Rust wrapper for metrics kernel
|
||||
- `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` — orchestrator (upload → step loop → metrics)
|
||||
- `crates/ml/src/hyperopt/gpu_evaluator.rs` — adapter integration
|
||||
|
||||
### Phase 2 (modified files)
|
||||
- `crates/ml/src/cuda_pipeline/mod.rs` — re-export new modules
|
||||
- `crates/ml/src/hyperopt/adapters/dqn.rs` — GPU eval path
|
||||
- `crates/ml/src/hyperopt/adapters/ppo.rs` — GPU eval path
|
||||
- `crates/ml/src/hyperopt/adapters/tft.rs` — GPU eval path (and remaining supervised adapters)
|
||||
|
||||
### Phase 3 (new files)
|
||||
- `crates/ml/src/cuda_pipeline/backtest_full_metrics_kernel.cu` — extended metrics (Sortino, VaR, etc.)
|
||||
- `crates/ml/src/cuda_pipeline/gpu_backtest_full_metrics.rs` — Rust wrapper for full metrics
|
||||
|
||||
### Phase 3 (modified files)
|
||||
- `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` — ensemble support, full metrics
|
||||
- Evaluate binary entry points
|
||||
|
||||
---
|
||||
|
||||
## Risks & Mitigations
|
||||
|
||||
| Risk | Mitigation |
|
||||
|------|-----------|
|
||||
| Candle tensor ↔ cudarc CudaSlice interop | Both use same CUDA context via `candle_core::cuda_backend::cudarc`; proven pattern in existing `gpu_experience_collector.rs` |
|
||||
| Numerical divergence GPU vs CPU | Validation tests comparing f32 GPU vs f64 CPU with documented tolerance bounds |
|
||||
| Kernel compilation time (NVRTC) | Cache compiled PTX in `GpuBacktestEvaluator::new()`, reuse across trials |
|
||||
| Variable-length windows | Pad to max length, use `window_lens` array + early `done_flag` exit per thread |
|
||||
| Model architecture diversity | Use Candle `.forward()` for all models — only env kernels are custom CUDA |
|
||||
|
||||
---
|
||||
|
||||
## Success Criteria
|
||||
|
||||
1. **Phase 1**: Zero `cudaStreamSynchronize` during DQN training epoch (only at epoch boundary for monitoring summary download)
|
||||
2. **Phase 2**: Hyperopt evaluation runs fully on GPU; per-trial eval time drops from 16-40s to 1-3s on H100
|
||||
3. **Phase 3**: `evaluate_baseline` binary uses GPU backtester when CUDA available, produces results within 0.1% of CPU path
|
||||
4. **All phases**: Existing CPU paths unchanged, no regressions in 2758+ lib tests
|
||||
Reference in New Issue
Block a user