From 6c06168230a15cbddfaed527f1e9787f42db2cd3 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 13 Apr 2026 17:25:10 +0200 Subject: [PATCH] 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) --- .../src/cuda_pipeline/experience_kernels.cu | 47 +++++++++++++------ .../cuda_pipeline/gpu_backtest_evaluator.rs | 1 + .../cuda_pipeline/gpu_experience_collector.rs | 1 + 3 files changed, 35 insertions(+), 14 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 06876a814..b71d2ad56 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/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++) { diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 428d1601a..987044043 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/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}" diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 822fd2a07..cfa3f6dd1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/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}"