56efd96cb2
Adds the DQN component of the integrated RL trainer per the plan at docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md. What this commit lands: - DqnHead: linear projection h_t [B, HIDDEN_DIM] -> atom logits [B, N_ACTIONS=9, Q_N_ATOMS=21] with parallel target-network weights. Xavier x 0.01 init (initial softmax-over-atoms approx uniform), scoped_init_seed-guarded per pearl_scoped_init_seed_for_reproducibility. - dqn_distributional_q.cu: forward (one block per (batch, action), one thread per atom) + Bellman categorical-CE backward against a pre-projected target distribution. Atom-softmax fused into backward. - ReplayBuffer (rl/replay.rs): capacity-bounded PER with priority^alpha sampling, random replacement, and TD-error priority update. O(N) cumulative-sum sampling; Phase E may upgrade to a GPU sum-tree once capacity profiling demands it. - rl_gamma_controller.cu: ISV[RL_GAMMA_INDEX=400] producer; gamma adapts toward 0.5^(1/mean_trade_duration) via Wiener-alpha blend (floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary), clamped to [0.90, 0.999]. Bootstrap gamma = 0.99 on sentinel. - rl_target_tau_controller.cu: ISV[RL_TARGET_TAU_INDEX=401] producer; tau adapts multiplicatively from Q-divergence ratio vs anchor 0.01, Wiener-alpha blend with floor 0.4, clamped to [0.001, 0.05]. Bootstrap tau = 0.005 on sentinel. - Action enum + try_from_u32 in rl/common.rs (matches existing ml DQN action grid for cross-system policy comparability). - C51 atom support constants Q_V_MIN / Q_V_MAX in rl/common.rs (kept for Phase E's projection kernel; backward in this commit operates in categorical domain on a pre-projected target). What this commit DEFERS to Phase E: - soft_update_target kernel (struct fields w_target_d / b_target_d are wired and read by the Bellman backward in this commit; the writer lives in Phase E alongside the training-loop tau driver). - Categorical projection kernel that reads gamma from ISV[400] and produces the target_dist input to the backward kernel. - Toy bandit test activation (tests/dqn_toy.rs is #[ignore]-gated; the type contract is locked here, the training loop wires in Phase E). - atomicAdd in the per-batch CE accumulator (Phase E replaces with the warp-shuffle + shared reduce pattern from aux_loss.cu when batches reach production sizes; B <= 32 toy contention is negligible). Per pearl_controller_anchors_isv_driven and feedback_isv_for_adaptive_bounds: gamma and tau are NOT hardcoded constants. They live in ISV[400] / ISV[401], emitted by the controller kernels above, and Phase E consumers read via __ldg(isv + INDEX). Bootstrap values (0.99, 0.005) appear only in the controller kernel as first-observation defaults, NOT baked into the loss kernel. Validation: - SQLX_OFFLINE=true cargo check -p ml-alpha --lib -> clean (1m 03s) - SQLX_OFFLINE=true cargo check --workspace --lib -> clean (42s) - SQLX_OFFLINE=true cargo test -p ml-alpha --lib rl::replay -> 3 pass - Cubins built for all 3 new kernels (sm_80 default). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
83 lines
3.8 KiB
Plaintext
83 lines
3.8 KiB
Plaintext
// rl_gamma_controller.cu — emits γ to ISV[RL_GAMMA_INDEX=400].
|
||
//
|
||
// Phase C of the integrated RL trainer
|
||
// (docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md).
|
||
//
|
||
// γ is the Bellman discount factor used by the categorical Bellman
|
||
// projection (Phase E) and by the PPO advantage estimator (Phase D).
|
||
// Per `pearl_controller_anchors_isv_driven` and
|
||
// `feedback_isv_for_adaptive_bounds`, γ is NOT a hardcoded constant; it
|
||
// adapts so that
|
||
//
|
||
// γ^(mean_trade_duration_events) ≈ 0.5
|
||
//
|
||
// i.e. roughly half of the discounted value remains by the time the
|
||
// average trade closes. This anchors the credit-assignment horizon to
|
||
// the actual trade timescale rather than a tuned constant; if the
|
||
// strategy lengthens its holding period, γ rises automatically.
|
||
//
|
||
// Bootstrap discipline (per `pearl_first_observation_bootstrap`): the
|
||
// ISV slot starts at 0.0 (sentinel "uninitialised"). On the first
|
||
// emit the kernel writes γ = 0.99 directly. Subsequent emits use a
|
||
// Wiener-α blend with floor 0.4 per
|
||
// `pearl_wiener_alpha_floor_for_nonstationary` (the target γ_target
|
||
// drifts as the strategy's holding period co-adapts with the policy,
|
||
// which violates the stationarity precondition of Wiener-optimal α).
|
||
//
|
||
// Bounds: γ ∈ [0.90, 0.999] enforced by clamp to prevent runaway
|
||
// (γ → 1 makes credit assignment infinite-horizon; γ → 0.5 makes
|
||
// the policy myopic).
|
||
|
||
#define RL_GAMMA_INDEX 400
|
||
#define GAMMA_MIN 0.90f
|
||
#define GAMMA_MAX 0.999f
|
||
#define GAMMA_BOOTSTRAP 0.99f
|
||
#define WIENER_ALPHA_FLOOR 0.4f
|
||
|
||
|
||
// ─────────────────────────────────────────────────────────────────────
|
||
// rl_gamma_controller:
|
||
// Single-thread controller — writes ONE float to isv[RL_GAMMA_INDEX].
|
||
//
|
||
// Inputs:
|
||
// isv [≥ RL_GAMMA_INDEX+1] — ISV bus
|
||
// alpha — Wiener-α from the controller's own
|
||
// signal stats (caller computes this
|
||
// upstream from γ-divergence variance).
|
||
// Floored at WIENER_ALPHA_FLOOR before
|
||
// the blend.
|
||
// mean_trade_duration_events — EMA of trade hold time in event count.
|
||
// Caller is responsible for the EMA
|
||
// (Phase E); we just read the scalar.
|
||
//
|
||
// Outputs:
|
||
// isv[RL_GAMMA_INDEX] — γ ∈ [GAMMA_MIN, GAMMA_MAX]
|
||
// ─────────────────────────────────────────────────────────────────────
|
||
extern "C" __global__ void rl_gamma_controller(
|
||
float* __restrict__ isv,
|
||
float alpha,
|
||
float mean_trade_duration_events
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
|
||
const float gamma_prev = isv[RL_GAMMA_INDEX];
|
||
// Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap.
|
||
if (gamma_prev == 0.0f) {
|
||
isv[RL_GAMMA_INDEX] = GAMMA_BOOTSTRAP;
|
||
return;
|
||
}
|
||
|
||
// Target: γ^d ≈ 0.5 ⇒ γ = 0.5^(1/d). Clamp d ≥ 1 so a single-event
|
||
// trade doesn't push γ to 0.5.
|
||
const float d = fmaxf(mean_trade_duration_events, 1.0f);
|
||
const float gamma_target = powf(0.5f, 1.0f / d);
|
||
|
||
// Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
|
||
const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR);
|
||
float gamma_new = (1.0f - a) * gamma_prev + a * gamma_target;
|
||
|
||
// Clamp into the bounded range; runaway protection.
|
||
gamma_new = fmaxf(GAMMA_MIN, fminf(gamma_new, GAMMA_MAX));
|
||
isv[RL_GAMMA_INDEX] = gamma_new;
|
||
}
|