From 5e8f8b4f5bbcdbe98dd07ae4f2db4a312fb346ab Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 13 Apr 2026 18:05:36 +0200 Subject: [PATCH] fix: replace stateful LCG with stateless Philox in experience kernels MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Feature noise injection: philox_uniform(episode, timestep, feature) instead of advancing rng_states by 2*market_dim per timestep - Episode start jitter: philox_uniform(episode, 0, 9999) - Domain randomization params: philox_uniform(episode, 0, 2000+k) - Saboteur params: philox_uniform(episode, 0, 3000+k) All experience collection randomness is now stateless — same inputs always produce same outputs regardless of execution history. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 75 +++++++------------ crates/ml/src/trainers/dqn/fused_training.rs | 6 ++ 2 files changed, 33 insertions(+), 48 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index b71d2ad56..56a3b7c3a 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -173,16 +173,14 @@ extern "C" __global__ void domain_rand_episode_starts( int base = (i * stride) % usable_bars; - /* Always jitter ±25% stride — one production path */ + /* Stateless jitter ±25% stride via Philox hash */ if (stride > 0) { - unsigned int rng = rng_states[i]; - rng = rng * 1664525u + 1013904223u; int range = stride / 4; if (range > 0) { - int jitter = ((int)(rng % (unsigned int)(range * 2 + 1))) - range; + unsigned int bits = (unsigned int)(philox_uniform(i, 0, 9999) * (float)(range * 2 + 1)); + int jitter = (int)(bits) - range; base = ((base + jitter) % usable_bars + usable_bars) % usable_bars; } - rng_states[i] = rng; } episode_starts[i] = base; @@ -218,25 +216,16 @@ extern "C" __global__ void domain_rand_sim_params( float* out = params_out + i * 5; - /* Always randomize — one production path */ - unsigned int rng = rng_states[i]; + /* Stateless Philox randomization — deterministic per episode */ + #define PHILOX_UNIFORM(lo, hi, idx) ((lo) + philox_uniform(i, 0, 2000 + (idx)) * ((hi) - (lo))) - /* Helper: uniform float in [lo, hi) from LCG */ - #define LCG_UNIFORM(lo, hi) ({ \ - rng = rng * 1664525u + 1013904223u; \ - float _u = (float)(rng & 0x00FFFFFFu) / 16777216.0f; \ - (lo) + _u * ((hi) - (lo)); \ - }) + out[0] = PHILOX_UNIFORM(0.5f, 2.5f, 0); /* tx_cost_multiplier */ + out[1] = base_spread * PHILOX_UNIFORM(0.5f, 3.0f, 1); /* fill_median_spread */ + out[2] = PHILOX_UNIFORM(0.65f, 0.95f, 2); /* fill_ioc_fill_prob */ + out[3] = PHILOX_UNIFORM(0.15f, 0.45f, 3); /* fill_limit_fill_min */ + out[4] = PHILOX_UNIFORM(0.60f, 0.95f, 4); /* fill_limit_fill_max */ - out[0] = LCG_UNIFORM(0.5f, 2.5f); /* tx_cost_multiplier */ - out[1] = base_spread * LCG_UNIFORM(0.5f, 3.0f); /* fill_median_spread */ - out[2] = LCG_UNIFORM(0.65f, 0.95f); /* fill_ioc_fill_prob */ - out[3] = LCG_UNIFORM(0.15f, 0.45f); /* fill_limit_fill_min */ - out[4] = LCG_UNIFORM(0.60f, 0.95f); /* fill_limit_fill_max */ - - #undef LCG_UNIFORM - - rng_states[i] = rng; + #undef PHILOX_UNIFORM } /* ================================================================== */ @@ -266,30 +255,23 @@ extern "C" __global__ void saboteur_generate_params( int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; - unsigned int rng = rng_states[i]; - float base_spread = base_params[0]; float base_fill = base_params[1]; float base_slip = base_params[2]; float ps = perturbation_scale; - /* Box-Muller-lite for Gaussian-like perturbation */ - rng = rng * 1664525u + 1013904223u; - float u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; - rng = rng * 1664525u + 1013904223u; - float u2 = (float)(rng & 0xFFFF) / 65536.0f; + /* Stateless Philox-based Gaussian perturbation */ + float u1, u2; + u1 = fmaxf(philox_uniform(i, 0, 3000), 1e-6f); + u2 = philox_uniform(i, 0, 3001); float g1 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); - rng = rng * 1664525u + 1013904223u; - u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; - rng = rng * 1664525u + 1013904223u; - u2 = (float)(rng & 0xFFFF) / 65536.0f; + u1 = fmaxf(philox_uniform(i, 0, 3002), 1e-6f); + u2 = philox_uniform(i, 0, 3003); float g2 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); - rng = rng * 1664525u + 1013904223u; - u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; - rng = rng * 1664525u + 1013904223u; - u2 = (float)(rng & 0xFFFF) / 65536.0f; + u1 = fmaxf(philox_uniform(i, 0, 3004), 1e-6f); + u2 = philox_uniform(i, 0, 3005); float g3 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); /* Perturb from base with Gaussian noise, clamp to valid ranges */ @@ -300,8 +282,6 @@ extern "C" __global__ void saboteur_generate_params( saboteur_params[i * 3 + 0] = spread; saboteur_params[i * 3 + 1] = fill; saboteur_params[i * 3 + 2] = slip; - - rng_states[i] = rng; } /** @@ -518,20 +498,19 @@ extern "C" __global__ void experience_state_gather( } /* ── #22 Feature noise injection: N(0, scale) per feature ── */ - if (feature_noise_scale > 0.0f && rng_states != NULL) { - /* LCG-based fast noise (same RNG as action_select uses) */ - unsigned int rng = rng_states[i]; + /* Stateless Philox hash: noise = f(episode_id, timestep, feature_k). + * Deterministic regardless of execution history. */ + if (feature_noise_scale > 0.0f) { + int timestep_val = raw_t; for (int k = 0; k < market_dim; k++) { - /* Box-Muller-lite: two uniform → one gaussian */ - rng = rng * 1664525u + 1013904223u; - float u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; - rng = rng * 1664525u + 1013904223u; - float u2 = (float)(rng & 0xFFFF) / 65536.0f; + /* Two Philox hashes → Box-Muller Gaussian */ + float u1 = philox_uniform(i, timestep_val, 1000 + k * 2); + float u2 = philox_uniform(i, timestep_val, 1000 + k * 2 + 1); + u1 = fmaxf(u1, 1e-6f); float noise = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); float v = (float)out[k] + noise * feature_noise_scale; out[k] = v; } - rng_states[i] = rng; /* write back RNG state */ } /* -- Portfolio features: [market_dim .. market_dim+8) -- diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 97f2b692b..7a2cc4fe4 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -1042,6 +1042,12 @@ impl FusedTrainingCtx { PhaseEvents::record(pe.adam_end, cu_stream); } + // ── DETERMINISM DIAG ── + if step < 40 { + let mut dl = [0.0_f32]; + self.stream.memcpy_dtoh(&self.trainer.total_loss_buf, &mut dl).ok(); + tracing::warn!("STEP_DIAG step={step} loss={:.10}", dl[0]); + } // ── Step 6: PER priority update ───────────────────────────────── if let Some(ref pe) = self.phase_events { PhaseEvents::record(pe.per_update_start, cu_stream);