From f74d8d018ce1b17f3c2790325fac880fc0cd6f39 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 14 Apr 2026 08:21:37 +0200 Subject: [PATCH] =?UTF-8?q?fix:=20episode=20starts=20use=20epoch=20counter?= =?UTF-8?q?=20in=20Philox=20seed=20=E2=80=94=20breaks=20deterministic=20re?= =?UTF-8?q?petition?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit domain_rand_episode_starts used constant seed (i, 0, 9999) producing identical episode starting bars every epoch. With deterministic starts and alternating weight updates, the OOS Sharpe oscillated perfectly between +0.85 and -1.0. Adding epoch to the Philox hash key diversifies starting positions across epochs. Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml/src/cuda_pipeline/experience_kernels.cu | 8 +++++--- crates/ml/src/cuda_pipeline/gpu_experience_collector.rs | 3 +++ crates/ml/src/trainers/dqn/trainer/training_loop.rs | 2 +- 3 files changed, 9 insertions(+), 4 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index cd5fa9d2c..f49fa5ab3 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -165,18 +165,20 @@ extern "C" __global__ void domain_rand_episode_starts( int* __restrict__ episode_starts, /* [N] output */ int N, int stride, - int usable_bars + int usable_bars, + int epoch /* epoch counter for seed diversity */ ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; int base = (i * stride) % usable_bars; - /* Stateless jitter ±25% stride via Philox hash */ + /* Stateless jitter ±25% stride via Philox hash — epoch in seed + so each epoch samples different starting points */ if (stride > 0) { int range = stride / 4; if (range > 0) { - unsigned int bits = (unsigned int)(philox_uniform(i, 0, 9999) * (float)(range * 2 + 1)); + unsigned int bits = (unsigned int)(philox_uniform(i, epoch, 9999) * (float)(range * 2 + 1)); int jitter = (int)(bits) - range; base = ((base + jitter) % usable_bars + usable_bars) % usable_bars; } diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 91806e6ba..cf60347b7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -1286,8 +1286,10 @@ impl GpuExperienceCollector { n_episodes: usize, stride: i32, usable_bars: i32, + epoch: usize, ) -> Result<(), MLError> { let n = n_episodes as i32; + let epoch_i32 = epoch as i32; let blocks = ((n as u32 + 255) / 256) as u32; unsafe { self.stream @@ -1296,6 +1298,7 @@ impl GpuExperienceCollector { .arg(&n) .arg(&stride) .arg(&usable_bars) + .arg(&epoch_i32) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 4654ae1c4..4779beb88 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -931,7 +931,7 @@ impl DQNTrainer { let stride = (usable_bars / n_episodes).max(1); collector.generate_episode_starts_gpu( - n_episodes as usize, stride, usable_bars, + n_episodes as usize, stride, usable_bars, self.current_epoch, ).map_err(|e| anyhow::anyhow!("GPU episode starts: {e}"))?; collector.generate_sim_params_gpu( n_episodes as usize,