From 6ded2c55c2011f39e5132cdacfb1fcd30976d1b7 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 31 May 2026 12:01:50 +0200 Subject: [PATCH] =?UTF-8?q?style(rl):=20F1.3=20code=20review=20fixes=20?= =?UTF-8?q?=E2=80=94=20comments=20+=20spec=20annotations?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Code-quality review found 1 Important + 4 Minor; applied all 5: - I1: restored `(issue #9, drift-free)` cross-link in Phase 4 header - M1: added explicit launch profile in header `Grid=(1,1,1), Block=(1,1,1), smem=0` - M2: comment explaining `count` reuse in Phase 3 (Welford state retained in registers) - M3: inline rationale for `kelly == 0.0f` and `cooldown == 0.0f` exact equality - M4: extended Welford-bootstrap comment with statistical floor justification (exempt from ISV-promotion per "physics/math constants" rule) Zero functional change — comments only; cubin rebuilds identical bytes. --- crates/ml-alpha/cuda/rl_regime_observer.cu | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) 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];