033906f213
While alpha-rl-8gtk2 (the B-7+B-8+B-9 20k+5k cluster run) was training, B-9's saturation diag exposed V_MAX_eff dropping from 856.82 (step 482) to 464.61 (step 817) within the same run. The kernel docstrings and slot doc-comments uniformly claimed "ratchet (monotone-grow)" semantics — contradicting the observation by 392 units. Root cause: wwcsz followup 2026-05-24 (commit landed in rl_reward_clamp_controller.cu Step 5 only) REPLACED the original ratchet with a slow symmetric EWMA (α=0.001, half-life ~700 steps) on `win_bound`/`loss_bound`. Static ratchet was wasting atom resolution on rare tails (avg rewards in [-5,+5] with span at [-60,+20] → Δz=4, Q couldn't distinguish "slightly winning" from "slightly losing"). The EWMA refocuses atom resolution on the ACTIVE range. The controller's own header was correctly updated at the time. The documentation drift was in 3 other places — fixed here: - bellman_target_projection.cu header (lines 47-49): "ratchet (monotone-grow)" → accurate EWMA description with cross-references + pearl_c51_v_max_freeze_required_for_surfer warning (V_MAX in 100-200 → trend-follower; past 1000 → degraded). - rl_atom_support_update.cu line 4: "Companion to the C51 atom-span ratchet" → "Companion to the C51 atom-span EWMA". - isv_slots.rs slot allocation table line 43: "C51 atom-span ratchet slots" → "C51 atom-span EWMA slots (α=0.001)". - isv_slots.rs RL_C51_V_MAX_INDEX / RL_C51_V_MIN_INDEX doc-comments (lines 649-668): replaced with accurate EWMA description, observed 857 → 465 drop example, and asymmetric tracking note (V_MIN_eff EWMAs -loss_bound NOT -V_MAX_eff — they only coincide when ratio≈1). Latent test bug also surfaced + fixed: c51_atom_saturation_diagnostic assumed `V_MIN_eff = -V_MAX_eff` (line 136). With observed Kelly EMAs avg_loss=$436 vs avg_win=$872 → ratio ≈ 0.5 → V_MIN_eff ≈ -0.5 × V_MAX_eff, the test's bot-rate consistency check would false-fail if saturation ever became non-zero. Dormant so far (sat=0% in both smoke and full run). Relaxed the bot-rate assertion to use the v_bound_floor (-1.0) as the necessary lower bound — correct for any asymmetric ratio and catches gross drift without false-failing on legitimate adaptation. No behavior change; pure docstring drift + dormant test bug repair per feedback_trust_code_not_docs. Compile clean. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
39 lines
1.7 KiB
Plaintext
39 lines
1.7 KiB
Plaintext
// rl_atom_support_update.cu — refresh `atom_supports_d` from the
|
||
// ISV-driven [V_MIN, V_MAX] span (audit 2026-05-24 second follow-up).
|
||
//
|
||
// Companion to the C51 atom-span EWMA in rl_reward_clamp_controller
|
||
// (α=0.001, ~700-step half-life — wwcsz followup 2026-05-24 replaced
|
||
// the original ratchet to focus atom resolution on the ACTIVE range).
|
||
// `atom_supports_d` is the device buffer of 21 float values that the
|
||
// non-projection C51 kernels (`argmax_expected_q`, `rl_action_kernel`,
|
||
// `dqn_distributional_q`) read instead of recomputing the per-atom
|
||
// values themselves. When V_MIN/V_MAX adapt, this buffer must be
|
||
// rewritten or those kernels see a stale span.
|
||
//
|
||
// One block, Q_N_ATOMS threads. Each thread writes one atom value:
|
||
// atom_supports[i] = V_MIN + i * (V_MAX - V_MIN) / (N_ATOMS - 1)
|
||
//
|
||
// Per `feedback_no_atomicadd`: no atomics. Per `feedback_cpu_is_read_only`:
|
||
// pure device write. Per `feedback_no_htod_htoh_only_mapped_pinned`: no
|
||
// host transfer — values come from ISV slots, written from the trainer
|
||
// per-step launch right after the reward clamp controller refreshes
|
||
// V_MIN/V_MAX.
|
||
|
||
#define Q_N_ATOMS 21
|
||
#define RL_C51_V_MAX_INDEX 484
|
||
#define RL_C51_V_MIN_INDEX 485
|
||
|
||
extern "C" __global__ void rl_atom_support_update(
|
||
const float* __restrict__ isv, // ≥ RL_C51_V_MAX_INDEX + 1
|
||
float* __restrict__ atom_supports // [Q_N_ATOMS]
|
||
) {
|
||
const int tid = threadIdx.x;
|
||
if (tid >= Q_N_ATOMS) return;
|
||
|
||
const float v_min = isv[RL_C51_V_MIN_INDEX];
|
||
const float v_max = isv[RL_C51_V_MAX_INDEX];
|
||
const float delta = (v_max - v_min) / (float)(Q_N_ATOMS - 1);
|
||
|
||
atom_supports[tid] = v_min + (float)tid * delta;
|
||
}
|