gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix: replace stateful LCG with stateless Philox in experience kernels

Browse Source 
- 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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-13 18:05:36 +02:00
parent 78970d37fe
commit 5e8f8b4f5b
2 changed files with 33 additions and 48 deletions
Download Patch File Download Diff File Expand all files Collapse all files

75
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

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

6
crates/ml/src/trainers/dqn/fused_training.rs
View File

@@ -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);