diff --git a/crates/ml-alpha/cuda/rl_regime_observer.cu b/crates/ml-alpha/cuda/rl_regime_observer.cu index a7e858473..449abccc5 100644 --- a/crates/ml-alpha/cuda/rl_regime_observer.cu +++ b/crates/ml-alpha/cuda/rl_regime_observer.cu @@ -8,6 +8,8 @@ // Per feedback_no_atomicadd: single-thread single-block, no atomics. // Per feedback_cpu_is_read_only: pure device kernel. // Per feedback_nvidia_grade_perf_for_kernels: no host branches in capture. +// +// Launch: Grid=(1,1,1), Block=(1,1,1), smem=0. // ─── ISV slot indices (mirror Rust constants from isv_slots.rs) ───── #define RL_KELLY_FRACTION_INDEX 676 @@ -39,6 +41,10 @@ extern "C" __global__ void rl_regime_observer( if (threadIdx.x != 0 || blockIdx.x != 0) return; // ─── Phase 1: dead-zone detection ────────────────────── + // kelly==0.0f: analytic Kelly clamp writes exactly 0.0 when fraction ≤ 0 + // (see rl_kelly_fraction_controller). Exact equality is the intended test. + // cooldown==0.0f: integer-valued step counter stored as float; exact 0 is + // the only safe comparison point. const float kelly = isv[RL_KELLY_FRACTION_INDEX]; // prior-step value const float cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX]; const int flat_count = flat_count_d[0]; @@ -84,11 +90,15 @@ extern "C" __global__ void rl_regime_observer( } // ─── Phase 3: tail-event detection (3σ threshold) ────── + // `count` is intentionally reused below from Phase 2 (in-register value + // post-conditional-increment); no re-read from ISV needed. const float var = isv[RL_REGIME_SESSION_PNL_VARIANCE_EMA_INDEX]; const float sigma_thr = isv[RL_REGIME_TAIL_SIGMA_THRESHOLD_INDEX]; const float sigma = sqrtf(var); - // Require count >= 10 before tail detection (Welford bootstrap) + // Require count >= 10 before tail detection. Statistical floor — Welford + // sample variance is undefined for N<2 and noisy for N<10; exempt from + // ISV-promotion per "physics/math constants" rule. const int sufficient_count = (count >= 10.0f) ? 1 : 0; const int is_tail = sufficient_count && (sigma > 0.0f) && (fabsf(dx) > sigma_thr * sigma); @@ -96,7 +106,7 @@ extern "C" __global__ void rl_regime_observer( if (is_tail) recency = 0.0f; else recency += 1.0f; isv[RL_REGIME_TAIL_EVENT_RECENCY_INDEX] = recency; - // ─── Phase 4: emit recovery_factor and eps_recovery_live (drift-free) ── + // ─── Phase 4: emit recovery_factor and eps_recovery_live (issue #9, drift-free) ── const float n_recovery = isv[RL_KELLY_EPS_RECOVERY_N_RECOVERY_INDEX]; const float eps_min = isv[RL_KELLY_EPS_RECOVERY_MIN_INDEX]; const float eps_max = isv[RL_KELLY_EPS_RECOVERY_MAX_INDEX];