fix: episode starts use epoch counter in Philox seed — breaks deterministic repetition

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) <noreply@anthropic.com>

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;
         }

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),

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,