fix: stateless Philox RNG for experience action selection (replace LCG)

experience_action_select kernel now uses philox_uniform(episode_id,
timestep, call_counter) instead of stateful lcg_random(&rng).
Eliminates branching-dependent RNG state divergence across processes.

Also passed timestep to backtest evaluator's action_select launch.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -111,6 +111,23 @@ __device__ __forceinline__ float lcg_random(unsigned int* state) {
     return (float)(*state & 0x00FFFFFFu) / 16777216.0f;
 }
 
+/**
+ * Stateless Philox-based uniform random in [0, 1).
+ * Deterministic given (episode_id, timestep, call_idx).
+ * No accumulated state — immune to branching-dependent divergence.
+ */
+__device__ __forceinline__ float philox_uniform(int episode_id, int timestep, int call_idx) {
+    unsigned int key = (unsigned int)episode_id;
+    unsigned int ctr = (unsigned int)(timestep * 37 + call_idx);
+    /* Single-round Philox-like hash */
+    ctr ^= key * 0x9E3779B9u;
+    ctr *= 0x85ebca6bu;
+    ctr ^= ctr >> 13;
+    ctr *= 0xc2b2ae35u;
+    ctr ^= ctr >> 16;
+    return (float)(ctr & 0x00FFFFFFu) / 16777216.0f;
+}
+
 /**
  * Argmax over a float array of length n.
  * Returns the index of the maximum element; ties broken in favour of the
@@ -737,12 +754,14 @@ extern "C" __global__ void experience_action_select(
     float max_position,
     float eps_exp_mult,    /* per-branch epsilon multiplier: exposure (dir+mag) */
     float eps_ord_mult,    /* per-branch epsilon multiplier: order type */
-    float eps_urg_mult     /* per-branch epsilon multiplier: urgency */
+    float eps_urg_mult,    /* per-branch epsilon multiplier: urgency */
+    int timestep           /* current timestep for stateless RNG */
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
 
-    unsigned int rng = rng_states[i];
+    unsigned int rng = rng_states[i]; /* kept for NoisyNet noise, not for action selection */
+    int rng_ctr = 0; /* stateless call counter for philox_uniform */
 
     /* Eval mode: eps_start==0 AND eps_end==0 AND all multipliers==1.
      * In eval mode, use pure greedy argmax — no Boltzmann, no epsilon floor.
@@ -835,8 +854,8 @@ extern "C" __global__ void experience_action_select(
      * conviction gating while allowing diverse exploration.
      *
      * Temperature tau = Q_range: max e:1 ratio between best and worst direction. */
-    if (!eval_mode && lcg_random(&rng) < eps_dir) {
-        int r = (int)(lcg_random(&rng) * (float)b0_size);
+    if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_dir) {
+        int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b0_size);
         dir_idx = (r >= b0_size) ? b0_size - 1 : r;
     } else if (eval_mode) {
         /* Eval mode: pure greedy argmax — deterministic, no RNG */
@@ -864,7 +883,7 @@ extern "C" __global__ void experience_action_select(
                 exps_d[a] = expf((qv - q_max_d) / tau_d);
                 sum_e += exps_d[a];
             }
-            float r = lcg_random(&rng) * sum_e;
+            float r = philox_uniform(i, timestep, rng_ctr++) * sum_e;
             float cum = 0.0f;
             dir_idx = b0_size - 1;
             for (int a = 0; a < b0_size; a++) {
@@ -901,8 +920,8 @@ extern "C" __global__ void experience_action_select(
                 float qv = (q_b1[a]);
                 if (qv > best_q) { best_q = qv; mag_idx = a; }
             }
-        } else if (lcg_random(&rng) < eps_mag) {
-            int r = (int)(lcg_random(&rng) * (float)b1_size);
+        } else if (philox_uniform(i, timestep, rng_ctr++) < eps_mag) {
+            int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b1_size);
             mag_idx = (r >= b1_size) ? b1_size - 1 : r;
         } else {
             /* Adaptive temperature: scale by Q-range so Boltzmann is meaningful
@@ -928,7 +947,7 @@ extern "C" __global__ void experience_action_select(
                 exps[a] = expf((qv - q_max_m) / tau);
                 sum_e += exps[a];
             }
-            float r = lcg_random(&rng) * sum_e;
+            float r = philox_uniform(i, timestep, rng_ctr++) * sum_e;
             float cum = 0.0f;
             mag_idx = b1_size - 1;
             for (int a = 0; a < b1_size; a++) {
@@ -946,8 +965,8 @@ extern "C" __global__ void experience_action_select(
         for (int a = 1; a < b2_size; a++) {
             if (q_b2[a] > best_q) { best_q = q_b2[a]; a2 = a; }
         }
-    } else if (lcg_random(&rng) < eps_ord) {
-        int r = (int)(lcg_random(&rng) * (float)b2_size);
+    } else if (philox_uniform(i, timestep, rng_ctr++) < eps_ord) {
+        int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b2_size);
         a2 = (r >= b2_size) ? b2_size - 1 : r;
     } else {
         /* Boltzmann softmax over order Q-values */
@@ -968,7 +987,7 @@ extern "C" __global__ void experience_action_select(
             exps_ord[a] = expf((q_b2[a] - q_max_ord) / tau_ord);
             sum_e += exps_ord[a];
         }
-        float ro = lcg_random(&rng) * sum_e;
+        float ro = philox_uniform(i, timestep, rng_ctr++) * sum_e;
         float cum = 0.0f;
         a2 = b2_size - 1;
         for (int a = 0; a < b2_size; a++) {
@@ -984,8 +1003,8 @@ extern "C" __global__ void experience_action_select(
         for (int a = 1; a < b3_size; a++) {
             if (q_b3[a] > best_q) { best_q = q_b3[a]; a3 = a; }
         }
-    } else if (lcg_random(&rng) < eps_urg) {
-        int r = (int)(lcg_random(&rng) * (float)b3_size);
+    } else if (philox_uniform(i, timestep, rng_ctr++) < eps_urg) {
+        int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b3_size);
         a3 = (r >= b3_size) ? b3_size - 1 : r;
     } else {
         /* Boltzmann softmax over urgency Q-values */
@@ -1001,7 +1020,7 @@ extern "C" __global__ void experience_action_select(
             exps_urg[a] = expf((q_b3[a] - q_max_urg) / tau_urg);
             sum_e += exps_urg[a];
         }
-        float ru = lcg_random(&rng) * sum_e;
+        float ru = philox_uniform(i, timestep, rng_ctr++) * sum_e;
         float cum = 0.0f;
         a3 = b3_size - 1;
         for (int a = 0; a < b3_size; a++) {

crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs

@@ -979,6 +979,7 @@ impl GpuBacktestEvaluator {
                     .arg(&1.0_f32)     // eps_exp_mult (no-op at eps=0)
                     .arg(&1.0_f32)     // eps_ord_mult (no-op at eps=0)
                     .arg(&1.0_f32)     // eps_urg_mult (no-op at eps=0)
+                    .arg(&(chunk_start as i32)) // timestep for stateless Philox RNG
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "chunked action_select chunk_start={chunk_start}: {e}"

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -2107,6 +2107,7 @@ impl GpuExperienceCollector {
                     .arg(&config.eps_exp_mult)          // per-branch epsilon multiplier: exposure
                     .arg(&config.eps_ord_mult)          // per-branch epsilon multiplier: order
                     .arg(&config.eps_urg_mult)          // per-branch epsilon multiplier: urgency
+                    .arg(&(t as i32))                   // timestep for stateless Philox RNG
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_action_select t={t}: {e}"