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,