Compare commits
8 Commits
ml-alpha-p
...
worktree-a
| Author | SHA1 | Date | |
|---|---|---|---|
|
a3dfcd63f5 | ||
|
|
6695785666 | ||
|
|
12635bd708 | ||
|
|
25f5ce99b6 | ||
|
|
acdafe508e | ||
|
|
13bf277cd6 | ||
|
|
af35bc778e | ||
|
|
fd31742627 |
1
Cargo.lock
generated
1
Cargo.lock
generated
@@ -6090,7 +6090,6 @@ dependencies = [
|
||||
"approx",
|
||||
"arrow 56.2.0",
|
||||
"bincode",
|
||||
"bytemuck",
|
||||
"clap",
|
||||
"cudarc",
|
||||
"data",
|
||||
|
||||
@@ -33,7 +33,6 @@ tracing-subscriber.workspace = true
|
||||
serde = { workspace = true, features = ["derive"] }
|
||||
serde_json.workspace = true
|
||||
bincode.workspace = true
|
||||
bytemuck = { workspace = true }
|
||||
thiserror.workspace = true
|
||||
clap = { workspace = true, features = ["derive"] }
|
||||
|
||||
|
||||
@@ -73,7 +73,6 @@ const KERNELS: &[&str] = &[
|
||||
"rl_kl_reference_grad",
|
||||
"rl_q_pi_distill_grad", // audit 2026-05-24 vj5f6 followup: KL(softmax(E_Q/τ) || π_new) gradient ADDED to pi_grad_logits — couples Q's improved C51 calibration to action selection (was decoupled per Option B)
|
||||
"action_entropy_per_step", // POST-gate action entropy EMA for SAC α/τ co-tuning
|
||||
"rl_drawdown_stop", // per-step drawdown penalty + hard stop-loss
|
||||
"rl_q_distill_lambda_controller",// audit 2026-05-24 rljzl followup: adaptive λ_distill via Schulman bounded step on KL_EMA vs target
|
||||
"rl_unit_state_update", // SP20 P1+P5 audit fix: per-unit trade state machine — detects open/close/reverse transitions, sets up unit slot 0 entry+trail
|
||||
"rl_trail_mutate", // SP20 P1+P5 audit fix (a7/a8 dead): TrailTighten/TrailLoosen mutate unit_trail_distance bounded MIN/MAX, symmetric reciprocal adjust
|
||||
@@ -97,6 +96,13 @@ const KERNELS: &[&str] = &[
|
||||
"rl_iqn_forward", // IQN distributional Q-head: quantile embedding + action-value projection; complementary to C51
|
||||
"rl_iqn_loss", // IQN quantile Huber loss: ρ_τ(δ) = |τ - 1(δ<0)| × Huber(δ, κ=1.0); forward + backward
|
||||
"rl_iqn_backward", // IQN backward through forward pass: grad_output → grad_w_out/b_out/w_embed/b_embed per-batch scratch
|
||||
"rl_dueling_q_forward", // Phase 4 (2026-05-30): Independent dueling Q head forward — V[B] + A[B,N] + composed_Q[B,N] = V + A − mean_a A; parallel to C51/IQN, zero shared state per spec 2026-05-30-phase4-independent-dueling-head-design.md
|
||||
"rl_dueling_q_bellman_target", // Phase 4: argmax over target composed_Q + Bellman target = r + γ^n × (1-done) × max_Q
|
||||
"rl_dueling_q_loss_and_grad", // Phase 4: Huber loss on (target − online_composed_Q[taken]) + grad_composed
|
||||
"rl_dueling_q_decompose_and_bwd", // Phase 4: decompose grad_composed → grad_V + grad_A via mean-subtraction Jacobian + per-batch weight gradients
|
||||
"rl_v_blend", // Phase 4.4 (2026-05-30): elementwise V_used = α V_scalar + (1−α) V_dq, α from ISV
|
||||
"rl_v_blend_alpha_controller", // Phase 4.4: ISV-adaptive Schulman-bounded controller on α from observed |V_dq − V_scalar| / |V_scalar| tracking ratio
|
||||
"rl_advantage_normalize", // Phase 4.5 (2026-05-30): per-batch advantage normalization (A − mean)/std with ε² variance floor — standard PPO practice, self-adaptive (no tuned params)
|
||||
"rl_ensemble_action_value", // C51+IQN ensemble: E_ensemble = α×E_C51 + (1-α)×E_IQN; α from ISV[544]
|
||||
"rl_noisy_linear_forward", // NoisyNet: factored noisy linear forward — y = (mu_w + sigma_w ⊙ eps_w) × x + (mu_b + sigma_b ⊙ eps_b); state-dependent exploration for C51/IQN final projection
|
||||
"rl_noisy_linear_backward", // NoisyNet: factored noisy linear backward — grad_mu_w/sigma_w/mu_b/sigma_b per-batch scratch for reduce_axis0
|
||||
@@ -112,7 +118,6 @@ const KERNELS: &[&str] = &[
|
||||
"rl_hindsight_track", // HER Phase 1: per-step mid-price ring + peak tracking for backward hindsight
|
||||
"rl_hindsight_inject", // HER Phase 2: backward inject — synthetic replay push on done if peak >> actual
|
||||
"rl_hindsight_forward", // HER Phase 3: forward continuation — evaluates closed trades after lookahead
|
||||
"rl_lr_from_mapped_pinned", // Mega-graph: LR controller reads loss from mapped-pinned dev_ptrs (eliminates host loss readback between reward + training graphs)
|
||||
"rl_increment_step", // device-resident step counter bump (ISV[548] += 1.0); graph-safe prereq — removes scalar current_step from all downstream kernel args
|
||||
"rl_fused_controllers", // fused kernel: 10 RL ISV controllers in one launch (gamma, tau, ppo_clip, entropy_coef, rollout_steps, per_alpha, reward_scale, ppo_ratio_clamp, gate_threshold, q_distill_lambda) — saves 9 kernel launch overheads (~40-80μs/step)
|
||||
"rl_popart_normalize", // PopArt: Welford-EMA reward normalization (replaces apply_reward_scale)
|
||||
|
||||
@@ -44,41 +44,3 @@ extern "C" __global__ void adamw_step(
|
||||
|
||||
theta[i] -= lr * (m_hat / (sqrtf(v_hat) + eps) + wd * theta[i]);
|
||||
}
|
||||
|
||||
// Mega-graph variant: reads LR from an ISV device pointer at a given
|
||||
// slot index instead of taking a scalar argument. This allows the LR
|
||||
// to vary across graph replays (the ISV slot is modified in-place by
|
||||
// the rl_lr_from_mapped_pinned controller kernel which is captured
|
||||
// earlier in the same graph). All other args are identical to adamw_step.
|
||||
extern "C" __global__ void adamw_step_isv_lr(
|
||||
float* __restrict__ theta,
|
||||
const float* __restrict__ grad,
|
||||
float* __restrict__ m,
|
||||
float* __restrict__ v,
|
||||
int n_params,
|
||||
const float* __restrict__ isv,
|
||||
int lr_slot,
|
||||
float beta1,
|
||||
float beta2,
|
||||
float eps,
|
||||
float wd,
|
||||
int step
|
||||
) {
|
||||
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i >= n_params) return;
|
||||
|
||||
const float lr = isv[lr_slot];
|
||||
|
||||
const float g = grad[i];
|
||||
const float m_n = beta1 * m[i] + (1.0f - beta1) * g;
|
||||
const float v_n = beta2 * v[i] + (1.0f - beta2) * g * g;
|
||||
m[i] = m_n;
|
||||
v[i] = v_n;
|
||||
|
||||
const float bc1 = 1.0f - powf(beta1, (float) step);
|
||||
const float bc2 = 1.0f - powf(beta2, (float) step);
|
||||
const float m_hat = m_n / fmaxf(bc1, 1e-12f);
|
||||
const float v_hat = v_n / fmaxf(bc2, 1e-12f);
|
||||
|
||||
theta[i] -= lr * (m_hat / (sqrtf(v_hat) + eps) + wd * theta[i]);
|
||||
}
|
||||
|
||||
@@ -392,10 +392,14 @@ extern "C" __global__ void heads_lr_multiplier_scale_kernel(
|
||||
// `channels_in_bucket[bucket][i]` populated at transition by
|
||||
// `channels_in_bucket_kernel`.
|
||||
//
|
||||
// Launch: grid = (N_HORIZONS, 1, 1), block = (32, 1, 1). Each block
|
||||
// Launch: grid = (N_HORIZONS, 1, 1), block = (128, 1, 1). Each block
|
||||
// reduces one bucket. Per `feedback_no_atomicadd`, reduction is
|
||||
// block-tree on shared memory (no atomicAdd).
|
||||
//
|
||||
// Block size = 128 covers MAX_BUCKET_DIM=96 (smallest pow2 ≥ 96).
|
||||
// Previous block_dim=32 silently dropped channels 32..bdim when bdim>32
|
||||
// — caught by `tests/bucket_transition_kernels.rs` (block_dim=43 case).
|
||||
//
|
||||
// Input `h_state` is `[B × HIDDEN_DIM]` (ORIGINAL channel layout). Each
|
||||
// block reads its bucket's `bucket_dim_k[bucket]` channels (via
|
||||
// `channels_in_bucket[bucket][0..bdim]`) per sample (across all B
|
||||
@@ -412,20 +416,20 @@ extern "C" __global__ void h_mag_per_bucket_kernel(
|
||||
if (bucket >= N_HORIZONS) return;
|
||||
int bdim = (int)bucket_dim_k[bucket];
|
||||
|
||||
// sdata sized for a single-warp reduction (32 lanes).
|
||||
__shared__ float sdata[32];
|
||||
// sdata sized for the 128-lane block reduction.
|
||||
__shared__ float sdata[128];
|
||||
int tid = threadIdx.x;
|
||||
|
||||
// Each thread sums |h| over its bucket-local index (tid), looking up
|
||||
// the ORIGINAL channel via channels_in_bucket. Threads with tid >= bdim
|
||||
// idle — uniform predicate, no warp divergence inside [0, 32).
|
||||
// contribute 0. With bdim ∈ [1, MAX_BUCKET_DIM=96] and block_dim=128,
|
||||
// tail lanes [bdim, 128) are always inactive — uniform predicate.
|
||||
float local_sum = 0.0f;
|
||||
if (tid < bdim) {
|
||||
unsigned int c = channels_in_bucket[bucket * MAX_BUCKET_DIM + tid];
|
||||
// Defensive: sentinel slot OR out-of-range channel index should
|
||||
// not contribute. Predicate is uniform across the warp because all
|
||||
// active threads (tid < bdim) hold valid entries by construction
|
||||
// of channels_in_bucket_kernel.
|
||||
// Defensive: out-of-range channel index should not contribute.
|
||||
// Predicate is uniform across active lanes (tid < bdim) — all hold
|
||||
// valid entries by construction of channels_in_bucket_kernel.
|
||||
if (c < (unsigned int)HIDDEN_DIM) {
|
||||
for (int b = 0; b < B; ++b) {
|
||||
float v = h_state[b * HIDDEN_DIM + c];
|
||||
@@ -433,11 +437,12 @@ extern "C" __global__ void h_mag_per_bucket_kernel(
|
||||
}
|
||||
}
|
||||
}
|
||||
sdata[tid] = (tid < 32) ? local_sum : 0.0f;
|
||||
sdata[tid] = local_sum;
|
||||
__syncthreads();
|
||||
|
||||
// Block-tree reduction over 32 lanes (single warp). No atomicAdd.
|
||||
for (int s = 16; s > 0; s >>= 1) {
|
||||
// Block-tree reduction over 128 lanes (multi-warp). No atomicAdd.
|
||||
// Stages: 64→32→16→8→4→2→1. Each stage halves the active lane count.
|
||||
for (int s = 64; s > 0; s >>= 1) {
|
||||
if (tid < s) sdata[tid] += sdata[tid + s];
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
@@ -19,13 +19,26 @@ extern "C" __global__ void compute_advantage_return(
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
const float gamma = isv[RL_GAMMA_INDEX];
|
||||
const float r = rewards[b];
|
||||
const float done = dones[b];
|
||||
const float vt = v_t[b];
|
||||
const float vtp1 = v_tp1[b];
|
||||
const float gamma = isv[RL_GAMMA_INDEX];
|
||||
const float r = rewards[b];
|
||||
const float is_done = (dones[b] > 0.5f);
|
||||
const float vt = v_t[b];
|
||||
const float vtp1 = v_tp1[b];
|
||||
|
||||
const float ret = r + gamma * (1.0f - done) * vtp1;
|
||||
returns[b] = ret;
|
||||
advantages[b] = done * (ret - vt);
|
||||
// 2026-05-29: branch-gate (not multiplication-gate) the V_tp1 term
|
||||
// and the (ret - vt) advantage. Multiplication-gating fails on IEEE
|
||||
// `0 * Inf = NaN`: if any batch's encoder produces a non-finite V
|
||||
// prediction early in training (random init outliers), the prior
|
||||
// `done * (ret - vt)` multiplication propagated NaN to ALL non-done
|
||||
// batches' advantages, which then broke compute_advantage_rms (sum
|
||||
// of A² → NaN) and PPO (A/RMS → NaN, ratio×A → NaN, l_pi → NaN).
|
||||
// Branch-gating ensures non-finite vtp1/vt are NEVER mixed into a
|
||||
// non-done batch's advantage. Validated by the smoke bisect at
|
||||
// 10d4614fb (deterministic NaN at step 4 with l_v=6.329 stable —
|
||||
// proving V REGRESSION is fine while V FORWARD has at least one
|
||||
// non-finite output that the 0*Inf trick was promoting to NaN
|
||||
// everywhere). Per `pearl_atomicadd_masks_v_instability`.
|
||||
const float ret = is_done ? r : (r + gamma * vtp1);
|
||||
returns[b] = ret;
|
||||
advantages[b] = is_done ? (ret - vt) : 0.0f;
|
||||
}
|
||||
|
||||
@@ -24,8 +24,6 @@
|
||||
|
||||
#define BOOK_LEVELS 10
|
||||
#define REGIME_DIM 6
|
||||
#define N_HORIZONS 3
|
||||
#define N_LABEL_SOURCES 5
|
||||
|
||||
// Must match #[repr(C)] Mbp10RawInput in snap_features.rs (216 bytes).
|
||||
struct __align__(8) Mbp10Raw {
|
||||
@@ -82,26 +80,11 @@ extern "C" __global__ void gpu_sample_and_gather(
|
||||
long long* __restrict__ ts_ns_soa, // [B*K]
|
||||
long long* __restrict__ prev_ts_ns_soa, // [B*K]
|
||||
// Global-memory outputs for the sampled (file_offset, anchor, file_idx)
|
||||
// per batch element. Read by gpu_gather_next and gpu_gather_frd_labels
|
||||
// to reuse the same sampling decision.
|
||||
// per batch element. Read by gpu_gather_next, gpu_gather_frd_labels,
|
||||
// and gpu_gather_bce_labels to reuse the same sampling decision.
|
||||
int* __restrict__ out_file_offset, // [B]
|
||||
int* __restrict__ out_anchor, // [B]
|
||||
int* __restrict__ out_file_idx, // [B]
|
||||
// Fused label gather — 5 label sources (BCE, prof_long, prof_short,
|
||||
// size_long, size_short), each [N_HORIZONS × total_snaps]. Reads in
|
||||
// the same pass as snapshot gather to avoid 5 separate random-access
|
||||
// kernel launches that consumed 95.6% of GPU time.
|
||||
int total_snaps,
|
||||
const float* __restrict__ labels_src_0, // [N_HORIZONS × total_snaps]
|
||||
const float* __restrict__ labels_src_1,
|
||||
const float* __restrict__ labels_src_2,
|
||||
const float* __restrict__ labels_src_3,
|
||||
const float* __restrict__ labels_src_4,
|
||||
float* __restrict__ labels_out_0, // [K, B, N_HORIZONS]
|
||||
float* __restrict__ labels_out_1,
|
||||
float* __restrict__ labels_out_2,
|
||||
float* __restrict__ labels_out_3,
|
||||
float* __restrict__ labels_out_4,
|
||||
int B
|
||||
) {
|
||||
int b = blockIdx.x;
|
||||
@@ -165,26 +148,9 @@ extern "C" __global__ void gpu_sample_and_gather(
|
||||
tc_soa[n] = (int)snap.trade_count;
|
||||
ts_ns_soa[n] = (long long)snap.ts_ns;
|
||||
prev_ts_ns_soa[n] = (long long)snap.prev_ts_ns;
|
||||
|
||||
// Fused label gather — same global_idx, zero extra random access.
|
||||
// Output layout: [K, B, N_HORIZONS] row-major.
|
||||
int lbl_base = (k * B + b) * N_HORIZONS;
|
||||
const float* srcs[N_LABEL_SOURCES] = {
|
||||
labels_src_0, labels_src_1, labels_src_2, labels_src_3, labels_src_4
|
||||
};
|
||||
float* outs[N_LABEL_SOURCES] = {
|
||||
labels_out_0, labels_out_1, labels_out_2, labels_out_3, labels_out_4
|
||||
};
|
||||
#pragma unroll
|
||||
for (int s = 0; s < N_LABEL_SOURCES; s++) {
|
||||
#pragma unroll
|
||||
for (int h = 0; h < N_HORIZONS; h++) {
|
||||
outs[s][lbl_base + h] = srcs[s][h * total_snaps + global_idx];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// ── Label gather kernel (STANDALONE — kept for backward compat) ─────
|
||||
// ── Label gather kernel ──────────────────────────────────────────────
|
||||
//
|
||||
// Runs AFTER gpu_sample_and_gather. Gathers per-horizon FRD labels at
|
||||
// the anchor position (rightmost K position = newest snapshot in the
|
||||
|
||||
@@ -36,20 +36,3 @@ extern "C" __global__ void grad_h_accumulate_scaled(
|
||||
if (i >= n) return;
|
||||
grad_h_encoder[i] += lambda * grad_h_head[i];
|
||||
}
|
||||
|
||||
// Mega-graph variant: reads lambda from ISV device pointer at given slot.
|
||||
// This allows the lambda to vary across graph replays (ISV is modified
|
||||
// in-place by controller kernels captured earlier in the same graph).
|
||||
extern "C" __global__ void grad_h_accumulate_scaled_isv(
|
||||
const float* __restrict__ grad_h_head,
|
||||
const float* __restrict__ isv,
|
||||
int lambda_slot,
|
||||
float batch_inv, // 1.0 / b_size
|
||||
int n,
|
||||
float* __restrict__ grad_h_encoder
|
||||
) {
|
||||
const int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i >= n) return;
|
||||
const float lambda = isv[lambda_slot] * batch_inv;
|
||||
grad_h_encoder[i] += lambda * grad_h_head[i];
|
||||
}
|
||||
|
||||
84
crates/ml-alpha/cuda/rl_advantage_normalize.cu
Normal file
84
crates/ml-alpha/cuda/rl_advantage_normalize.cu
Normal file
@@ -0,0 +1,84 @@
|
||||
// rl_advantage_normalize.cu — Phase 4.5 per-batch advantage normalization (2026-05-30).
|
||||
//
|
||||
// Standard PPO practice (Schulman et al. 2017): normalize per-batch
|
||||
// advantages before the PPO surrogate computation:
|
||||
//
|
||||
// mean = (1/B) Σ_b advantage[b]
|
||||
// var = (1/B) Σ_b (advantage[b] − mean)²
|
||||
// std = sqrt(var + ε²)
|
||||
// advantage_norm[b] = (advantage[b] − mean) / std (in-place)
|
||||
//
|
||||
// Why: PPO surrogate = ratio × A. Without normalization, |A| can vary
|
||||
// 1000× across batches → l_pi magnitude unstable → Adam's per-parameter
|
||||
// scaling still functional, but gradient direction confidence drops
|
||||
// when advantage magnitudes are wildly inconsistent.
|
||||
//
|
||||
// In Phase 4.3, V_dq baseline produced advantages with ~100× more
|
||||
// variance than Plan A v2's V_scalar baseline → l_pi grew to 1.6e9
|
||||
// vs Plan A v2's 3.6e5. Adam internally normalizes, but the policy
|
||||
// updates have higher variance per step → pnl trajectory chops.
|
||||
//
|
||||
// Per pearl_adaptive_not_tuned: this normalization is self-adaptive
|
||||
// (uses observed per-batch statistics, no tuned hyperparameters).
|
||||
// Per pearl_blend_formulas_must_have_permanent_floor: ε² floor on
|
||||
// variance prevents div-by-zero when all advantages are identical.
|
||||
//
|
||||
// Block layout: grid=(1, 1, 1), block=(BLOCK_X=1024, 1, 1). Single
|
||||
// block does parallel reduction for mean → variance → normalize.
|
||||
// Suitable for batch sizes up to B=1024 (current production scale).
|
||||
|
||||
#define BLOCK_X 1024
|
||||
#define ADVANTAGE_VAR_FLOOR 1e-6f
|
||||
|
||||
extern "C" __global__ void rl_advantage_normalize(
|
||||
float* __restrict__ advantage, // [B] in-place
|
||||
int B
|
||||
) {
|
||||
const int tid = threadIdx.x;
|
||||
if (tid >= BLOCK_X) return;
|
||||
|
||||
// ── Pass 1: compute mean ──
|
||||
__shared__ float s_sum[BLOCK_X];
|
||||
float sum_partial = 0.0f;
|
||||
for (int b = tid; b < B; b += BLOCK_X) {
|
||||
sum_partial += advantage[b];
|
||||
}
|
||||
s_sum[tid] = sum_partial;
|
||||
__syncthreads();
|
||||
|
||||
for (int stride = BLOCK_X / 2; stride > 0; stride >>= 1) {
|
||||
if (tid < stride) s_sum[tid] += s_sum[tid + stride];
|
||||
__syncthreads();
|
||||
}
|
||||
__shared__ float s_mean;
|
||||
if (tid == 0) s_mean = s_sum[0] / (float)B;
|
||||
__syncthreads();
|
||||
const float mean = s_mean;
|
||||
|
||||
// ── Pass 2: compute variance ──
|
||||
__shared__ float s_var[BLOCK_X];
|
||||
float var_partial = 0.0f;
|
||||
for (int b = tid; b < B; b += BLOCK_X) {
|
||||
const float d = advantage[b] - mean;
|
||||
var_partial += d * d;
|
||||
}
|
||||
s_var[tid] = var_partial;
|
||||
__syncthreads();
|
||||
|
||||
for (int stride = BLOCK_X / 2; stride > 0; stride >>= 1) {
|
||||
if (tid < stride) s_var[tid] += s_var[tid + stride];
|
||||
__syncthreads();
|
||||
}
|
||||
__shared__ float s_inv_std;
|
||||
if (tid == 0) {
|
||||
const float var = s_var[0] / (float)B;
|
||||
s_inv_std = rsqrtf(var + ADVANTAGE_VAR_FLOOR); // 1/sqrt(var + ε²)
|
||||
}
|
||||
__syncthreads();
|
||||
const float inv_std = s_inv_std;
|
||||
|
||||
// ── Pass 3: normalize in-place ──
|
||||
for (int b = tid; b < B; b += BLOCK_X) {
|
||||
advantage[b] = (advantage[b] - mean) * inv_std;
|
||||
}
|
||||
}
|
||||
@@ -34,8 +34,6 @@
|
||||
#define RL_HOLD_TARGET_FRAC_INDEX 575
|
||||
#define RL_HOLD_FRAC_EMA_INDEX 576
|
||||
#define RL_CONF_GATE_MAX_HOLD_FRAC_INDEX 584
|
||||
#define RL_STOP_LOSS_THRESHOLD_INDEX 587
|
||||
#define RL_MEAN_ABS_PNL_EMA_INDEX 423
|
||||
#define THRESHOLD_MIN 0.05f
|
||||
#define THRESHOLD_MAX 0.95f
|
||||
#define HOLD_EMA_ALPHA 0.1f
|
||||
@@ -53,35 +51,6 @@ extern "C" __global__ void rl_confidence_gate(
|
||||
const int b = blockIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
// ── Hard stop-loss: override action to Flat when unrealized loss
|
||||
// exceeds ISV-driven threshold. Fires BEFORE gate logic so the
|
||||
// lobsim actually closes the position (no state mismatch).
|
||||
// unrealized_loss = |vwap - mid| × |lots|; normalized by mean_abs_pnl.
|
||||
{
|
||||
const unsigned char* p = pos_state + b * pos_bytes;
|
||||
const int lots = *(const int*)(p + 0);
|
||||
if (lots != 0) {
|
||||
const float stop_thresh = isv[RL_STOP_LOSS_THRESHOLD_INDEX];
|
||||
const float mean_abs = fmaxf(isv[RL_MEAN_ABS_PNL_EMA_INDEX], 1e-6f);
|
||||
// Approximate mid from the first bid/ask level would need
|
||||
// book data; instead use the realized_pnl delta as a proxy
|
||||
// for position health. Simpler: use vwap vs realized_pnl trend.
|
||||
// For now, use the unit_entry_price-based unrealized_r from
|
||||
// trade_context — but that runs AFTER this kernel.
|
||||
//
|
||||
// Practical approach: read peak_equity (offset 12) and
|
||||
// realized_pnl (offset 8). Drawdown from peak = peak - realized.
|
||||
const float realized = *(const float*)(p + 8);
|
||||
const float peak = *(const float*)(p + 12);
|
||||
const float drawdown = peak - realized;
|
||||
const float normalized_dd = drawdown / mean_abs;
|
||||
if (normalized_dd > stop_thresh) {
|
||||
actions[b] = (lots > 0) ? 3 : 4; // FlatFromLong / FlatFromShort
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX];
|
||||
const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX];
|
||||
if (current_step < warmup) return;
|
||||
|
||||
@@ -1,49 +0,0 @@
|
||||
// rl_drawdown_stop.cu — per-step drawdown penalty + hard stop-loss.
|
||||
//
|
||||
// Runs AFTER rl_trade_context_update (which computes unrealized_r).
|
||||
// Reads unrealized_r from trade_context_d and applies:
|
||||
//
|
||||
// 1. Per-step drawdown penalty: adds min(0, unrealized_r) × PENALTY_RATE
|
||||
// to rewards[b]. Creates continuous exit gradient on losing positions.
|
||||
// Penalty-only (never positive) = no exposure incentive.
|
||||
//
|
||||
// 2. Hard stop-loss: when unrealized_r < -STOP_THRESHOLD, force-close
|
||||
// by setting dones[b] = 1. The realized PnL at this point becomes
|
||||
// the final trade reward. Caps max loss per trade.
|
||||
//
|
||||
// Both thresholds are ISV-driven per feedback_isv_for_adaptive_bounds.
|
||||
// Grid: ceil(B/32), Block: 32.
|
||||
|
||||
#define RL_DRAWDOWN_PENALTY_RATE_INDEX 586
|
||||
#define RL_STOP_LOSS_THRESHOLD_INDEX 587
|
||||
#define TRADE_CONTEXT_STRIDE 5
|
||||
#define UNREALIZED_R_OFFSET 1
|
||||
|
||||
extern "C" __global__ void rl_drawdown_stop(
|
||||
const float* __restrict__ trade_context, // [B, TRADE_CONTEXT_STRIDE]
|
||||
float* __restrict__ rewards, // [B] IN/OUT
|
||||
float* __restrict__ dones, // [B] IN/OUT
|
||||
float* __restrict__ isv,
|
||||
int b_size
|
||||
) {
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
const float unrealized_r = trade_context[b * TRADE_CONTEXT_STRIDE + UNREALIZED_R_OFFSET];
|
||||
const float penalty_rate = isv[RL_DRAWDOWN_PENALTY_RATE_INDEX];
|
||||
const float stop_thresh = isv[RL_STOP_LOSS_THRESHOLD_INDEX];
|
||||
|
||||
// Per-step drawdown penalty: only fires when losing (unrealized_r < 0).
|
||||
// min(0, x) * rate is always <= 0.
|
||||
if (unrealized_r < 0.0f) {
|
||||
rewards[b] += unrealized_r * penalty_rate;
|
||||
}
|
||||
|
||||
// Hard stop-loss DISABLED: setting dones=1 here creates a state
|
||||
// mismatch — the lobsim position stays open while Q sees a false
|
||||
// close. The done signal must come from actual position changes
|
||||
// in the lobsim, not synthetic overrides. To implement hard
|
||||
// stop-loss properly, the action must be overridden to FlatL/FlatS
|
||||
// BEFORE the lobsim step, not after.
|
||||
(void) stop_thresh;
|
||||
}
|
||||
57
crates/ml-alpha/cuda/rl_dueling_q_bellman_target.cu
Normal file
57
crates/ml-alpha/cuda/rl_dueling_q_bellman_target.cu
Normal file
@@ -0,0 +1,57 @@
|
||||
// rl_dueling_q_bellman_target.cu — Phase 4 Bellman target build.
|
||||
//
|
||||
// Picks argmax_a' over the target net's composed_Q at s_{t+1}, then
|
||||
// builds the scalar Bellman target:
|
||||
//
|
||||
// a*_b = argmax_a' target_composed_Q[b, a']
|
||||
// target_value[b] = r[b] + γ × (1 − done[b]) × target_composed_Q[b, a*_b]
|
||||
//
|
||||
// γ read from ISV bus at runtime — same source as C51 / IQN Bellman
|
||||
// target kernels (RL_GAMMA_INDEX=400). For n-step returns, the
|
||||
// per-batch n_step_gammas are passed (matches existing PER convention).
|
||||
//
|
||||
// Block layout:
|
||||
// grid = (B, 1, 1)
|
||||
// block = (N_ACTIONS, 1, 1)
|
||||
// One block per batch. Each thread holds one action's value.
|
||||
// Tree-reduce in shared mem to find argmax. Thread 0 computes the
|
||||
// final target.
|
||||
//
|
||||
// Per feedback_no_atomicadd: sole-writer per output cell.
|
||||
|
||||
#define N_ACTIONS 11
|
||||
|
||||
extern "C" __global__ void rl_dueling_q_bellman_target_build(
|
||||
const float* __restrict__ target_composed_q, // [B × N_ACTIONS]
|
||||
const float* __restrict__ rewards, // [B]
|
||||
const float* __restrict__ dones, // [B]
|
||||
const float* __restrict__ n_step_gammas, // [B] (γ^n_step per sample)
|
||||
int B,
|
||||
float* __restrict__ target_value_out // [B]
|
||||
) {
|
||||
const int b = blockIdx.x;
|
||||
const int a = threadIdx.x;
|
||||
if (b >= B) return;
|
||||
if (a >= N_ACTIONS) return;
|
||||
|
||||
__shared__ float s_q[N_ACTIONS];
|
||||
s_q[a] = target_composed_q[b * N_ACTIONS + a];
|
||||
__syncthreads();
|
||||
|
||||
if (a == 0) {
|
||||
// Find argmax serially (small N).
|
||||
float best = s_q[0];
|
||||
#pragma unroll
|
||||
for (int i = 1; i < N_ACTIONS; ++i) {
|
||||
if (s_q[i] > best) best = s_q[i];
|
||||
}
|
||||
|
||||
const float r = rewards[b];
|
||||
const float done = dones[b];
|
||||
const float gamma_n = n_step_gammas[b];
|
||||
|
||||
// Standard Bellman with done masking:
|
||||
// target = r + γ^n × (1 − done) × max_q
|
||||
target_value_out[b] = r + gamma_n * (1.0f - done) * best;
|
||||
}
|
||||
}
|
||||
89
crates/ml-alpha/cuda/rl_dueling_q_decompose_and_bwd.cu
Normal file
89
crates/ml-alpha/cuda/rl_dueling_q_decompose_and_bwd.cu
Normal file
@@ -0,0 +1,89 @@
|
||||
// rl_dueling_q_decompose_and_bwd.cu — Phase 4 decompose grad + weight grad.
|
||||
//
|
||||
// Decompose:
|
||||
// composed_Q[b, a] = V[b] + A[b, a] − (1/N) Σ_a' A[b, a']
|
||||
//
|
||||
// Chain rule (only taken action has nonzero grad_composed[b, a]):
|
||||
// grad_V[b] = Σ_a grad_composed[b, a] = grad_composed[b, a_taken]
|
||||
// grad_A[b, a] = grad_composed[b, a] − (1/N) × grad_V[b]
|
||||
//
|
||||
// For a == a_taken: grad_A[b, a] = (1 − 1/N) × grad_composed[b, a_taken]
|
||||
// For a ≠ a_taken: grad_A[b, a] = −(1/N) × grad_composed[b, a_taken]
|
||||
//
|
||||
// After decompose, compute per-batch weight gradients via outer
|
||||
// product with h_t:
|
||||
// grad_w_v_pb[b, c] = grad_V[b] × h_t[b, c]
|
||||
// grad_b_v_pb[b] = grad_V[b]
|
||||
// grad_w_a_pb[b, c, a] = grad_A[b, a] × h_t[b, c]
|
||||
// grad_b_a_pb[b, a] = grad_A[b, a]
|
||||
//
|
||||
// Caller reduces per-batch grads via reduce_axis0 to final shapes:
|
||||
// grad_w_v [HIDDEN_DIM]
|
||||
// grad_b_v [1]
|
||||
// grad_w_a [HIDDEN_DIM × N_ACTIONS]
|
||||
// grad_b_a [N_ACTIONS]
|
||||
//
|
||||
// Block layout:
|
||||
// grid = (B, 1, 1)
|
||||
// block = (HIDDEN_DIM, 1, 1) — one thread per hidden-dim index
|
||||
// Each thread loops over N_ACTIONS to write A weights, plus the
|
||||
// single V weight. Thread 0 additionally writes grad_b_v_pb +
|
||||
// grad_b_a_pb (small).
|
||||
//
|
||||
// Per feedback_no_atomicadd: sole-writer per (b, c, a) and (b, c) cells.
|
||||
|
||||
#define HIDDEN_DIM 128
|
||||
#define N_ACTIONS 11
|
||||
|
||||
extern "C" __global__ void rl_dueling_q_decompose_and_weight_grad(
|
||||
const float* __restrict__ h_t, // [B × HIDDEN_DIM]
|
||||
const float* __restrict__ grad_composed, // [B × N_ACTIONS] (only taken cell nonzero)
|
||||
const int* __restrict__ actions_taken, // [B]
|
||||
int B,
|
||||
float* __restrict__ grad_w_v_pb, // [B × HIDDEN_DIM]
|
||||
float* __restrict__ grad_b_v_pb, // [B]
|
||||
float* __restrict__ grad_w_a_pb, // [B × HIDDEN_DIM × N_ACTIONS]
|
||||
float* __restrict__ grad_b_a_pb // [B × N_ACTIONS]
|
||||
) {
|
||||
const int b = blockIdx.x;
|
||||
const int c = threadIdx.x;
|
||||
if (b >= B) return;
|
||||
if (c >= HIDDEN_DIM) return;
|
||||
|
||||
int a_t = actions_taken[b];
|
||||
if (a_t < 0) a_t = 0;
|
||||
if (a_t >= N_ACTIONS) a_t = 0;
|
||||
|
||||
// grad_composed is nonzero only at a == a_t.
|
||||
const float gc_taken = grad_composed[b * N_ACTIONS + a_t];
|
||||
|
||||
// grad_V[b] = Σ_a grad_composed[b, a] = gc_taken (others are 0).
|
||||
const float grad_v = gc_taken;
|
||||
|
||||
// h_t[b, c].
|
||||
const float h_bc = h_t[b * HIDDEN_DIM + c];
|
||||
|
||||
// Per-batch V weight grad.
|
||||
grad_w_v_pb[b * HIDDEN_DIM + c] = grad_v * h_bc;
|
||||
|
||||
// Per-batch A weight grad (one thread writes N_ACTIONS values).
|
||||
// grad_A[b, a] = grad_composed[b, a] − (1/N) × grad_V[b]
|
||||
// = (a == a_t ? gc_taken : 0) − (1/N) × gc_taken
|
||||
const float inv_N = 1.0f / (float)N_ACTIONS;
|
||||
#pragma unroll
|
||||
for (int a = 0; a < N_ACTIONS; ++a) {
|
||||
const float grad_a = (a == a_t ? gc_taken : 0.0f) - inv_N * gc_taken;
|
||||
// Layout: grad_w_a_pb[b, c, a] = h_bc * grad_a
|
||||
grad_w_a_pb[b * HIDDEN_DIM * N_ACTIONS + c * N_ACTIONS + a] = h_bc * grad_a;
|
||||
}
|
||||
|
||||
// Thread 0 writes biases (small).
|
||||
if (c == 0) {
|
||||
grad_b_v_pb[b] = grad_v;
|
||||
#pragma unroll
|
||||
for (int a = 0; a < N_ACTIONS; ++a) {
|
||||
const float grad_a = (a == a_t ? gc_taken : 0.0f) - inv_N * gc_taken;
|
||||
grad_b_a_pb[b * N_ACTIONS + a] = grad_a;
|
||||
}
|
||||
}
|
||||
}
|
||||
121
crates/ml-alpha/cuda/rl_dueling_q_forward.cu
Normal file
121
crates/ml-alpha/cuda/rl_dueling_q_forward.cu
Normal file
@@ -0,0 +1,121 @@
|
||||
// rl_dueling_q_forward.cu — Phase 4 Independent Dueling Q head forward.
|
||||
//
|
||||
// Per spec docs/superpowers/specs/2026-05-30-phase4-independent-dueling-head-design.md.
|
||||
//
|
||||
// Architecture: parallel head to C51/IQN/π/value_head with ZERO shared
|
||||
// state with downstream consumers (ensemble, distill, action selection).
|
||||
// Trains its own V + A weights via Bellman loss on composed_Q at taken
|
||||
// action. V output feeds PPO advantage baseline (Phase 4.3) — composed_Q
|
||||
// is internal to this head's loss path and never consumed elsewhere.
|
||||
//
|
||||
// Why this design (vs Phase 2 v2 / Phase 3.x failures):
|
||||
// - Phase 2 v2 (scalar V + categorical CE): N/A here — uses scalar
|
||||
// Bellman loss, no softmax math eating V.
|
||||
// - Phase 3.2 (V_IQN cross-architecture calibration): V_dq trained
|
||||
// on same Bellman reward signal as C51, scales align by construction.
|
||||
// - Phase 3.1 (composed Q perturbs ensemble): composed_Q_dq never
|
||||
// feeds ensemble. Only feeds its own Bellman loss + diag.
|
||||
// - Phase 3.1-fix (gradient structure mismatch contaminates weights):
|
||||
// DuelingQHead has its OWN weights — mean-zero A grad pattern only
|
||||
// affects DuelingQHead's training, no shared weights with C51/IQN.
|
||||
//
|
||||
// Forward computation:
|
||||
// V[b] = Σ_c w_v[c] × h_t[b, c] + b_v[0]
|
||||
// A[b, a] = Σ_c w_a[c, a] × h_t[b, c] + b_a[a] for a in 0..N
|
||||
// composed_Q[b, a] = V[b] + A[b, a] − (1/N) Σ_a' A[b, a']
|
||||
//
|
||||
// Block layout:
|
||||
// grid = (B, 1, 1)
|
||||
// block = (HIDDEN_DIM = 128, 1, 1)
|
||||
// One block per batch. Each thread computes one hidden-dim term and
|
||||
// participates in tree-reduce. Then thread 0 broadcasts to A
|
||||
// computation. Each thread computes one action's matmul contribution.
|
||||
// Mean reduction over N_ACTIONS done in shared mem.
|
||||
//
|
||||
// Memory:
|
||||
// shared float s_h[HIDDEN_DIM] — cached h_t row
|
||||
// shared float s_v — scalar V value
|
||||
// shared float s_a[N_ACTIONS] — A values (post-bias)
|
||||
//
|
||||
// Per feedback_no_atomicadd: sole-writer per output cell.
|
||||
// Per feedback_cpu_is_read_only: pure device kernel.
|
||||
|
||||
#define HIDDEN_DIM 128
|
||||
#define N_ACTIONS 11
|
||||
|
||||
extern "C" __global__ void rl_dueling_q_forward(
|
||||
const float* __restrict__ h_t, // [B × HIDDEN_DIM]
|
||||
const float* __restrict__ w_v, // [HIDDEN_DIM]
|
||||
const float* __restrict__ b_v, // [1]
|
||||
const float* __restrict__ w_a, // [HIDDEN_DIM × N_ACTIONS], row-major
|
||||
const float* __restrict__ b_a, // [N_ACTIONS]
|
||||
int B,
|
||||
float* __restrict__ v_out, // [B]
|
||||
float* __restrict__ a_out, // [B × N_ACTIONS]
|
||||
float* __restrict__ q_composed_out // [B × N_ACTIONS]
|
||||
) {
|
||||
const int b = blockIdx.x;
|
||||
const int c = threadIdx.x;
|
||||
if (b >= B) return;
|
||||
if (c >= HIDDEN_DIM) return;
|
||||
|
||||
extern __shared__ float s_h[]; // [HIDDEN_DIM], sized by smem arg
|
||||
|
||||
// ── Load h_t[b] into shared mem ──
|
||||
s_h[c] = h_t[b * HIDDEN_DIM + c];
|
||||
__syncthreads();
|
||||
|
||||
// ── V projection (block-reduce) ──
|
||||
// V[b] = Σ_c w_v[c] × s_h[c] + b_v[0]
|
||||
// Tree reduction over HIDDEN_DIM threads.
|
||||
__shared__ float s_v_partial[HIDDEN_DIM];
|
||||
s_v_partial[c] = w_v[c] * s_h[c];
|
||||
__syncthreads();
|
||||
// Tree reduce
|
||||
for (int stride = HIDDEN_DIM / 2; stride > 0; stride >>= 1) {
|
||||
if (c < stride) {
|
||||
s_v_partial[c] += s_v_partial[c + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
__shared__ float s_v;
|
||||
if (c == 0) {
|
||||
s_v = s_v_partial[0] + b_v[0];
|
||||
v_out[b] = s_v;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// ── A projection (each thread handles one action) ──
|
||||
// A[b, a] = Σ_c w_a[c, a] × s_h[c] + b_a[a]
|
||||
// Threads 0..N_ACTIONS-1 compute one action each.
|
||||
// Other threads idle for this section.
|
||||
__shared__ float s_a[N_ACTIONS];
|
||||
if (c < N_ACTIONS) {
|
||||
float acc = 0.0f;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < HIDDEN_DIM; ++i) {
|
||||
acc += w_a[i * N_ACTIONS + c] * s_h[i];
|
||||
}
|
||||
acc += b_a[c];
|
||||
s_a[c] = acc;
|
||||
a_out[b * N_ACTIONS + c] = acc;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// ── Mean over actions (thread 0 only — small N) ──
|
||||
__shared__ float s_mean_a;
|
||||
if (c == 0) {
|
||||
float sum = 0.0f;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < N_ACTIONS; ++i) {
|
||||
sum += s_a[i];
|
||||
}
|
||||
s_mean_a = sum * (1.0f / (float)N_ACTIONS);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// ── Compose Q (each thread for one action writes the result) ──
|
||||
if (c < N_ACTIONS) {
|
||||
q_composed_out[b * N_ACTIONS + c] = s_v + s_a[c] - s_mean_a;
|
||||
}
|
||||
}
|
||||
83
crates/ml-alpha/cuda/rl_dueling_q_loss_and_grad.cu
Normal file
83
crates/ml-alpha/cuda/rl_dueling_q_loss_and_grad.cu
Normal file
@@ -0,0 +1,83 @@
|
||||
// rl_dueling_q_loss_and_grad.cu — Phase 4 Bellman loss + decompose backward.
|
||||
//
|
||||
// Computes the scalar Huber loss on (target − online_composed_Q[taken])
|
||||
// AND emits grad_composed[B × N_ACTIONS] in one fused kernel.
|
||||
//
|
||||
// Loss (per-batch):
|
||||
// δ_b = target_value[b] − online_composed_Q[b, a_taken[b]]
|
||||
// L_b = Huber(δ_b, κ=1.0) / B (mean over batch)
|
||||
// loss_per_batch[b] = L_b
|
||||
//
|
||||
// Gradient w.r.t. online_composed_Q (mostly zero, nonzero at taken):
|
||||
// For a_taken: dL/d_composed[b, a_taken] = −(1/B) × Huber'(δ_b)
|
||||
// For a ≠ a_taken: dL/d_composed[b, a] = 0
|
||||
//
|
||||
// Huber'(δ) = δ if |δ| ≤ κ
|
||||
// = κ × sign(δ) if |δ| > κ
|
||||
//
|
||||
// Block layout:
|
||||
// grid = (B, 1, 1)
|
||||
// block = (N_ACTIONS, 1, 1)
|
||||
// Each thread writes one (b, a) cell of grad_composed; thread 0
|
||||
// computes loss + the nonzero gradient scalar.
|
||||
//
|
||||
// Per feedback_no_atomicadd: sole-writer per cell.
|
||||
|
||||
#define N_ACTIONS 11
|
||||
#define HUBER_KAPPA 1.0f
|
||||
|
||||
extern "C" __global__ void rl_dueling_q_loss_and_grad(
|
||||
const float* __restrict__ online_composed_q, // [B × N_ACTIONS]
|
||||
const float* __restrict__ target_value, // [B]
|
||||
const int* __restrict__ actions_taken, // [B]
|
||||
int B,
|
||||
float* __restrict__ loss_per_batch, // [B]
|
||||
float* __restrict__ grad_composed // [B × N_ACTIONS]
|
||||
) {
|
||||
const int b = blockIdx.x;
|
||||
const int a = threadIdx.x;
|
||||
if (b >= B) return;
|
||||
if (a >= N_ACTIONS) return;
|
||||
|
||||
// Defensive clamp on actions_taken (trainer should never produce
|
||||
// out-of-range, but guard against corrupt indices).
|
||||
int a_t = actions_taken[b];
|
||||
if (a_t < 0) a_t = 0;
|
||||
if (a_t >= N_ACTIONS) a_t = 0;
|
||||
|
||||
__shared__ float s_grad_scalar;
|
||||
|
||||
if (a == 0) {
|
||||
const float q_taken = online_composed_q[b * N_ACTIONS + a_t];
|
||||
const float target = target_value[b];
|
||||
const float delta = target - q_taken;
|
||||
const float abs_d = fabsf(delta);
|
||||
const float inv_B = 1.0f / (float)B;
|
||||
|
||||
// Huber loss.
|
||||
float l;
|
||||
if (abs_d <= HUBER_KAPPA) {
|
||||
l = 0.5f * delta * delta;
|
||||
} else {
|
||||
l = HUBER_KAPPA * (abs_d - 0.5f * HUBER_KAPPA);
|
||||
}
|
||||
loss_per_batch[b] = l * inv_B;
|
||||
|
||||
// Huber'(δ) — gradient of l w.r.t. delta.
|
||||
float huber_grad;
|
||||
if (abs_d <= HUBER_KAPPA) {
|
||||
huber_grad = delta;
|
||||
} else {
|
||||
huber_grad = HUBER_KAPPA * ((delta > 0.0f) ? 1.0f : -1.0f);
|
||||
}
|
||||
// dL/d_composed[a_taken] = (dL/dδ) × (dδ/d_composed[a_taken])
|
||||
// = (−1/B) × huber_grad (δ = target − Q)
|
||||
s_grad_scalar = -inv_B * huber_grad;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// All threads write grad_composed. Only the taken action gets a
|
||||
// nonzero value (CE on a single Q value).
|
||||
const int idx = b * N_ACTIONS + a;
|
||||
grad_composed[idx] = (a == a_t) ? s_grad_scalar : 0.0f;
|
||||
}
|
||||
@@ -1,134 +0,0 @@
|
||||
// rl_lr_from_mapped_pinned.cu — GPU-side LR controller that reads loss
|
||||
// observations from mapped-pinned device pointers instead of host-
|
||||
// supplied scalar arguments. This eliminates the host-side loss readback
|
||||
// between the reward graph and training graph, enabling mega-graph
|
||||
// capture of the entire per-step pipeline.
|
||||
//
|
||||
// The mapped-pinned buffers (ss_q_loss_mapped, ss_pi_loss_mapped,
|
||||
// ss_v_loss_sum_mapped) have stable device pointers. The GPU writes
|
||||
// loss values during backward kernels; this kernel reads them from the
|
||||
// same physical memory via the device-side pointer. The reads are
|
||||
// coherent because this kernel launches AFTER the backward kernels on
|
||||
// the same stream (stream ordering guarantees visibility).
|
||||
//
|
||||
// Internally delegates to the same plateau_decay_head logic as
|
||||
// rl_lr_controller.cu. The only difference is the loss observation
|
||||
// source: pointer dereference instead of scalar argument.
|
||||
|
||||
#define RL_LR_BCE_INDEX 412
|
||||
#define RL_LR_Q_INDEX 413
|
||||
#define RL_LR_PI_INDEX 414
|
||||
#define RL_LR_V_INDEX 415
|
||||
#define RL_LR_AUX_INDEX 416
|
||||
|
||||
#define RL_LR_BOOTSTRAP_INDEX 462
|
||||
#define RL_LR_MIN_INDEX 463
|
||||
#define RL_LR_MAX_INDEX 464
|
||||
#define RL_LR_LOSS_EMA_ALPHA_INDEX 465
|
||||
#define RL_LR_DECAY_FACTOR_INDEX 466
|
||||
|
||||
#define RL_IMPROVEMENT_THRESHOLD_INDEX 455
|
||||
#define RL_PLATEAU_PATIENCE_INDEX 456
|
||||
#define RL_LR_WARMUP_STEPS_INDEX 461
|
||||
|
||||
__device__ __forceinline__ void plateau_decay_head(
|
||||
float* isv,
|
||||
int lr_idx,
|
||||
float observed_loss,
|
||||
int loss_ema_slot,
|
||||
int best_slot,
|
||||
int counter_slot,
|
||||
int warmup_slot
|
||||
) {
|
||||
const float lr_prev = isv[lr_idx];
|
||||
const float lr_bootstrap = isv[RL_LR_BOOTSTRAP_INDEX];
|
||||
|
||||
if (lr_prev == 0.0f) {
|
||||
isv[lr_idx] = lr_bootstrap;
|
||||
return;
|
||||
}
|
||||
|
||||
if (observed_loss == 0.0f) return;
|
||||
if (loss_ema_slot < 0) return;
|
||||
|
||||
const float loss_ema_prev = isv[loss_ema_slot];
|
||||
const float loss_ema_alpha = isv[RL_LR_LOSS_EMA_ALPHA_INDEX];
|
||||
float loss_ema_new;
|
||||
if (loss_ema_prev == 0.0f) {
|
||||
loss_ema_new = observed_loss;
|
||||
} else {
|
||||
loss_ema_new = (1.0f - loss_ema_alpha) * loss_ema_prev
|
||||
+ loss_ema_alpha * observed_loss;
|
||||
}
|
||||
isv[loss_ema_slot] = loss_ema_new;
|
||||
|
||||
const float warmup_prev = isv[warmup_slot];
|
||||
const float warmup_target = isv[RL_LR_WARMUP_STEPS_INDEX];
|
||||
if (warmup_prev < warmup_target) {
|
||||
isv[best_slot] = loss_ema_new;
|
||||
isv[counter_slot] = 0.0f;
|
||||
isv[warmup_slot] = warmup_prev + 1.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
const float improvement_threshold = isv[RL_IMPROVEMENT_THRESHOLD_INDEX];
|
||||
const float plateau_patience = isv[RL_PLATEAU_PATIENCE_INDEX];
|
||||
const float best_prev = isv[best_slot];
|
||||
|
||||
if (loss_ema_new < best_prev * improvement_threshold) {
|
||||
isv[best_slot] = loss_ema_new;
|
||||
isv[counter_slot] = 0.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
const float counter_next = isv[counter_slot] + 1.0f;
|
||||
if (counter_next >= plateau_patience) {
|
||||
const float lr_min = isv[RL_LR_MIN_INDEX];
|
||||
const float decay_factor = isv[RL_LR_DECAY_FACTOR_INDEX];
|
||||
const float lr_new = fmaxf(lr_min, lr_prev * decay_factor);
|
||||
isv[lr_idx] = lr_new;
|
||||
isv[counter_slot] = 0.0f;
|
||||
} else {
|
||||
isv[counter_slot] = counter_next;
|
||||
}
|
||||
}
|
||||
|
||||
extern "C" __global__ void rl_lr_from_mapped_pinned(
|
||||
float* __restrict__ isv,
|
||||
const float* __restrict__ q_loss_ptr, // mapped-pinned dev_ptr (1 float)
|
||||
const float* __restrict__ pi_loss_ptr, // mapped-pinned dev_ptr (1 float)
|
||||
const float* __restrict__ v_loss_sum_ptr, // mapped-pinned dev_ptr (1 float)
|
||||
int b_size, // batch size for V loss normalization
|
||||
int q_loss_ema_slot,
|
||||
int q_best_slot,
|
||||
int q_counter_slot,
|
||||
int q_warmup_slot,
|
||||
int pi_loss_ema_slot,
|
||||
int pi_best_slot,
|
||||
int pi_counter_slot,
|
||||
int pi_warmup_slot,
|
||||
int v_loss_ema_slot,
|
||||
int v_best_slot,
|
||||
int v_counter_slot,
|
||||
int v_warmup_slot
|
||||
) {
|
||||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||||
|
||||
// Read losses from mapped-pinned device pointers. These are the
|
||||
// same physical pages the backward kernels wrote to — coherent
|
||||
// because this kernel is stream-ordered after them.
|
||||
const float observed_loss_q = *q_loss_ptr;
|
||||
const float observed_loss_pi = *pi_loss_ptr;
|
||||
// V loss is stored as a sum across the batch; normalize to per-sample.
|
||||
const float observed_loss_v = (b_size > 0) ? (*v_loss_sum_ptr / (float)b_size) : 0.0f;
|
||||
|
||||
// BCE and AUX: perception-owned heads, pass slot=-1 to skip.
|
||||
plateau_decay_head(isv, RL_LR_BCE_INDEX, 0.0f, -1, -1, -1, -1);
|
||||
plateau_decay_head(isv, RL_LR_Q_INDEX, observed_loss_q,
|
||||
q_loss_ema_slot, q_best_slot, q_counter_slot, q_warmup_slot);
|
||||
plateau_decay_head(isv, RL_LR_PI_INDEX, observed_loss_pi,
|
||||
pi_loss_ema_slot, pi_best_slot, pi_counter_slot, pi_warmup_slot);
|
||||
plateau_decay_head(isv, RL_LR_V_INDEX, observed_loss_v,
|
||||
v_loss_ema_slot, v_best_slot, v_counter_slot, v_warmup_slot);
|
||||
plateau_decay_head(isv, RL_LR_AUX_INDEX, 0.0f, -1, -1, -1, -1);
|
||||
}
|
||||
@@ -1,23 +1,25 @@
|
||||
/**
|
||||
* rl_per_push_flush.cu — Two-kernel coordinated coalesced replay write.
|
||||
* rl_per_push_flush.cu — Coordinated coalesced replay write.
|
||||
*
|
||||
* Split from the original single-kernel volatile-spin design into two
|
||||
* kernels with stream-ordered dependency (graph-safe):
|
||||
* Second kernel in the split rl_per_push pipeline. After rl_per_push_ring has
|
||||
* written data to the n-step ring and set flush_flags[b] for completed n-step
|
||||
* transitions, this kernel coordinates slot allocation and performs the actual
|
||||
* coalesced write to the replay buffer.
|
||||
*
|
||||
* Kernel 1: rl_per_flush_prefix_sum
|
||||
* Grid=(1), Block=(1). Single thread computes exclusive prefix-sum of
|
||||
* flush_flags, advances write_head and replay_len, writes per-batch
|
||||
* offsets + base into write_offsets[0..b_size+1].
|
||||
* Architecture:
|
||||
* - Grid=(b_size, 1, 1), Block=(128, 1, 1). One block per batch element.
|
||||
* - Block 0 runs the prefix-sum to assign write slots (single-threaded, O(B)).
|
||||
* - All other blocks spin on a volatile global-memory ready flag (no atomics).
|
||||
* - Once ready, each flushing block does a coalesced 128-thread copy of h_t
|
||||
* and h_tp1, then thread 0 computes the n-step return and writes scalars +
|
||||
* priority + resets the ring count.
|
||||
*
|
||||
* Kernel 2: rl_per_flush_write
|
||||
* Grid=(b_size), Block=(128). Each block reads its pre-computed offset
|
||||
* from write_offsets, then does the coalesced h_t/h_tp1 copy and
|
||||
* scalar write. No spin — the prefix-sum kernel completed before this
|
||||
* launch (stream ordering).
|
||||
*
|
||||
* No atomicAdd — slot assignment is deterministic serial scan.
|
||||
* No volatile spin — correctness from stream ordering / graph node deps.
|
||||
* Coordination: block-0 prefix-sum + volatile ready flag.
|
||||
* No atomicAdd — slot assignment is deterministic serial scan (feedback_no_atomicadd).
|
||||
* Pre-compiled cubin only (feedback_no_nvrtc).
|
||||
*
|
||||
* Launch: CUDA_HOME=/usr/local/cuda nvcc -cubin -arch sm_86 -O3 \
|
||||
* --generate-line-info rl_per_push_flush.cu -o rl_per_push_flush.cubin
|
||||
*/
|
||||
|
||||
#define HIDDEN_DIM 128
|
||||
@@ -28,89 +30,105 @@
|
||||
#define RL_GAMMA_INDEX 400
|
||||
#define RL_N_STEP_INDEX 403
|
||||
|
||||
/* =====================================================================
|
||||
* Kernel 1: prefix-sum + write_head advance.
|
||||
*
|
||||
* Grid=(1,1,1), Block=(1,1,1). Single thread.
|
||||
*
|
||||
* Reads flush_flags[0..b_size], computes exclusive prefix-sum into
|
||||
* write_offsets[0..b_size], stores base write_head at write_offsets[b_size],
|
||||
* and advances the replay buffer's write_head + replay_len.
|
||||
* ===================================================================== */
|
||||
extern "C" __global__ void rl_per_flush_prefix_sum(
|
||||
const int* __restrict__ flush_flags, /* [B] */
|
||||
int* __restrict__ write_offsets, /* [B+1] output: offsets + base */
|
||||
unsigned int* __restrict__ write_head, /* [1] */
|
||||
unsigned int* __restrict__ replay_len, /* [1] */
|
||||
int b_size,
|
||||
int capacity
|
||||
)
|
||||
{
|
||||
unsigned int base = write_head[0];
|
||||
extern "C" __global__ void rl_per_push_flush(
|
||||
/* ── Flush coordination ── */
|
||||
const int* __restrict__ flush_flags, /* [B] from ring kernel: 1=flush, 0=skip */
|
||||
int* __restrict__ write_offsets, /* [B+2] output: per-batch offsets, base, ready flag */
|
||||
|
||||
/* Exclusive prefix-sum */
|
||||
int total = 0;
|
||||
for (int i = 0; i < b_size; ++i) {
|
||||
write_offsets[i] = total;
|
||||
total += flush_flags[i];
|
||||
}
|
||||
|
||||
/* Advance write head and replay length */
|
||||
if (total > 0) {
|
||||
write_head[0] = (base + (unsigned int)total) % (unsigned int)capacity;
|
||||
unsigned int len = replay_len[0] + (unsigned int)total;
|
||||
if (len > (unsigned int)capacity) len = (unsigned int)capacity;
|
||||
replay_len[0] = len;
|
||||
}
|
||||
|
||||
/* Store base write head for the write kernel */
|
||||
write_offsets[b_size] = (int)base;
|
||||
}
|
||||
|
||||
/* =====================================================================
|
||||
* Kernel 2: coalesced replay write.
|
||||
*
|
||||
* Grid=(b_size,1,1), Block=(128,1,1). One block per batch element.
|
||||
*
|
||||
* Reads pre-computed offset from write_offsets (populated by kernel 1).
|
||||
* No spin — stream ordering guarantees kernel 1 completed before this.
|
||||
* ===================================================================== */
|
||||
extern "C" __global__ void rl_per_flush_write(
|
||||
const int* __restrict__ flush_flags, /* [B] */
|
||||
const int* __restrict__ write_offsets, /* [B+1] */
|
||||
|
||||
/* N-step ring (read source) */
|
||||
const float* __restrict__ nstep_scalars, /* [B * N_STEP_MAX * 3] */
|
||||
/* ── N-step ring (read source) ── */
|
||||
const float* __restrict__ nstep_scalars, /* [B * N_STEP_MAX * 3]: (r_scaled, r_raw, done) per ring slot */
|
||||
const float* __restrict__ nstep_h_t, /* [B * N_STEP_MAX * HIDDEN_DIM] */
|
||||
const unsigned int* __restrict__ nstep_write_idx, /* [B] ring write cursor */
|
||||
const unsigned int* __restrict__ nstep_count, /* [B] items in ring */
|
||||
|
||||
/* Current step data (for h_tp1) */
|
||||
/* ── Current step data (for h_tp1) ── */
|
||||
const float* __restrict__ h_tp1_current, /* [B * HIDDEN_DIM] */
|
||||
const float* __restrict__ log_pi_old, /* [B] */
|
||||
const int* __restrict__ actions, /* [B] */
|
||||
const float* __restrict__ isv, /* ISV array */
|
||||
|
||||
/* Replay storage (write destination) */
|
||||
/* ── Replay storage (write destination) ── */
|
||||
float* __restrict__ replay_h_t, /* [capacity * HIDDEN_DIM] */
|
||||
float* __restrict__ replay_h_tp1, /* [capacity * HIDDEN_DIM] */
|
||||
float* __restrict__ replay_scalars, /* [capacity * SCALARS_PER_TRANSITION] */
|
||||
float* __restrict__ priority_tree, /* [2 * capacity] (leaf layer at offset capacity) */
|
||||
float* __restrict__ max_priority, /* [1] running max priority */
|
||||
unsigned int* __restrict__ write_head, /* [1] circular write cursor */
|
||||
unsigned int* __restrict__ replay_len, /* [1] current buffer occupancy */
|
||||
float* __restrict__ max_priority, /* [1] running max priority for new inserts */
|
||||
|
||||
/* N-step count reset (written on flush) */
|
||||
unsigned int* __restrict__ nstep_count_out, /* [B] */
|
||||
/* ── N-step count reset (written on flush) ── */
|
||||
unsigned int* __restrict__ nstep_count_out, /* [B] — set to 0 for flushed batches */
|
||||
|
||||
int b_size,
|
||||
int capacity
|
||||
)
|
||||
{
|
||||
/* ====================================================================
|
||||
* PHASE 1: Block 0 coordination (thread 0 only).
|
||||
*
|
||||
* Computes prefix-sum of flush_flags to assign contiguous write slots,
|
||||
* advances write_head and replay_len, stores base for other blocks,
|
||||
* and signals ready via volatile write.
|
||||
* ==================================================================== */
|
||||
if (blockIdx.x == 0 && threadIdx.x == 0) {
|
||||
unsigned int base = write_head[0];
|
||||
|
||||
/* Exclusive prefix-sum: write_offsets[i] = number of flushes before batch i */
|
||||
int total = 0;
|
||||
for (int i = 0; i < b_size; ++i) {
|
||||
write_offsets[i] = total;
|
||||
total += flush_flags[i];
|
||||
}
|
||||
|
||||
/* Advance write head and replay length */
|
||||
if (total > 0) {
|
||||
write_head[0] = (base + (unsigned int)total) % (unsigned int)capacity;
|
||||
unsigned int len = replay_len[0] + (unsigned int)total;
|
||||
if (len > (unsigned int)capacity) len = (unsigned int)capacity;
|
||||
replay_len[0] = len;
|
||||
}
|
||||
|
||||
/* Store base write head for other blocks at slot [b_size] */
|
||||
write_offsets[b_size] = (int)base;
|
||||
|
||||
/* Fence: ensure all writes above are globally visible before signaling */
|
||||
__threadfence();
|
||||
|
||||
/* Signal ready at slot [b_size + 1] */
|
||||
write_offsets[b_size + 1] = 1;
|
||||
__threadfence();
|
||||
}
|
||||
|
||||
/* ====================================================================
|
||||
* PHASE 2: All blocks wait for ready signal.
|
||||
*
|
||||
* Non-block-0 blocks spin on a volatile read of the ready flag.
|
||||
* Block 0 already wrote it, so it proceeds immediately.
|
||||
* ==================================================================== */
|
||||
if (threadIdx.x == 0 && blockIdx.x != 0) {
|
||||
volatile int* ready = (volatile int*)&write_offsets[b_size + 1];
|
||||
while (*ready == 0) {
|
||||
/* spin — volatile load, no atomic */
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
/* ====================================================================
|
||||
* PHASE 3: Coalesced replay write (128 threads per block).
|
||||
*
|
||||
* Each flushing block writes:
|
||||
* - h_t[HIDDEN_DIM]: from oldest ring entry (coalesced, all 128 threads)
|
||||
* - h_tp1[HIDDEN_DIM]: from current step (coalesced, all 128 threads)
|
||||
* - scalars[7]: n-step return + metadata (thread 0 only)
|
||||
* - priority: max_priority into leaf layer (thread 0 only)
|
||||
* - reset: nstep_count_out[b] = 0 (thread 0 only)
|
||||
* ==================================================================== */
|
||||
const int b = blockIdx.x;
|
||||
|
||||
/* Early exit: nothing to flush for this batch */
|
||||
if (flush_flags[b] == 0) return;
|
||||
|
||||
/* Compute target replay slot from pre-computed offsets */
|
||||
/* Compute target replay slot */
|
||||
const unsigned int base = (unsigned int)write_offsets[b_size];
|
||||
const unsigned int my_offset = (unsigned int)write_offsets[b];
|
||||
const unsigned int my_slot = (base + my_offset) % (unsigned int)capacity;
|
||||
|
||||
@@ -1,11 +1,11 @@
|
||||
/* =====================================================================
|
||||
* rl_per_sample.cu — GPU-resident PER: proportional sampling via sum-tree
|
||||
*
|
||||
* Grid=(b_size), Block=(128). One block per sample.
|
||||
* Grid=(b_size), Block=(1). One thread per sample.
|
||||
*
|
||||
* Thread 0 does the tree walk + scalar unpack. All 128 threads cooperate
|
||||
* on the coalesced h_t and h_tp1 gathers (128 floats = HIDDEN_DIM per
|
||||
* sample, one float per thread, single 512-byte coalesced transaction).
|
||||
* Each thread samples a leaf from the priority sum-tree using stratified
|
||||
* sampling (segment per thread) with xorshift32 PRNG, then gathers the
|
||||
* transition data from replay storage.
|
||||
*
|
||||
* ISV reads: none (alpha used only at priority-update time)
|
||||
* ===================================================================== */
|
||||
@@ -44,87 +44,75 @@ extern "C" __global__ void rl_per_sample(
|
||||
int capacity
|
||||
)
|
||||
{
|
||||
const int b = blockIdx.x;
|
||||
const int tid = threadIdx.x;
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
/* ── Shared memory for broadcasting the sampled leaf index ────────── */
|
||||
__shared__ int s_safe_leaf;
|
||||
const unsigned int len = replay_len[0];
|
||||
const float total_priority = priority_tree[1]; /* root */
|
||||
|
||||
/* ── Thread 0: tree walk + scalar reads ──────────────────────────── */
|
||||
if (tid == 0) {
|
||||
const unsigned int len = replay_len[0];
|
||||
const float total_priority = priority_tree[1]; /* root */
|
||||
|
||||
/* Guard: empty or zero-priority replay */
|
||||
if (total_priority < 1e-9f || len == 0) {
|
||||
s_safe_leaf = -1; /* sentinel: zero-fill outputs */
|
||||
} else {
|
||||
/* Seed PRNG if cold */
|
||||
if (prng_state[b] == 0) {
|
||||
prng_state[b] = (unsigned int)(b + 1) * 2654435761u;
|
||||
}
|
||||
|
||||
/* Stratified sampling: draw u in [b*segment, (b+1)*segment) */
|
||||
const float segment = total_priority / (float)b_size;
|
||||
unsigned int rng = xorshift32(&prng_state[b]);
|
||||
const float u_frac = (float)(rng & 0x00FFFFFFu) / (float)0x01000000u;
|
||||
float u = ((float)b + u_frac) * segment;
|
||||
|
||||
/* Clamp to avoid floating-point overshoot */
|
||||
if (u >= total_priority) u = total_priority - 1e-6f;
|
||||
if (u < 0.0f) u = 0.0f;
|
||||
|
||||
/* Walk tree top-down */
|
||||
int idx = 1;
|
||||
while (idx < capacity) {
|
||||
const int left = 2 * idx;
|
||||
const float left_val = priority_tree[left];
|
||||
if (u <= left_val) {
|
||||
idx = left;
|
||||
} else {
|
||||
u -= left_val;
|
||||
idx = left + 1;
|
||||
}
|
||||
}
|
||||
const int leaf = idx - capacity;
|
||||
|
||||
/* Clamp leaf to valid range */
|
||||
const int safe_leaf = (leaf >= 0 && leaf < (int)len) ? leaf : 0;
|
||||
s_safe_leaf = safe_leaf;
|
||||
sample_indices[b] = (unsigned int)safe_leaf;
|
||||
|
||||
/* Unpack scalars (thread 0 only — small data, not worth parallelising) */
|
||||
const int sc_base = safe_leaf * SCALARS_PER_TRANSITION;
|
||||
sampled_actions[b] = (int)replay_scalars[sc_base + 0];
|
||||
sampled_rewards[b] = replay_scalars[sc_base + 1];
|
||||
sampled_dones[b] = replay_scalars[sc_base + 3];
|
||||
sampled_log_pi_old[b] = replay_scalars[sc_base + 4];
|
||||
sampled_n_step_gammas[b] = replay_scalars[sc_base + 5];
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
const int safe_leaf = s_safe_leaf;
|
||||
|
||||
if (safe_leaf < 0) {
|
||||
/* Empty replay — zero-fill h_t and h_tp1 (coalesced, all 128 threads) */
|
||||
sampled_h_t[b * HIDDEN_DIM + tid] = 0.0f;
|
||||
sampled_h_tp1[b * HIDDEN_DIM + tid] = 0.0f;
|
||||
if (tid == 0) {
|
||||
sampled_rewards[b] = 0.0f;
|
||||
sampled_dones[b] = 0.0f;
|
||||
sampled_log_pi_old[b] = 0.0f;
|
||||
sampled_n_step_gammas[b] = 0.0f;
|
||||
sampled_actions[b] = 0;
|
||||
sample_indices[b] = 0;
|
||||
/* Guard: empty or zero-priority replay */
|
||||
if (total_priority < 1e-9f || len == 0) {
|
||||
for (int i = 0; i < HIDDEN_DIM; ++i) {
|
||||
sampled_h_t[b * HIDDEN_DIM + i] = 0.0f;
|
||||
sampled_h_tp1[b * HIDDEN_DIM + i] = 0.0f;
|
||||
}
|
||||
sampled_rewards[b] = 0.0f;
|
||||
sampled_dones[b] = 0.0f;
|
||||
sampled_log_pi_old[b] = 0.0f;
|
||||
sampled_n_step_gammas[b] = 0.0f;
|
||||
sampled_actions[b] = 0;
|
||||
sample_indices[b] = 0;
|
||||
return;
|
||||
}
|
||||
|
||||
/* ── Coalesced h_t gather (128 threads, 1 float each) ────────────── */
|
||||
sampled_h_t[b * HIDDEN_DIM + tid] = replay_h_t[safe_leaf * HIDDEN_DIM + tid];
|
||||
/* ── Seed PRNG if cold ────────────────────────────────────────────── */
|
||||
if (prng_state[b] == 0) {
|
||||
prng_state[b] = (unsigned int)(b + 1) * 2654435761u;
|
||||
}
|
||||
|
||||
/* ── Coalesced h_tp1 gather (128 threads, 1 float each) ──────────── */
|
||||
sampled_h_tp1[b * HIDDEN_DIM + tid] = replay_h_tp1[safe_leaf * HIDDEN_DIM + tid];
|
||||
/* ── Stratified sampling: draw u in [b*segment, (b+1)*segment) ────── */
|
||||
const float segment = total_priority / (float)b_size;
|
||||
unsigned int rng = xorshift32(&prng_state[b]);
|
||||
const float u_frac = (float)(rng & 0x00FFFFFFu) / (float)0x01000000u; /* [0, 1) */
|
||||
float u = ((float)b + u_frac) * segment;
|
||||
|
||||
/* Clamp to avoid floating-point overshoot */
|
||||
if (u >= total_priority) u = total_priority - 1e-6f;
|
||||
if (u < 0.0f) u = 0.0f;
|
||||
|
||||
/* ── Walk tree top-down ───────────────────────────────────────────── */
|
||||
int idx = 1;
|
||||
while (idx < capacity) {
|
||||
const int left = 2 * idx;
|
||||
const float left_val = priority_tree[left];
|
||||
if (u <= left_val) {
|
||||
idx = left;
|
||||
} else {
|
||||
u -= left_val;
|
||||
idx = left + 1;
|
||||
}
|
||||
}
|
||||
const int leaf = idx - capacity;
|
||||
|
||||
/* Clamp leaf to valid range */
|
||||
const int safe_leaf = (leaf >= 0 && leaf < (int)len) ? leaf : 0;
|
||||
sample_indices[b] = (unsigned int)safe_leaf;
|
||||
|
||||
/* ── Gather h_t ──────────────────────────────────────────────────── */
|
||||
for (int i = 0; i < HIDDEN_DIM; ++i) {
|
||||
sampled_h_t[b * HIDDEN_DIM + i] = replay_h_t[safe_leaf * HIDDEN_DIM + i];
|
||||
}
|
||||
|
||||
/* ── Gather h_tp1 ────────────────────────────────────────────────── */
|
||||
for (int i = 0; i < HIDDEN_DIM; ++i) {
|
||||
sampled_h_tp1[b * HIDDEN_DIM + i] = replay_h_tp1[safe_leaf * HIDDEN_DIM + i];
|
||||
}
|
||||
|
||||
/* ── Unpack scalars ──────────────────────────────────────────────── */
|
||||
const int sc_base = safe_leaf * SCALARS_PER_TRANSITION;
|
||||
sampled_actions[b] = (int)replay_scalars[sc_base + 0];
|
||||
sampled_rewards[b] = replay_scalars[sc_base + 1]; /* n-step scaled return */
|
||||
sampled_dones[b] = replay_scalars[sc_base + 3];
|
||||
sampled_log_pi_old[b] = replay_scalars[sc_base + 4];
|
||||
sampled_n_step_gammas[b] = replay_scalars[sc_base + 5];
|
||||
}
|
||||
|
||||
@@ -1,30 +1,44 @@
|
||||
/* =====================================================================
|
||||
* rl_per_tree_rebuild.cu — GPU-resident PER: per-level sum-tree rebuild
|
||||
* rl_per_tree_rebuild.cu — GPU-resident PER: bottom-up parallel sum-tree rebuild
|
||||
*
|
||||
* One kernel invocation per tree level. The host launches log2(capacity)
|
||||
* times, from bottom (level 0 = parents of leaves) to top (root).
|
||||
* Stream ordering between launches provides the device-wide barrier that
|
||||
* the previous __threadfence() approach lacked — each level's writes are
|
||||
* fully complete before the next level reads them.
|
||||
* Grid=(128), Block=(256). Total 32768 threads = capacity.
|
||||
*
|
||||
* Graph-safe: each level is a separate kernel node in the CUDA graph.
|
||||
* The graph engine respects stream-order dependencies between nodes.
|
||||
* Rebuilds the entire sum-tree from leaf priorities (at indices
|
||||
* [capacity..2*capacity)) up to the root (index 1). Each level is
|
||||
* processed in parallel with __threadfence() device-wide barriers
|
||||
* between levels.
|
||||
*
|
||||
* No atomicAdd — each internal node is written by exactly one thread.
|
||||
* ===================================================================== */
|
||||
|
||||
extern "C" __global__ void rl_per_tree_rebuild_level(
|
||||
extern "C" __global__ void rl_per_tree_rebuild(
|
||||
float* __restrict__ priority_tree, /* [2 * capacity] */
|
||||
int start, /* first node index at this level */
|
||||
int nodes_at_level /* number of nodes to process at this level */
|
||||
int capacity
|
||||
)
|
||||
{
|
||||
const int tid_global = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
const int total_threads = gridDim.x * blockDim.x;
|
||||
|
||||
/* Grid-stride loop: each thread processes one or more nodes */
|
||||
for (int i = tid_global; i < nodes_at_level; i += total_threads) {
|
||||
const int node = start + i;
|
||||
priority_tree[node] = priority_tree[2 * node] + priority_tree[2 * node + 1];
|
||||
/* Bottom-up: level 0 = parents of leaves, level (log2(cap)-1) = root */
|
||||
/* At level L, there are capacity >> (L+1) internal nodes. */
|
||||
/* Node indices at level L: [capacity >> (L+1) .. capacity >> L) */
|
||||
|
||||
int nodes_at_level = capacity >> 1; /* level 0: cap/2 nodes */
|
||||
int start = nodes_at_level; /* first node index at this level */
|
||||
|
||||
while (nodes_at_level >= 1) {
|
||||
/* Each thread processes multiple nodes via grid-stride loop */
|
||||
for (int i = tid_global; i < nodes_at_level; i += total_threads) {
|
||||
const int node = start + i;
|
||||
priority_tree[node] = priority_tree[2 * node] + priority_tree[2 * node + 1];
|
||||
}
|
||||
|
||||
/* Device-wide fence: all writes at this level visible before
|
||||
* any thread reads them at the next level */
|
||||
__threadfence();
|
||||
|
||||
/* Move up one level */
|
||||
nodes_at_level >>= 1;
|
||||
start >>= 1;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -230,15 +230,6 @@ extern "C" __global__ void rl_reward_clamp_controller(
|
||||
// Suppress unused warnings on diagnostic-only state.
|
||||
(void) margin;
|
||||
|
||||
// Adaptive LOSS clamp from observed neg/pos ratio. WIN stays
|
||||
// structural at 1.0; LOSS adapts to actual loss magnitude to
|
||||
// give Q accurate resolution without over-allocating to rare tails.
|
||||
if (ema_new > 0.0f && neg_new > 0.0f) {
|
||||
const float observed_ratio = neg_new / ema_new;
|
||||
const float adaptive_loss = fmaxf(1.0f, fminf(observed_ratio * 1.1f, 3.0f));
|
||||
isv[RL_REWARD_CLAMP_LOSS_INDEX] = adaptive_loss;
|
||||
}
|
||||
|
||||
// ── Step 5: C51 atom span adaptation from observed reward EMAs. ──
|
||||
//
|
||||
// G.2 disabled this because atom-span growth + clamp-bound growth
|
||||
|
||||
37
crates/ml-alpha/cuda/rl_v_blend.cu
Normal file
37
crates/ml-alpha/cuda/rl_v_blend.cu
Normal file
@@ -0,0 +1,37 @@
|
||||
// rl_v_blend.cu — Phase 4.4 adaptive V baseline blend (2026-05-30).
|
||||
//
|
||||
// V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b]
|
||||
//
|
||||
// α read on-device from ISV[alpha_slot], emitted by
|
||||
// rl_v_blend_alpha_controller (separate kernel) which adapts α
|
||||
// based on observed |V_dq − V_scalar| / |V_scalar| tracking ratio.
|
||||
//
|
||||
// α = 1.0: pure Plan A v2 behavior (V_scalar drives PPO advantage)
|
||||
// α = 0.0: pure Phase 4.3 behavior (V_dq drives PPO advantage)
|
||||
// Anywhere in between: adaptive blend
|
||||
//
|
||||
// Per feedback_cpu_is_read_only: pure device kernel; α computed
|
||||
// device-side by the controller.
|
||||
// Per pearl_no_host_branches_in_captured_graph: graph-safe (reads
|
||||
// ISV pointer, no host params).
|
||||
// Per feedback_no_atomicadd: sole-writer per cell.
|
||||
//
|
||||
// Block layout: grid=(ceil(B/256), 1, 1), block=(256, 1, 1). Pure
|
||||
// elementwise op.
|
||||
|
||||
extern "C" __global__ void rl_v_blend(
|
||||
const float* __restrict__ v_scalar, // [B]
|
||||
const float* __restrict__ v_dq, // [B]
|
||||
const float* __restrict__ isv, // ISV bus
|
||||
int B,
|
||||
int alpha_slot,
|
||||
float* __restrict__ v_blended // [B]
|
||||
) {
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= B) return;
|
||||
// Defensive clamp to [0, 1] — controller should keep α bounded
|
||||
// (per pearl_audit_unboundedness_for_implicit_asymmetry) but a
|
||||
// kernel-side guard protects against any controller bug.
|
||||
const float a = fminf(1.0f, fmaxf(0.0f, isv[alpha_slot]));
|
||||
v_blended[b] = a * v_scalar[b] + (1.0f - a) * v_dq[b];
|
||||
}
|
||||
121
crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu
Normal file
121
crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu
Normal file
@@ -0,0 +1,121 @@
|
||||
// rl_v_blend_alpha_controller.cu — Phase 4.4 ISV-adaptive V blend (2026-05-30).
|
||||
//
|
||||
// Drives α ∈ [0, 1] for `V_used = α × V_scalar + (1−α) × V_dq` based
|
||||
// on observed V_dq vs V_scalar tracking ratio:
|
||||
//
|
||||
// track_ratio = EMA(|V_dq − V_scalar|) / EMA(|V_scalar|)
|
||||
//
|
||||
// if track_ratio > 1.5 × TARGET: α ← min(α + step, 1.0) ↑ V_scalar
|
||||
// if track_ratio < TARGET / 1.5: α ← max(α - step, 0.0) ↑ V_dq
|
||||
// else: hold α
|
||||
//
|
||||
// Per pearl_wiener_alpha_floor_for_nonstationary: Schulman bounded
|
||||
// discrete step, no Wiener-α blending of the controller variable
|
||||
// itself (α is the controlled quantity).
|
||||
//
|
||||
// Per pearl_first_observation_bootstrap: bootstrap α = 1.0 on
|
||||
// sentinel input (ISV[alpha_slot] == 0), EMAs use first observation
|
||||
// directly. After bootstrap, α never naturally returns to exactly 0
|
||||
// because Schulman step (0.01) is unlikely to land on it; if it
|
||||
// does, controller re-bootstraps harmlessly.
|
||||
//
|
||||
// Per pearl_blend_formulas_must_have_permanent_floor: dead-signal
|
||||
// guard — if EMA(|V_scalar|) < FLOOR, hold α (no V signal to
|
||||
// calibrate against; the trainer hasn't seen meaningful rewards yet).
|
||||
//
|
||||
// Per feedback_cpu_is_read_only: pure device kernel; reads V_scalar,
|
||||
// V_dq, ISV; emits α + EMAs to ISV. No host control.
|
||||
//
|
||||
// Block layout: grid=(1, 1, 1), block=(BLOCK_X=1024, 1, 1). Single
|
||||
// block does parallel reduction over batch up to B=1024. Thread 0
|
||||
// performs the controller update.
|
||||
|
||||
#define BLOCK_X 1024
|
||||
#define EMA_ALPHA 0.01f
|
||||
#define TARGET_TRACK_RATIO 0.10f
|
||||
#define SCHULMAN_STEP 0.01f
|
||||
#define DEAD_SIGNAL_FLOOR 1e-4f
|
||||
#define BOOTSTRAP_ALPHA 1.0f
|
||||
|
||||
extern "C" __global__ void rl_v_blend_alpha_controller(
|
||||
const float* __restrict__ v_scalar, // [B]
|
||||
const float* __restrict__ v_dq, // [B]
|
||||
float* __restrict__ isv, // ISV bus
|
||||
int B,
|
||||
int alpha_slot, // ISV[α]
|
||||
int trackerr_ema_slot, // ISV[|V_dq − V_scalar|_ema]
|
||||
int v_scalar_mag_ema_slot // ISV[|V_scalar|_ema] (dead-signal floor)
|
||||
) {
|
||||
const int tid = threadIdx.x;
|
||||
if (tid >= BLOCK_X) return;
|
||||
|
||||
// ── Per-thread partial sums over strided batch ──
|
||||
float track_partial = 0.0f;
|
||||
float mag_partial = 0.0f;
|
||||
for (int b = tid; b < B; b += BLOCK_X) {
|
||||
const float vs = v_scalar[b];
|
||||
const float vd = v_dq[b];
|
||||
track_partial += fabsf(vd - vs);
|
||||
mag_partial += fabsf(vs);
|
||||
}
|
||||
|
||||
// ── Tree-reduce in shared mem ──
|
||||
__shared__ float s_t[BLOCK_X];
|
||||
__shared__ float s_m[BLOCK_X];
|
||||
s_t[tid] = track_partial;
|
||||
s_m[tid] = mag_partial;
|
||||
__syncthreads();
|
||||
|
||||
for (int stride = BLOCK_X / 2; stride > 0; stride >>= 1) {
|
||||
if (tid < stride) {
|
||||
s_t[tid] += s_t[tid + stride];
|
||||
s_m[tid] += s_m[tid + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
if (tid == 0) {
|
||||
const float inv_B = 1.0f / (float)B;
|
||||
const float track_mean = s_t[0] * inv_B;
|
||||
const float mag_mean = s_m[0] * inv_B;
|
||||
|
||||
// ── Bootstrap α on sentinel ──
|
||||
float alpha = isv[alpha_slot];
|
||||
if (alpha == 0.0f) {
|
||||
alpha = BOOTSTRAP_ALPHA;
|
||||
}
|
||||
|
||||
// ── First-observation bootstrap on EMAs ──
|
||||
float prev_track_ema = isv[trackerr_ema_slot];
|
||||
float prev_mag_ema = isv[v_scalar_mag_ema_slot];
|
||||
const float track_ema = (prev_track_ema == 0.0f)
|
||||
? track_mean
|
||||
: (1.0f - EMA_ALPHA) * prev_track_ema + EMA_ALPHA * track_mean;
|
||||
const float mag_ema = (prev_mag_ema == 0.0f)
|
||||
? mag_mean
|
||||
: (1.0f - EMA_ALPHA) * prev_mag_ema + EMA_ALPHA * mag_mean;
|
||||
|
||||
// ── Dead-signal guard: V_scalar magnitude too small to calibrate against ──
|
||||
if (mag_ema < DEAD_SIGNAL_FLOOR) {
|
||||
isv[trackerr_ema_slot] = track_ema;
|
||||
isv[v_scalar_mag_ema_slot] = mag_ema;
|
||||
isv[alpha_slot] = alpha; // hold (write bootstrap if needed)
|
||||
return;
|
||||
}
|
||||
|
||||
// ── Track ratio + Schulman-bounded step on α ──
|
||||
const float track_ratio = track_ema / mag_ema;
|
||||
if (track_ratio > 1.5f * TARGET_TRACK_RATIO) {
|
||||
// V_dq diverged from V_scalar → raise α toward V_scalar
|
||||
alpha = fminf(alpha + SCHULMAN_STEP, 1.0f);
|
||||
} else if (track_ratio < TARGET_TRACK_RATIO / 1.5f) {
|
||||
// V_dq tracks V_scalar well → lower α toward V_dq
|
||||
alpha = fmaxf(alpha - SCHULMAN_STEP, 0.0f);
|
||||
}
|
||||
// else: hold α (within band)
|
||||
|
||||
isv[alpha_slot] = alpha;
|
||||
isv[trackerr_ema_slot] = track_ema;
|
||||
isv[v_scalar_mag_ema_slot] = mag_ema;
|
||||
}
|
||||
}
|
||||
@@ -26,6 +26,24 @@
|
||||
#define VSN_FEATURE_DIM 40
|
||||
#define VSN_BLOCK 64 // round up to warp-multiple; threads i >= FEATURE_DIM idle.
|
||||
|
||||
// 2026-05-29 stride-mismatch fix.
|
||||
// VSN's input buffer (window_tensor_d) is allocated [B, K, ENCODER_INPUT_DIM=56]
|
||||
// by perception.rs (snap features [0..40) + per-batch broadcast context
|
||||
// [40..56) written by rl_encoder_context_broadcast). VSN only processes the
|
||||
// first VSN_FEATURE_DIM=40 features per row (snap features), but the input
|
||||
// rows are spaced 56 floats apart, not 40. The kernel originally indexed x
|
||||
// with stride VSN_FEATURE_DIM=40 — correct ONLY for row 0; every subsequent
|
||||
// row read mixed broadcast-context + snap features across the [B, K, 56]
|
||||
// row boundaries. Symptoms: intermittent step-4 NaN as accumulating trade
|
||||
// context magnitudes overflowed VSN's softmax via the bleed.
|
||||
//
|
||||
// Fix: use VSN_X_ROW_STRIDE=56 for reading x in both forward and backward.
|
||||
// Output (gates, y) and gradient outputs (grad_W, grad_b, grad_x) remain at
|
||||
// VSN_FEATURE_DIM=40 because the downstream consumers (Mamba2 L1 with
|
||||
// in_dim=40) read at compact 40-stride. grad_x is unused downstream (see
|
||||
// `vsn_grad_x_d` audit — write-only), so its stride doesn't matter.
|
||||
#define VSN_X_ROW_STRIDE 56 // = ENCODER_INPUT_DIM in heads.rs / perception.rs
|
||||
|
||||
extern "C" __global__ void variable_selection_fwd(
|
||||
const float* __restrict__ W_vsn, // [FEATURE_DIM, FEATURE_DIM]
|
||||
const float* __restrict__ b_vsn, // [FEATURE_DIM]
|
||||
@@ -38,7 +56,7 @@ extern "C" __global__ void variable_selection_fwd(
|
||||
int tid = threadIdx.x;
|
||||
if (row >= n_rows) return;
|
||||
|
||||
const float* x_row = x + (long long)row * VSN_FEATURE_DIM;
|
||||
const float* x_row = x + (long long)row * VSN_X_ROW_STRIDE;
|
||||
|
||||
// Shared mem: gate_logit + max-reduce scratch + sum-reduce scratch.
|
||||
__shared__ float s_logit[VSN_FEATURE_DIM];
|
||||
@@ -170,7 +188,7 @@ extern "C" __global__ void variable_selection_bwd(
|
||||
float dy_i = 0.0f;
|
||||
if (tid < VSN_FEATURE_DIM) {
|
||||
gates_i = gates[(long long)row * VSN_FEATURE_DIM + tid];
|
||||
x_i = x[(long long)row * VSN_FEATURE_DIM + tid];
|
||||
x_i = x[(long long)row * VSN_X_ROW_STRIDE + tid];
|
||||
dy_i = grad_y[(long long)row * VSN_FEATURE_DIM + tid];
|
||||
s_gates[tid] = gates_i;
|
||||
s_dgates[tid] = dy_i * x_i;
|
||||
@@ -201,7 +219,7 @@ extern "C" __global__ void variable_selection_bwd(
|
||||
const float dl_t = s_dlogit[tid];
|
||||
#pragma unroll
|
||||
for (int j = 0; j < VSN_FEATURE_DIM; ++j) {
|
||||
const float xj = x[(long long)row * VSN_FEATURE_DIM + j];
|
||||
const float xj = x[(long long)row * VSN_X_ROW_STRIDE + j];
|
||||
grad_W_vsn_scratch[row_FF + (long long)tid * VSN_FEATURE_DIM + j]
|
||||
+= dl_t * xj;
|
||||
}
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -188,7 +188,6 @@ impl GpuDataLoader {
|
||||
soa: &SoaBufferPtrs,
|
||||
seq_len: usize,
|
||||
batch_size: usize,
|
||||
label_outs: [u64; 5],
|
||||
) -> Result<()> {
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(dataset.snapshots_d.raw_ptr());
|
||||
@@ -211,18 +210,6 @@ impl GpuDataLoader {
|
||||
args.push_ptr(self.sample_file_offset_d.raw_ptr());
|
||||
args.push_ptr(self.sample_anchor_d.raw_ptr());
|
||||
args.push_ptr(self.sample_file_idx_d.raw_ptr());
|
||||
// Fused label gather arguments
|
||||
args.push_i32(dataset.total_snapshots as i32);
|
||||
args.push_ptr(dataset.labels_d.raw_ptr());
|
||||
args.push_ptr(dataset.outcome_prof_long_d.raw_ptr());
|
||||
args.push_ptr(dataset.outcome_prof_short_d.raw_ptr());
|
||||
args.push_ptr(dataset.outcome_size_long_d.raw_ptr());
|
||||
args.push_ptr(dataset.outcome_size_short_d.raw_ptr());
|
||||
args.push_ptr(label_outs[0]);
|
||||
args.push_ptr(label_outs[1]);
|
||||
args.push_ptr(label_outs[2]);
|
||||
args.push_ptr(label_outs[3]);
|
||||
args.push_ptr(label_outs[4]);
|
||||
args.push_i32(batch_size as i32);
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
|
||||
@@ -547,12 +547,6 @@ impl Mamba2Block {
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow!("Mamba2Block: cublasSetWorkspace_v2: {e:?}"))?;
|
||||
cudarc::cublas::sys::cublasSetMathMode(
|
||||
*cublas.handle(),
|
||||
cudarc::cublas::sys::cublasMath_t::CUBLAS_TF32_TENSOR_OP_MATH,
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow!("Mamba2Block: cublasSetMathMode TF32: {e:?}"))?;
|
||||
}
|
||||
|
||||
// ── Parameter allocation + initialisation. ──────────────────────
|
||||
|
||||
@@ -311,12 +311,6 @@ impl DqnHead {
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow::anyhow!("DqnHead: cublasSetWorkspace_v2: {e:?}"))?;
|
||||
cudarc::cublas::sys::cublasSetMathMode(
|
||||
*cublas.handle(),
|
||||
cudarc::cublas::sys::cublasMath_t::CUBLAS_TF32_TENSOR_OP_MATH,
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow::anyhow!("DqnHead: cublasSetMathMode TF32: {e:?}"))?;
|
||||
}
|
||||
|
||||
// Bias-add and reduce-sum-axis0 kernel handles from ml-core.
|
||||
|
||||
530
crates/ml-alpha/src/rl/dueling_q.rs
Normal file
530
crates/ml-alpha/src/rl/dueling_q.rs
Normal file
@@ -0,0 +1,530 @@
|
||||
//! Phase 4 — Independent Dueling Q head (2026-05-30).
|
||||
//!
|
||||
//! Parallel head to C51/IQN/π/value_head with ZERO shared state with
|
||||
//! downstream consumers (ensemble, distill, action selection). Trains
|
||||
//! its own V + A weights via Bellman loss on composed_Q at taken action.
|
||||
//!
|
||||
//! V output feeds PPO advantage baseline (Phase 4.3) — composed_Q is
|
||||
//! internal to this head's loss path and never consumed elsewhere.
|
||||
//!
|
||||
//! Per spec docs/superpowers/specs/2026-05-30-phase4-independent-dueling-head-design.md.
|
||||
//!
|
||||
//! ## Why a separate head (vs Phase 3 modifications to IQN)
|
||||
//!
|
||||
//! All 5 prior dueling attempts (Phase 2 v2, Phase 3.1, 3.2, 3.2c,
|
||||
//! 3.1-fix) failed because they modified existing architectures'
|
||||
//! weights or outputs in ways that perturbed downstream consumers. See
|
||||
//! the four session pearls captured in the spec's §10.
|
||||
//!
|
||||
//! Phase 4's design is **structurally encapsulated**: DuelingQHead has
|
||||
//! ITS OWN weights (w_v, b_v, w_a, b_a) and its OWN Bellman loss.
|
||||
//! Mean-zero A gradient pattern only affects DuelingQHead's training,
|
||||
//! never C51's or IQN's weights. composed_Q never feeds ensemble or
|
||||
//! distill or action selection.
|
||||
//!
|
||||
//! ## Forward
|
||||
//!
|
||||
//! ```text
|
||||
//! V[b] = Σ_c w_v[c] × h_t[b, c] + b_v[0]
|
||||
//! A[b, a] = Σ_c w_a[c, a] × h_t[b, c] + b_a[a]
|
||||
//! composed_Q[b, a] = V[b] + A[b, a] − (1/N) Σ_a' A[b, a']
|
||||
//! ```
|
||||
//!
|
||||
//! Single fused kernel (`rl_dueling_q_forward`) computes all three.
|
||||
//!
|
||||
//! ## Constraints honoured
|
||||
//!
|
||||
//! * `feedback_no_atomicadd` — sole-writer per output cell.
|
||||
//! * `feedback_cpu_is_read_only` — pure device kernels.
|
||||
//! * `feedback_no_htod_htoh_only_mapped_pinned` — weight uploads stage
|
||||
//! through `MappedF32Buffer`.
|
||||
//! * `feedback_no_nvrtc` — pre-compiled cubins via `build.rs`.
|
||||
//! * `pearl_scoped_init_seed_for_reproducibility` — `new()` installs
|
||||
//! `scoped_init_seed` before drawing Xavier samples.
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use anyhow::{Context, Result};
|
||||
use cudarc::driver::{CudaFunction, CudaModule, CudaSlice, CudaStream};
|
||||
use cudarc::driver::sys::CUstream;
|
||||
use ml_core::cuda_autograd::init::scoped_init_seed;
|
||||
use ml_core::device::MlDevice;
|
||||
use rand::{Rng, SeedableRng};
|
||||
use rand_chacha::ChaCha8Rng;
|
||||
|
||||
use crate::heads::HIDDEN_DIM;
|
||||
use crate::pinned_mem::MappedF32Buffer;
|
||||
use crate::rl::common::N_ACTIONS;
|
||||
use crate::trainer::raw_launch::{RawArgs, raw_launch};
|
||||
|
||||
const DUELING_Q_FORWARD_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"),
|
||||
"/rl_dueling_q_forward.cubin"
|
||||
));
|
||||
const DUELING_Q_BELLMAN_TARGET_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"),
|
||||
"/rl_dueling_q_bellman_target.cubin"
|
||||
));
|
||||
const DUELING_Q_LOSS_AND_GRAD_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"),
|
||||
"/rl_dueling_q_loss_and_grad.cubin"
|
||||
));
|
||||
const DUELING_Q_DECOMPOSE_AND_BWD_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"),
|
||||
"/rl_dueling_q_decompose_and_bwd.cubin"
|
||||
));
|
||||
/// Reuse the existing dqn_target_soft_update kernel — it's a generic
|
||||
/// element-wise `target = (1−τ) target + τ online` blend that works
|
||||
/// for any tensor size. τ read from ISV[RL_TARGET_TAU_INDEX=401].
|
||||
const DQN_TARGET_SOFT_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(
|
||||
env!("OUT_DIR"),
|
||||
"/dqn_target_soft_update.cubin"
|
||||
));
|
||||
|
||||
/// Construction config for [`DuelingQHead`].
|
||||
#[derive(Clone, Debug)]
|
||||
pub struct DuelingQHeadConfig {
|
||||
/// Encoder hidden dimension `h_t [B, HIDDEN_DIM]` feeding the head.
|
||||
/// Must equal the kernel-side `HIDDEN_DIM` define (128).
|
||||
pub hidden_dim: usize,
|
||||
/// Random seed for Xavier init. Distinct from C51 / IQN / V_scalar
|
||||
/// seeds so initial draws are independent.
|
||||
pub seed: u64,
|
||||
}
|
||||
|
||||
impl Default for DuelingQHeadConfig {
|
||||
fn default() -> Self {
|
||||
Self {
|
||||
hidden_dim: HIDDEN_DIM,
|
||||
seed: 0x4DEAD_1234,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// Independent dueling Q head — V + A decomposition with scalar
|
||||
/// Bellman loss. Trains in parallel to C51 / IQN; supplies V_dq for
|
||||
/// PPO advantage baseline use (Phase 4.3).
|
||||
pub struct DuelingQHead {
|
||||
#[allow(dead_code)]
|
||||
cfg: DuelingQHeadConfig,
|
||||
stream: Arc<CudaStream>,
|
||||
raw_stream: CUstream,
|
||||
|
||||
// ── Kernel handles ───────────────────────────────────────────────
|
||||
_fwd_module: Arc<CudaModule>,
|
||||
pub forward_fn: CudaFunction,
|
||||
_bellman_module: Arc<CudaModule>,
|
||||
pub bellman_target_fn: CudaFunction,
|
||||
_loss_module: Arc<CudaModule>,
|
||||
pub loss_and_grad_fn: CudaFunction,
|
||||
_bwd_module: Arc<CudaModule>,
|
||||
pub decompose_and_bwd_fn: CudaFunction,
|
||||
_soft_update_module: Arc<CudaModule>,
|
||||
pub soft_update_fn: CudaFunction,
|
||||
|
||||
// ── Online weights ───────────────────────────────────────────────
|
||||
/// V projection weights `[HIDDEN_DIM]`.
|
||||
pub w_v_d: CudaSlice<f32>,
|
||||
/// V projection bias `[1]`.
|
||||
pub b_v_d: CudaSlice<f32>,
|
||||
/// A projection weights `[HIDDEN_DIM × N_ACTIONS]`, row-major
|
||||
/// (kernel reads `w_a[c * N_ACTIONS + a]`).
|
||||
pub w_a_d: CudaSlice<f32>,
|
||||
/// A projection bias `[N_ACTIONS]`.
|
||||
pub b_a_d: CudaSlice<f32>,
|
||||
|
||||
// ── Target-network weights (soft-updated each step) ──────────────
|
||||
pub w_v_target_d: CudaSlice<f32>,
|
||||
pub b_v_target_d: CudaSlice<f32>,
|
||||
pub w_a_target_d: CudaSlice<f32>,
|
||||
pub b_a_target_d: CudaSlice<f32>,
|
||||
}
|
||||
|
||||
impl DuelingQHead {
|
||||
/// Allocate device weights, load cubins, cache kernel handles.
|
||||
pub fn new(dev: &MlDevice, cfg: DuelingQHeadConfig) -> Result<Self> {
|
||||
let stream: Arc<CudaStream> = dev
|
||||
.cuda_stream()
|
||||
.context("dueling_q_head stream")?
|
||||
.clone();
|
||||
let ctx = dev.cuda_context().context("dueling_q_head ctx")?;
|
||||
|
||||
// ── Load forward cubin ───────────────────────────────────────
|
||||
let fwd_module = ctx
|
||||
.load_cubin(DUELING_Q_FORWARD_CUBIN.to_vec())
|
||||
.context("load rl_dueling_q_forward cubin")?;
|
||||
let forward_fn = fwd_module
|
||||
.load_function("rl_dueling_q_forward")
|
||||
.context("load rl_dueling_q_forward")?;
|
||||
let bellman_module = ctx
|
||||
.load_cubin(DUELING_Q_BELLMAN_TARGET_CUBIN.to_vec())
|
||||
.context("load rl_dueling_q_bellman_target cubin")?;
|
||||
let bellman_target_fn = bellman_module
|
||||
.load_function("rl_dueling_q_bellman_target_build")
|
||||
.context("load rl_dueling_q_bellman_target_build")?;
|
||||
let loss_module = ctx
|
||||
.load_cubin(DUELING_Q_LOSS_AND_GRAD_CUBIN.to_vec())
|
||||
.context("load rl_dueling_q_loss_and_grad cubin")?;
|
||||
let loss_and_grad_fn = loss_module
|
||||
.load_function("rl_dueling_q_loss_and_grad")
|
||||
.context("load rl_dueling_q_loss_and_grad")?;
|
||||
let bwd_module = ctx
|
||||
.load_cubin(DUELING_Q_DECOMPOSE_AND_BWD_CUBIN.to_vec())
|
||||
.context("load rl_dueling_q_decompose_and_bwd cubin")?;
|
||||
let decompose_and_bwd_fn = bwd_module
|
||||
.load_function("rl_dueling_q_decompose_and_weight_grad")
|
||||
.context("load rl_dueling_q_decompose_and_weight_grad")?;
|
||||
// Reuse dqn_target_soft_update kernel — generic element-wise blend.
|
||||
let soft_update_module = ctx
|
||||
.load_cubin(DQN_TARGET_SOFT_UPDATE_CUBIN.to_vec())
|
||||
.context("load dqn_target_soft_update cubin for dueling")?;
|
||||
let soft_update_fn = soft_update_module
|
||||
.load_function("dqn_target_soft_update")
|
||||
.context("load dqn_target_soft_update for dueling")?;
|
||||
|
||||
// ── Weight init (Xavier uniform, scaled by 0.01 like sibling heads) ──
|
||||
// Per pearl_scoped_init_seed_for_reproducibility.
|
||||
let _seed_guard = scoped_init_seed(cfg.seed);
|
||||
let mut rng = ChaCha8Rng::seed_from_u64(cfg.seed);
|
||||
|
||||
// V projection: HIDDEN_DIM → 1
|
||||
let v_scale = 0.01_f32 * (6.0_f32 / (cfg.hidden_dim + 1) as f32).sqrt();
|
||||
let w_v_host: Vec<f32> = (0..cfg.hidden_dim)
|
||||
.map(|_| rng.gen_range(-v_scale..v_scale))
|
||||
.collect();
|
||||
let b_v_host: Vec<f32> = vec![0.0_f32; 1];
|
||||
|
||||
// A projection: HIDDEN_DIM → N_ACTIONS
|
||||
let a_scale =
|
||||
0.01_f32 * (6.0_f32 / (cfg.hidden_dim + N_ACTIONS) as f32).sqrt();
|
||||
let w_a_host: Vec<f32> = (0..cfg.hidden_dim * N_ACTIONS)
|
||||
.map(|_| rng.gen_range(-a_scale..a_scale))
|
||||
.collect();
|
||||
let b_a_host: Vec<f32> = vec![0.0_f32; N_ACTIONS];
|
||||
|
||||
// Upload online weights.
|
||||
let w_v_d = upload(&stream, &w_v_host)?;
|
||||
let b_v_d = upload(&stream, &b_v_host)?;
|
||||
let w_a_d = upload(&stream, &w_a_host)?;
|
||||
let b_a_d = upload(&stream, &b_a_host)?;
|
||||
|
||||
// Upload target weights (identical init — soft-updated by trainer).
|
||||
let w_v_target_d = upload(&stream, &w_v_host)?;
|
||||
let b_v_target_d = upload(&stream, &b_v_host)?;
|
||||
let w_a_target_d = upload(&stream, &w_a_host)?;
|
||||
let b_a_target_d = upload(&stream, &b_a_host)?;
|
||||
|
||||
let raw_stream = stream.cu_stream();
|
||||
|
||||
Ok(Self {
|
||||
cfg,
|
||||
stream,
|
||||
raw_stream,
|
||||
_fwd_module: fwd_module,
|
||||
forward_fn,
|
||||
_bellman_module: bellman_module,
|
||||
bellman_target_fn,
|
||||
_loss_module: loss_module,
|
||||
loss_and_grad_fn,
|
||||
_bwd_module: bwd_module,
|
||||
decompose_and_bwd_fn,
|
||||
_soft_update_module: soft_update_module,
|
||||
soft_update_fn,
|
||||
w_v_d,
|
||||
b_v_d,
|
||||
w_a_d,
|
||||
b_a_d,
|
||||
w_v_target_d,
|
||||
b_v_target_d,
|
||||
w_a_target_d,
|
||||
b_a_target_d,
|
||||
})
|
||||
}
|
||||
|
||||
/// Stream used to launch all kernels owned by this head.
|
||||
pub fn stream(&self) -> &Arc<CudaStream> {
|
||||
&self.stream
|
||||
}
|
||||
|
||||
/// cuBLAS-free forward pass with online weights.
|
||||
///
|
||||
/// Computes V[B], A[B × N_ACTIONS], and composed_Q[B × N_ACTIONS]
|
||||
/// in a single fused kernel.
|
||||
///
|
||||
/// Block layout: grid=(B, 1, 1), block=(HIDDEN_DIM, 1, 1),
|
||||
/// shared_mem = HIDDEN_DIM × 4 bytes (for s_h cache).
|
||||
pub fn forward(
|
||||
&self,
|
||||
h_t: &CudaSlice<f32>,
|
||||
b_size: usize,
|
||||
v_out: &mut CudaSlice<f32>,
|
||||
a_out: &mut CudaSlice<f32>,
|
||||
q_composed_out: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
self.forward_inner(
|
||||
h_t,
|
||||
&self.w_v_d, &self.b_v_d,
|
||||
&self.w_a_d, &self.b_a_d,
|
||||
b_size, v_out, a_out, q_composed_out,
|
||||
)
|
||||
}
|
||||
|
||||
/// cuBLAS-free forward pass with target-network weights.
|
||||
pub fn forward_target(
|
||||
&self,
|
||||
h_t: &CudaSlice<f32>,
|
||||
b_size: usize,
|
||||
v_target_out: &mut CudaSlice<f32>,
|
||||
a_target_out: &mut CudaSlice<f32>,
|
||||
q_target_composed_out: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
self.forward_inner(
|
||||
h_t,
|
||||
&self.w_v_target_d, &self.b_v_target_d,
|
||||
&self.w_a_target_d, &self.b_a_target_d,
|
||||
b_size, v_target_out, a_target_out, q_target_composed_out,
|
||||
)
|
||||
}
|
||||
|
||||
/// Internal forward — parameterised over weight slices so online /
|
||||
/// target paths share code.
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
fn forward_inner(
|
||||
&self,
|
||||
h_t: &CudaSlice<f32>,
|
||||
w_v: &CudaSlice<f32>,
|
||||
b_v: &CudaSlice<f32>,
|
||||
w_a: &CudaSlice<f32>,
|
||||
b_a: &CudaSlice<f32>,
|
||||
b_size: usize,
|
||||
v_out: &mut CudaSlice<f32>,
|
||||
a_out: &mut CudaSlice<f32>,
|
||||
q_composed_out: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
let hd = self.cfg.hidden_dim;
|
||||
debug_assert_eq!(h_t.len(), b_size * hd);
|
||||
debug_assert_eq!(w_v.len(), hd);
|
||||
debug_assert_eq!(b_v.len(), 1);
|
||||
debug_assert_eq!(w_a.len(), hd * N_ACTIONS);
|
||||
debug_assert_eq!(b_a.len(), N_ACTIONS);
|
||||
debug_assert_eq!(v_out.len(), b_size);
|
||||
debug_assert_eq!(a_out.len(), b_size * N_ACTIONS);
|
||||
debug_assert_eq!(q_composed_out.len(), b_size * N_ACTIONS);
|
||||
|
||||
let b_i = b_size as i32;
|
||||
let smem = (hd * std::mem::size_of::<f32>()) as u32;
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(h_t.raw_ptr());
|
||||
args.push_ptr(w_v.raw_ptr());
|
||||
args.push_ptr(b_v.raw_ptr());
|
||||
args.push_ptr(w_a.raw_ptr());
|
||||
args.push_ptr(b_a.raw_ptr());
|
||||
args.push_i32(b_i);
|
||||
args.push_ptr(v_out.raw_ptr());
|
||||
args.push_ptr(a_out.raw_ptr());
|
||||
args.push_ptr(q_composed_out.raw_ptr());
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.forward_fn.cu_function(),
|
||||
(b_size as u32, 1, 1),
|
||||
(hd as u32, 1, 1),
|
||||
smem,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("rl_dueling_q_forward: {:?}", e))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Build the scalar Bellman target from the target network's
|
||||
/// composed_Q at s_{t+1}:
|
||||
/// a*[b] = argmax_a' target_composed_Q[b, a']
|
||||
/// target_value[b] = r[b] + γ^n × (1 − done[b]) × target_composed_Q[b, a*[b]]
|
||||
///
|
||||
/// Per-sample γ^n is passed (matches PER n-step convention used by
|
||||
/// C51 / IQN target builds).
|
||||
pub fn build_bellman_target(
|
||||
&self,
|
||||
target_composed_q: &CudaSlice<f32>,
|
||||
rewards: &CudaSlice<f32>,
|
||||
dones: &CudaSlice<f32>,
|
||||
n_step_gammas: &CudaSlice<f32>,
|
||||
b_size: usize,
|
||||
target_value_out: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
debug_assert_eq!(target_composed_q.len(), b_size * N_ACTIONS);
|
||||
debug_assert_eq!(rewards.len(), b_size);
|
||||
debug_assert_eq!(dones.len(), b_size);
|
||||
debug_assert_eq!(n_step_gammas.len(), b_size);
|
||||
debug_assert_eq!(target_value_out.len(), b_size);
|
||||
|
||||
let b_i = b_size as i32;
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(target_composed_q.raw_ptr());
|
||||
args.push_ptr(rewards.raw_ptr());
|
||||
args.push_ptr(dones.raw_ptr());
|
||||
args.push_ptr(n_step_gammas.raw_ptr());
|
||||
args.push_i32(b_i);
|
||||
args.push_ptr(target_value_out.raw_ptr());
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.bellman_target_fn.cu_function(),
|
||||
(b_size as u32, 1, 1),
|
||||
(N_ACTIONS as u32, 1, 1),
|
||||
0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("rl_dueling_q_bellman_target_build: {:?}", e))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Scalar Huber loss on (target − online_composed_Q[taken]).
|
||||
/// Emits per-batch loss and grad_composed (only taken action has
|
||||
/// nonzero gradient — single-Q regression).
|
||||
pub fn compute_loss_and_grad(
|
||||
&self,
|
||||
online_composed_q: &CudaSlice<f32>,
|
||||
target_value: &CudaSlice<f32>,
|
||||
actions_taken: &CudaSlice<i32>,
|
||||
b_size: usize,
|
||||
loss_per_batch: &mut CudaSlice<f32>,
|
||||
grad_composed_out: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
debug_assert_eq!(online_composed_q.len(), b_size * N_ACTIONS);
|
||||
debug_assert_eq!(target_value.len(), b_size);
|
||||
debug_assert_eq!(actions_taken.len(), b_size);
|
||||
debug_assert_eq!(loss_per_batch.len(), b_size);
|
||||
debug_assert_eq!(grad_composed_out.len(), b_size * N_ACTIONS);
|
||||
|
||||
let b_i = b_size as i32;
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(online_composed_q.raw_ptr());
|
||||
args.push_ptr(target_value.raw_ptr());
|
||||
args.push_ptr(actions_taken.raw_ptr());
|
||||
args.push_i32(b_i);
|
||||
args.push_ptr(loss_per_batch.raw_ptr());
|
||||
args.push_ptr(grad_composed_out.raw_ptr());
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.loss_and_grad_fn.cu_function(),
|
||||
(b_size as u32, 1, 1),
|
||||
(N_ACTIONS as u32, 1, 1),
|
||||
0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("rl_dueling_q_loss_and_grad: {:?}", e))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Decompose grad_composed → grad_V + grad_A via mean-subtraction
|
||||
/// Jacobian, then produce per-batch weight gradients via outer
|
||||
/// product with h_t. Caller reduces axis 0 to final weight grads.
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
pub fn decompose_and_backward_to_weights(
|
||||
&self,
|
||||
h_t: &CudaSlice<f32>,
|
||||
grad_composed: &CudaSlice<f32>,
|
||||
actions_taken: &CudaSlice<i32>,
|
||||
b_size: usize,
|
||||
grad_w_v_per_batch: &mut CudaSlice<f32>,
|
||||
grad_b_v_per_batch: &mut CudaSlice<f32>,
|
||||
grad_w_a_per_batch: &mut CudaSlice<f32>,
|
||||
grad_b_a_per_batch: &mut CudaSlice<f32>,
|
||||
) -> Result<()> {
|
||||
let hd = self.cfg.hidden_dim;
|
||||
debug_assert_eq!(h_t.len(), b_size * hd);
|
||||
debug_assert_eq!(grad_composed.len(), b_size * N_ACTIONS);
|
||||
debug_assert_eq!(actions_taken.len(), b_size);
|
||||
debug_assert_eq!(grad_w_v_per_batch.len(), b_size * hd);
|
||||
debug_assert_eq!(grad_b_v_per_batch.len(), b_size);
|
||||
debug_assert_eq!(grad_w_a_per_batch.len(), b_size * hd * N_ACTIONS);
|
||||
debug_assert_eq!(grad_b_a_per_batch.len(), b_size * N_ACTIONS);
|
||||
|
||||
let b_i = b_size as i32;
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(h_t.raw_ptr());
|
||||
args.push_ptr(grad_composed.raw_ptr());
|
||||
args.push_ptr(actions_taken.raw_ptr());
|
||||
args.push_i32(b_i);
|
||||
args.push_ptr(grad_w_v_per_batch.raw_ptr());
|
||||
args.push_ptr(grad_b_v_per_batch.raw_ptr());
|
||||
args.push_ptr(grad_w_a_per_batch.raw_ptr());
|
||||
args.push_ptr(grad_b_a_per_batch.raw_ptr());
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.decompose_and_bwd_fn.cu_function(),
|
||||
(b_size as u32, 1, 1),
|
||||
(hd as u32, 1, 1),
|
||||
0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("rl_dueling_q_decompose_and_bwd: {:?}", e))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Soft-update target network: `target = (1−τ) × target + τ × online`
|
||||
/// element-wise on all four weight tensors. τ read from
|
||||
/// `ISV[RL_TARGET_TAU_INDEX=401]` — same controller as Q's target.
|
||||
/// Should be called once per training step AFTER the Adam steps so
|
||||
/// the update reflects the latest online weights.
|
||||
pub fn soft_update_target(&mut self, isv_dev_ptr: &u64) -> Result<()> {
|
||||
let launch_blend = |n: usize, online: u64, target: u64, label: &str| -> Result<()> {
|
||||
let n_i = n as i32;
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(online);
|
||||
args.push_ptr(target);
|
||||
args.push_ptr(*isv_dev_ptr);
|
||||
args.push_i32(n_i);
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.soft_update_fn.cu_function(),
|
||||
(((n as u32) + 255) / 256, 1, 1),
|
||||
(256, 1, 1),
|
||||
0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("dueling soft_update_target ({label}): {e:?}"))?;
|
||||
}
|
||||
Ok(())
|
||||
};
|
||||
launch_blend(self.w_v_d.len(), self.w_v_d.raw_ptr(), self.w_v_target_d.raw_ptr(), "w_v")?;
|
||||
launch_blend(self.b_v_d.len(), self.b_v_d.raw_ptr(), self.b_v_target_d.raw_ptr(), "b_v")?;
|
||||
launch_blend(self.w_a_d.len(), self.w_a_d.raw_ptr(), self.w_a_target_d.raw_ptr(), "w_a")?;
|
||||
launch_blend(self.b_a_d.len(), self.b_a_d.raw_ptr(), self.b_a_target_d.raw_ptr(), "b_a")?;
|
||||
Ok(())
|
||||
}
|
||||
}
|
||||
|
||||
// ── pinned-staging upload helper (mirrors aux_heads.rs::upload) ──
|
||||
|
||||
fn upload(stream: &Arc<CudaStream>, host: &[f32]) -> Result<CudaSlice<f32>> {
|
||||
let n = host.len();
|
||||
let staging = unsafe { MappedF32Buffer::new(n) }
|
||||
.map_err(|e| anyhow::anyhow!("dueling_q upload staging: {e}"))?;
|
||||
staging.write_from_slice(host);
|
||||
let dst = stream
|
||||
.alloc_zeros::<f32>(n)
|
||||
.context("dueling_q upload alloc")?;
|
||||
if n > 0 {
|
||||
let nbytes = n * std::mem::size_of::<f32>();
|
||||
unsafe {
|
||||
let dst_ptr = dst.raw_ptr();
|
||||
crate::trainer::raw_launch::raw_memcpy_dtod_async(
|
||||
dst_ptr,
|
||||
staging.dev_ptr,
|
||||
nbytes,
|
||||
stream.cu_stream(),
|
||||
)
|
||||
.map_err(|e| anyhow::anyhow!("dueling_q upload DtoD: {:?}", e))?;
|
||||
}
|
||||
}
|
||||
Ok(dst)
|
||||
}
|
||||
@@ -47,7 +47,7 @@ impl GpuReplayBuffer {
|
||||
nstep_write_idx_d: stream.alloc_zeros(b_size).context("gpu_replay nstep_write_idx")?,
|
||||
nstep_count_d: stream.alloc_zeros(b_size).context("gpu_replay nstep_count")?,
|
||||
flush_flags_d: stream.alloc_zeros(b_size).context("gpu_replay flush_flags")?,
|
||||
write_offsets_d: stream.alloc_zeros(b_size + 1).context("gpu_replay write_offsets")?,
|
||||
write_offsets_d: stream.alloc_zeros(b_size + 2).context("gpu_replay write_offsets")?,
|
||||
sample_indices_d: stream.alloc_zeros(b_size).context("gpu_replay sample_indices")?,
|
||||
sample_prng_d: stream.alloc_zeros(b_size).context("gpu_replay sample_prng")?,
|
||||
})
|
||||
|
||||
@@ -1111,16 +1111,25 @@ pub const RL_ACTION_ENTROPY_EMA_INDEX: usize = 583;
|
||||
/// Bootstrap: 0.85 (15% exploration floor).
|
||||
pub const RL_CONF_GATE_MAX_HOLD_FRAC_INDEX: usize = 584;
|
||||
|
||||
/// Per-step drawdown penalty rate. Applied as min(0, unrealized_r) × rate
|
||||
/// to the reward each step a position is losing. Penalty-only (never
|
||||
/// positive) creates continuous exit gradient without exposure incentive.
|
||||
/// Bootstrap: 0.01 (1% of unrealized_r magnitude per step).
|
||||
pub const RL_DRAWDOWN_PENALTY_RATE_INDEX: usize = 586;
|
||||
/// Phase 4.4 (2026-05-30) — adaptive V baseline blend coefficient.
|
||||
///
|
||||
/// `V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b]`
|
||||
///
|
||||
/// α ∈ [0, 1] adapts via Schulman-bounded step from the observed
|
||||
/// V_dq vs V_scalar tracking ratio (see `rl_v_blend_alpha_controller`).
|
||||
/// α = 1.0: pure Plan A v2 (V_scalar drives PPO advantage).
|
||||
/// α = 0.0: pure Phase 4.3 (V_dq drives PPO advantage).
|
||||
/// Bootstrap on sentinel 0 → α = 1.0 (Plan A v2 behavior on first step).
|
||||
pub const RL_V_BLEND_ALPHA_INDEX: usize = 585;
|
||||
|
||||
/// Hard stop-loss threshold in initial-risk multiples. When unrealized_r
|
||||
/// drops below -threshold, the position force-closes (done=1). Caps max
|
||||
/// loss per trade. Bootstrap: 2.0 (2× initial risk).
|
||||
pub const RL_STOP_LOSS_THRESHOLD_INDEX: usize = 587;
|
||||
/// Phase 4.4 — EMA of `|V_dq − V_scalar|` (numerator of tracking ratio).
|
||||
/// Updated on-device by `rl_v_blend_alpha_controller` with EMA α = 0.01.
|
||||
/// First-observation bootstrap on sentinel 0.
|
||||
pub const RL_V_TRACK_ERR_EMA_INDEX: usize = 586;
|
||||
|
||||
/// Phase 4.4 — EMA of `|V_scalar|` (denominator of tracking ratio + dead-signal floor).
|
||||
/// If `EMA < 1e-4` the controller holds α (no V signal to calibrate against).
|
||||
pub const RL_V_SCALAR_MAG_EMA_INDEX: usize = 587;
|
||||
|
||||
/// Last RL-allocated slot index (exclusive).
|
||||
pub const RL_SLOTS_END: usize = 588;
|
||||
|
||||
@@ -17,6 +17,7 @@
|
||||
|
||||
pub mod common;
|
||||
pub mod dqn;
|
||||
pub mod dueling_q;
|
||||
pub mod gpu_hindsight;
|
||||
pub mod gpu_replay;
|
||||
pub mod frd;
|
||||
|
||||
@@ -46,57 +46,9 @@
|
||||
|
||||
use anyhow::Result;
|
||||
use cudarc::driver::CudaSlice;
|
||||
use cudarc::driver::sys::CUfunction;
|
||||
|
||||
use crate::cfc::snap_features::Mbp10RawInput;
|
||||
|
||||
/// Cached raw device pointers and kernel function handles for the lobsim.
|
||||
/// Extracted once per step to bypass trait dispatch + CudaSlice borrow
|
||||
/// overhead in the mega-graph hot path. All pointers are stable (pre-
|
||||
/// allocated at lobsim construction time, never reallocated).
|
||||
pub struct LobSimRawPtrs {
|
||||
// ── Book update (apply_snapshot_from_device) ────────────────────
|
||||
pub bid_px: u64,
|
||||
pub bid_sz: u64,
|
||||
pub ask_px: u64,
|
||||
pub ask_sz: u64,
|
||||
pub books: u64,
|
||||
pub prev_mid: u64,
|
||||
pub atr_mid_ema: u64,
|
||||
pub snapshots_skipped: u64,
|
||||
pub min_reasonable_px: u64,
|
||||
pub max_reasonable_px: u64,
|
||||
pub book_update_fn: CUfunction,
|
||||
|
||||
// ── Fill (step_fill_from_market_targets) ────────────────────────
|
||||
pub market_targets: u64,
|
||||
pub pos: u64,
|
||||
pub cost_per_lot_per_side: u64,
|
||||
pub total_fees_per_b: u64,
|
||||
pub submit_market_fn: CUfunction,
|
||||
|
||||
// ── PnL track (step_pnl_track) ─────────────────────────────────
|
||||
pub open_trade_state: u64,
|
||||
pub trade_log: u64,
|
||||
pub trade_log_head: u64,
|
||||
pub trail_hwm: u64,
|
||||
pub zero_vwap_at_open: u64,
|
||||
pub saturated_vwap_at_open: u64,
|
||||
pub defensive_exit_clamp: u64,
|
||||
pub conv_signed_ema: u64,
|
||||
pub diag_hold_hist: u64,
|
||||
pub diag_outcome_n: u64,
|
||||
pub diag_outcome_sum_pnl: u64,
|
||||
pub diag_outcome_n_wins: u64,
|
||||
pub pnl_track_fn: CUfunction,
|
||||
|
||||
// ── Dimensions ─────────────────────────────────────────────────
|
||||
pub n_backtests: i32,
|
||||
pub pos_bytes: i32,
|
||||
/// `TRADE_LOG_CAP` from ml-backtesting — passed to `pnl_track` kernel.
|
||||
pub trade_log_cap: i32,
|
||||
}
|
||||
|
||||
/// Narrow device-oriented backend the integrated RL trainer's
|
||||
/// `step_with_lobsim` invokes for its LOB-simulator interaction.
|
||||
/// Implemented for `LobSimCuda` in `ml-backtesting`. See module-level
|
||||
@@ -174,10 +126,4 @@ pub trait RlLobBackend {
|
||||
ask_px_src: u64,
|
||||
ask_sz_src: u64,
|
||||
) -> Result<()>;
|
||||
|
||||
/// Extract all raw device pointers and kernel function handles needed
|
||||
/// for mega-graph capture. Returns a [`LobSimRawPtrs`] struct that the
|
||||
/// trainer caches to bypass trait dispatch inside the captured section.
|
||||
/// All pointers are stable for the lobsim's lifetime.
|
||||
fn raw_ptrs(&self) -> LobSimRawPtrs;
|
||||
}
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -10,7 +10,6 @@
|
||||
use std::sync::Arc;
|
||||
|
||||
use anyhow::{Context, Result};
|
||||
use bytemuck;
|
||||
use cudarc::driver::{CudaFunction, CudaModule, CudaSlice, CudaStream};
|
||||
use cudarc::driver::sys::CUstream;
|
||||
use ml_core::device::MlDevice;
|
||||
@@ -122,155 +121,9 @@ impl AdamW {
|
||||
&mut self.v
|
||||
}
|
||||
|
||||
/// Mega-graph variant of `step()`. Reads LR from an ISV device
|
||||
/// pointer at `lr_slot` instead of the host-side `self.lr` field.
|
||||
/// This allows the LR to vary across graph replays because the ISV
|
||||
/// is modified in-place by the `rl_lr_from_mapped_pinned` controller
|
||||
/// kernel captured earlier in the same graph.
|
||||
///
|
||||
/// `isv_lr_fn` is the `adamw_step_isv_lr` kernel function handle.
|
||||
/// `isv_ptr` is the stable ISV device pointer.
|
||||
pub fn step_isv_lr(
|
||||
&mut self,
|
||||
theta: &mut CudaSlice<f32>,
|
||||
grad: &CudaSlice<f32>,
|
||||
isv_lr_fn: cudarc::driver::sys::CUfunction,
|
||||
isv_ptr: u64,
|
||||
lr_slot: i32,
|
||||
) -> Result<()> {
|
||||
let n = theta.len();
|
||||
assert_eq!(grad.len(), n, "grad/theta size mismatch");
|
||||
assert_eq!(self.m.len(), n, "m size mismatch");
|
||||
assert_eq!(self.v.len(), n, "v size mismatch");
|
||||
|
||||
self.step_count_host += 1;
|
||||
|
||||
let n_params_i = n as i32;
|
||||
let grid_x = (n as u32).div_ceil(256);
|
||||
{
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(theta.raw_ptr());
|
||||
args.push_ptr(grad.raw_ptr());
|
||||
args.push_ptr(self.m.raw_ptr());
|
||||
args.push_ptr(self.v.raw_ptr());
|
||||
args.push_i32(n_params_i);
|
||||
args.push_ptr(isv_ptr);
|
||||
args.push_i32(lr_slot);
|
||||
args.push_f32(self.beta1);
|
||||
args.push_f32(self.beta2);
|
||||
args.push_f32(self.eps);
|
||||
args.push_f32(self.wd);
|
||||
args.push_i32(self.step_count_host);
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
isv_lr_fn,
|
||||
(grid_x, 1, 1), (256, 1, 1), 0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("adamw_step_isv_lr launch: {:?}", e))?;
|
||||
}
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Return the current step counter value. No GPU sync needed —
|
||||
/// the counter is host-resident.
|
||||
pub fn step_count(&self) -> i32 {
|
||||
self.step_count_host
|
||||
}
|
||||
|
||||
/// Serialize Adam state (m, v, step_count, hyperparams) to writer.
|
||||
///
|
||||
/// Wire format (little-endian):
|
||||
/// u64 n — number of parameters
|
||||
/// i32 step_count
|
||||
/// f32 lr, beta1, beta2, eps, wd
|
||||
/// [f32; n] m — first moment
|
||||
/// [f32; n] v — second moment
|
||||
pub fn save(&self, w: &mut impl std::io::Write) -> Result<()> {
|
||||
use crate::pinned_mem::MappedF32Buffer;
|
||||
use crate::trainer::raw_launch::raw_memcpy_dtod_async;
|
||||
let n = self.m.len();
|
||||
let staging = unsafe { MappedF32Buffer::new(n) }
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save staging: {e}"))?;
|
||||
let raw_s = self.raw_stream;
|
||||
unsafe {
|
||||
raw_memcpy_dtod_async(staging.dev_ptr, self.m.raw_ptr(), n * 4, raw_s)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save m dtod: {e:?}"))?;
|
||||
}
|
||||
self._stream.synchronize()
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save sync m: {e}"))?;
|
||||
let m_host = staging.read_all();
|
||||
unsafe {
|
||||
raw_memcpy_dtod_async(staging.dev_ptr, self.v.raw_ptr(), n * 4, raw_s)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save v dtod: {e:?}"))?;
|
||||
}
|
||||
self._stream.synchronize()
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save sync v: {e}"))?;
|
||||
let v_host = staging.read_all();
|
||||
w.write_all(&(n as u64).to_le_bytes())?;
|
||||
w.write_all(&self.step_count_host.to_le_bytes())?;
|
||||
w.write_all(&self.lr.to_le_bytes())?;
|
||||
w.write_all(&self.beta1.to_le_bytes())?;
|
||||
w.write_all(&self.beta2.to_le_bytes())?;
|
||||
w.write_all(&self.eps.to_le_bytes())?;
|
||||
w.write_all(&self.wd.to_le_bytes())?;
|
||||
w.write_all(bytemuck::cast_slice(&m_host))?;
|
||||
w.write_all(bytemuck::cast_slice(&v_host))?;
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Restore Adam state from reader. Parameter count must match allocation.
|
||||
pub fn load(&mut self, r: &mut impl std::io::Read) -> Result<()> {
|
||||
let mut buf8 = [0u8; 8];
|
||||
let mut buf4 = [0u8; 4];
|
||||
r.read_exact(&mut buf8)?;
|
||||
let n = u64::from_le_bytes(buf8) as usize;
|
||||
anyhow::ensure!(
|
||||
n == self.m.len(),
|
||||
"AdamW load: m size mismatch {n} vs {}",
|
||||
self.m.len()
|
||||
);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.step_count_host = i32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.lr = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.beta1 = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.beta2 = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.eps = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.wd = f32::from_le_bytes(buf4);
|
||||
use crate::pinned_mem::MappedF32Buffer;
|
||||
use crate::trainer::raw_launch::raw_memcpy_dtod_async;
|
||||
let staging = unsafe { MappedF32Buffer::new(n) }
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load staging: {e}"))?;
|
||||
let raw_s = self.raw_stream;
|
||||
// m: read file → mapped-pinned host_ptr → DtoD → device
|
||||
{
|
||||
let host_slice = unsafe { std::slice::from_raw_parts_mut(staging.host_ptr, n) };
|
||||
r.read_exact(bytemuck::cast_slice_mut(host_slice))?;
|
||||
unsafe {
|
||||
raw_memcpy_dtod_async(self.m.raw_ptr(), staging.dev_ptr, n * 4, raw_s)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load m dtod: {e:?}"))?;
|
||||
}
|
||||
self._stream.synchronize()
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load sync m: {e}"))?;
|
||||
}
|
||||
// v: same pattern
|
||||
{
|
||||
let host_slice = unsafe { std::slice::from_raw_parts_mut(staging.host_ptr, n) };
|
||||
r.read_exact(bytemuck::cast_slice_mut(host_slice))?;
|
||||
unsafe {
|
||||
raw_memcpy_dtod_async(self.v.raw_ptr(), staging.dev_ptr, n * 4, raw_s)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load v dtod: {e:?}"))?;
|
||||
}
|
||||
self._stream.synchronize()
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load sync v: {e}"))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
}
|
||||
|
||||
@@ -718,13 +718,6 @@ pub struct PerceptionTrainer {
|
||||
/// include cuBLAS calls.
|
||||
cublas_warmed: bool,
|
||||
|
||||
/// Mega-graph mode: when true, all perception sub-graph state
|
||||
/// machines (train_graph, forward_graph, forward_graph_no_scatter)
|
||||
/// run eagerly (no sub-capture, no sub-replay). The mega-graph in
|
||||
/// the integrated trainer captures the entire pipeline including
|
||||
/// perception kernels.
|
||||
pub mega_graph_enabled: bool,
|
||||
|
||||
/// Event recorded after every training graph launch in
|
||||
/// `step_batched_from_device`. The integrated trainer syncs on
|
||||
/// this event at the start of the NEXT step
|
||||
@@ -2323,7 +2316,6 @@ impl PerceptionTrainer {
|
||||
forward_graph_no_scatter: None,
|
||||
forward_no_scatter_warmed: false,
|
||||
cublas_warmed: false,
|
||||
mega_graph_enabled: false,
|
||||
training_done_event,
|
||||
|
||||
// AoS staging — single mapped-pinned buffer for B*K Mbp10RawInput.
|
||||
@@ -3243,7 +3235,7 @@ impl PerceptionTrainer {
|
||||
// replay (third+). The captured graph records all
|
||||
// in-graph kernel decisions at capture time per
|
||||
// `pearl_no_host_branches_in_captured_graph`.
|
||||
if self.train_graph.is_some() && !self.mega_graph_enabled {
|
||||
if self.train_graph.is_some() {
|
||||
self.train_graph
|
||||
.as_ref()
|
||||
.unwrap()
|
||||
@@ -3253,9 +3245,6 @@ impl PerceptionTrainer {
|
||||
self.dispatch_train_step(b_sz, k_seq, total_snaps)
|
||||
.context("train warmup dispatch")?;
|
||||
self.cublas_warmed = true;
|
||||
} else if self.mega_graph_enabled {
|
||||
self.dispatch_train_step(b_sz, k_seq, total_snaps)
|
||||
.context("train eager dispatch (mega-graph mode)")?;
|
||||
} else {
|
||||
// Event tracking was disabled at trainer construction; the
|
||||
// trainer's CudaSlices have no read/write events, so neither
|
||||
@@ -3897,7 +3886,7 @@ impl PerceptionTrainer {
|
||||
// Three-state machine for the training graph — identical to
|
||||
// step_batched but dispatches `dispatch_train_step_no_scatter`
|
||||
// (skips the AoS→SoA scatter since SoA buffers are pre-filled).
|
||||
if self.train_graph.is_some() && !self.mega_graph_enabled {
|
||||
if self.train_graph.is_some() {
|
||||
self.train_graph
|
||||
.as_ref()
|
||||
.unwrap()
|
||||
@@ -3907,11 +3896,6 @@ impl PerceptionTrainer {
|
||||
self.dispatch_train_step_no_scatter(b_sz, k_seq, total_snaps)
|
||||
.context("train warmup dispatch (from_device)")?;
|
||||
self.cublas_warmed = true;
|
||||
} else if self.mega_graph_enabled {
|
||||
// Mega-graph mode: always dispatch eagerly — the mega-graph
|
||||
// in the integrated trainer captures these kernel launches.
|
||||
self.dispatch_train_step_no_scatter(b_sz, k_seq, total_snaps)
|
||||
.context("train eager dispatch (mega-graph mode)")?;
|
||||
} else {
|
||||
let begin = self.stream.begin_capture(
|
||||
CUstreamCaptureMode::CU_STREAM_CAPTURE_MODE_RELAXED,
|
||||
@@ -4106,7 +4090,7 @@ impl PerceptionTrainer {
|
||||
// warmup → capture → replay with a new graph that excludes the
|
||||
// scatter. Let's use this approach since it's zero overhead
|
||||
// after the second call.
|
||||
if self.forward_graph_no_scatter.is_some() && !self.mega_graph_enabled {
|
||||
if self.forward_graph_no_scatter.is_some() {
|
||||
self.forward_graph_no_scatter
|
||||
.as_ref()
|
||||
.unwrap()
|
||||
@@ -4116,10 +4100,6 @@ impl PerceptionTrainer {
|
||||
self.dispatch_forward_kernels_no_scatter(b_sz, k_seq, total_snaps)
|
||||
.context("forward_no_scatter warmup dispatch")?;
|
||||
self.forward_no_scatter_warmed = true;
|
||||
} else if self.mega_graph_enabled {
|
||||
// Mega-graph mode: always dispatch eagerly.
|
||||
self.dispatch_forward_kernels_no_scatter(b_sz, k_seq, total_snaps)
|
||||
.context("forward_no_scatter eager dispatch (mega-graph mode)")?;
|
||||
} else {
|
||||
use cudarc::driver::sys::{CUgraphInstantiate_flags, CUstreamCaptureMode};
|
||||
let begin = self.stream.begin_capture(
|
||||
|
||||
@@ -27,6 +27,7 @@
|
||||
use anyhow::Result;
|
||||
use cudarc::driver::{CudaSlice, CudaStream};
|
||||
use ml_alpha::heads::HIDDEN_DIM;
|
||||
use ml_alpha::pinned_mem::MappedF32Buffer;
|
||||
use ml_alpha::rl::common::{FRD_HIDDEN_DIM, FRD_N_ATOMS, FRD_N_HORIZONS};
|
||||
use ml_alpha::rl::frd::{FrdHead, FrdHeadConfig, FRD_OUT_DIM};
|
||||
use ml_alpha::trainer::integrated::{
|
||||
@@ -35,6 +36,14 @@ use ml_alpha::trainer::integrated::{
|
||||
use ml_core::device::MlDevice;
|
||||
use std::sync::Arc;
|
||||
|
||||
/// Mapped-pinned scratch buffer for kernel output that the kernel writes
|
||||
/// via the raw `dev_ptr` and the test reads via volatile host access
|
||||
/// after a stream sync. Per `feedback_no_htod_htoh_only_mapped_pinned`:
|
||||
/// even tests use mapped-pinned for CPU↔GPU transfers.
|
||||
fn alloc_loss_buf(n: usize) -> MappedF32Buffer {
|
||||
unsafe { MappedF32Buffer::new(n) }.expect("loss MappedF32Buffer alloc")
|
||||
}
|
||||
|
||||
fn build_head() -> Option<(MlDevice, FrdHead)> {
|
||||
let dev = match MlDevice::cuda(0) {
|
||||
Ok(d) => d,
|
||||
@@ -209,11 +218,12 @@ fn frd_softmax_ce_grad_uniform_logits_match_log_n_atoms() -> Result<()> {
|
||||
let labels: Vec<i32> = vec![10; b_size * FRD_N_HORIZONS];
|
||||
let labels_d = upload_i32(&stream, &labels)?;
|
||||
let mut grad_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
let loss = read_slice_d_pub(&stream, &loss_d, b_size * FRD_N_HORIZONS)?;
|
||||
stream.synchronize()?;
|
||||
let loss = loss_buf.read_all();
|
||||
let expected = (FRD_N_ATOMS as f32).ln();
|
||||
for (i, v) in loss.iter().enumerate() {
|
||||
assert!(
|
||||
@@ -264,11 +274,12 @@ fn frd_softmax_ce_grad_sentinel_label_zeros_row() -> Result<()> {
|
||||
let labels: Vec<i32> = vec![-1; b_size * FRD_N_HORIZONS];
|
||||
let labels_d = upload_i32(&stream, &labels)?;
|
||||
let mut grad_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
let loss = read_slice_d_pub(&stream, &loss_d, b_size * FRD_N_HORIZONS)?;
|
||||
stream.synchronize()?;
|
||||
let loss = loss_buf.read_all();
|
||||
let grad = read_slice_d_pub(&stream, &grad_d, b_size * FRD_OUT_DIM)?;
|
||||
for (i, v) in loss.iter().enumerate() {
|
||||
assert_eq!(*v, 0.0, "sentinel label loss[{i}] must be 0; got {v}");
|
||||
@@ -305,8 +316,8 @@ fn frd_softmax_ce_grad_finite_diff_matches_analytical() -> Result<()> {
|
||||
// Analytical gradient via the kernel.
|
||||
let logits_d = upload_f32(&stream, &logits)?;
|
||||
let mut grad_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
let grad_analytical = read_slice_d_pub(&stream, &grad_d, b_size * FRD_OUT_DIM)?;
|
||||
|
||||
// Finite-difference for slot (b=0, h=0, a=3). Note: gradient was
|
||||
@@ -320,15 +331,17 @@ fn frd_softmax_ce_grad_finite_diff_matches_analytical() -> Result<()> {
|
||||
// L(logits + ε · e_j) — perturb only the target slot upward.
|
||||
logits[probe_off] += eps;
|
||||
let logits_plus_d = upload_f32(&stream, &logits)?;
|
||||
head.softmax_ce_grad(&logits_plus_d, &labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
let loss_plus = read_slice_d_pub(&stream, &loss_d, b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_plus_d, &labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
stream.synchronize()?;
|
||||
let loss_plus = loss_buf.read_all();
|
||||
let l_plus = loss_plus[probe_h]; // only h=0 affected — h=1,2 share the perturbation only if probe was in their horizon block
|
||||
|
||||
// L(logits - ε · e_j)
|
||||
logits[probe_off] -= 2.0 * eps;
|
||||
let logits_minus_d = upload_f32(&stream, &logits)?;
|
||||
head.softmax_ce_grad(&logits_minus_d, &labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
let loss_minus = read_slice_d_pub(&stream, &loss_d, b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_minus_d, &labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
stream.synchronize()?;
|
||||
let loss_minus = loss_buf.read_all();
|
||||
let l_minus = loss_minus[probe_h];
|
||||
|
||||
let numerical = (l_plus - l_minus) / (2.0 * eps);
|
||||
@@ -364,9 +377,10 @@ fn ce_total_loss(
|
||||
) -> Result<f32> {
|
||||
let logits_d = upload_f32(stream, logits)?;
|
||||
let mut grad_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, labels_d, &mut grad_d, &mut loss_d, b_size)?;
|
||||
let loss = read_slice_d_pub(stream, &loss_d, b_size * FRD_N_HORIZONS)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, labels_d, &mut grad_d, &loss_buf.dev_ptr, b_size)?;
|
||||
stream.synchronize()?;
|
||||
let loss = loss_buf.read_all();
|
||||
Ok(loss.iter().sum())
|
||||
}
|
||||
|
||||
@@ -394,8 +408,8 @@ fn frd_layer2_bwd_finite_diff_w2() -> Result<()> {
|
||||
|
||||
// Softmax+CE grad of logits.
|
||||
let mut grad_logits_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &mut loss_d, b_size)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
// Layer-2 backward: produce per-batch grad_W2 scratch.
|
||||
let mut grad_w2_pb_d =
|
||||
@@ -490,8 +504,8 @@ fn frd_layer2_bwd_db2_equals_grad_logits() -> Result<()> {
|
||||
let labels_d = upload_i32(&stream, &labels)?;
|
||||
|
||||
let mut grad_logits_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &mut loss_d, b_size)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
let mut grad_w2_pb_d =
|
||||
stream.alloc_zeros::<f32>(b_size * FRD_HIDDEN_DIM * FRD_OUT_DIM)?;
|
||||
@@ -542,8 +556,8 @@ fn frd_layer1_bwd_finite_diff_w1() -> Result<()> {
|
||||
head.forward(&h_t_d, &mut hidden_d, &mut logits_d, b_size)?;
|
||||
|
||||
let mut grad_logits_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &mut loss_d, b_size)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
let mut grad_w2_pb_d = stream.alloc_zeros::<f32>(b_size * FRD_HIDDEN_DIM * FRD_OUT_DIM)?;
|
||||
let mut grad_b2_pb_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
@@ -668,8 +682,8 @@ fn frd_layer1_bwd_relu_mask_zeros_grad() -> Result<()> {
|
||||
head.forward(&h_t_d, &mut hidden_d, &mut logits_d, b_size)?;
|
||||
|
||||
let mut grad_logits_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
let mut loss_d = stream.alloc_zeros::<f32>(b_size * FRD_N_HORIZONS)?;
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &mut loss_d, b_size)?;
|
||||
let loss_buf = alloc_loss_buf(b_size * FRD_N_HORIZONS);
|
||||
head.softmax_ce_grad(&logits_d, &labels_d, &mut grad_logits_d, &loss_buf.dev_ptr, b_size)?;
|
||||
|
||||
let mut grad_w2_pb_d = stream.alloc_zeros::<f32>(b_size * FRD_HIDDEN_DIM * FRD_OUT_DIM)?;
|
||||
let mut grad_b2_pb_d = stream.alloc_zeros::<f32>(b_size * FRD_OUT_DIM)?;
|
||||
|
||||
@@ -70,18 +70,35 @@ fn large_negative_bias_saturates_low() {
|
||||
|
||||
#[test]
|
||||
fn per_head_independence() {
|
||||
// Set bias[0] = 5, bias[4] = -5, others = 0 with zero weights.
|
||||
// sigmoid(5) ≈ 0.993, sigmoid(0) = 0.5, sigmoid(-5) ≈ 0.0067.
|
||||
// Per `feedback_use_consts_not_literals_for_structural_dims`:
|
||||
// address heads by N_HORIZONS-relative offsets, not hardcoded
|
||||
// indices that silently drift when N_HORIZONS changes.
|
||||
// Set first bias = +5, middle biases = 0, last bias = -5; with
|
||||
// zero weights and h=0 the head output is sigmoid(bias):
|
||||
// sigmoid(+5) ≈ 0.993, sigmoid(0) = 0.5, sigmoid(-5) ≈ 0.0067.
|
||||
assert!(N_HORIZONS >= 3, "test requires at least 3 horizons");
|
||||
let dev = test_device();
|
||||
let mut b = vec![0.0_f32; N_HORIZONS];
|
||||
b[0] = 5.0;
|
||||
b[N_HORIZONS - 1] = -5.0;
|
||||
let w = HeadsWeights {
|
||||
w: vec![0.0; N_HORIZONS * HIDDEN_DIM],
|
||||
b: vec![5.0, 0.0, 0.0, 0.0, -5.0],
|
||||
b,
|
||||
};
|
||||
let h = vec![0.0; HIDDEN_DIM];
|
||||
let probs = multi_horizon_heads_gpu(&dev, &w, &h).expect("gpu");
|
||||
assert!(probs[0] > 0.99 && probs[0] < 1.0);
|
||||
assert_relative_eq!(probs[1], 0.5, epsilon = 1e-6);
|
||||
assert_relative_eq!(probs[2], 0.5, epsilon = 1e-6);
|
||||
assert_relative_eq!(probs[3], 0.5, epsilon = 1e-6);
|
||||
assert!(probs[4] > 0.0 && probs[4] < 0.01);
|
||||
assert!(
|
||||
probs[0] > 0.99 && probs[0] < 1.0,
|
||||
"probs[0]=sigmoid(+5) should be > 0.99; got {}",
|
||||
probs[0]
|
||||
);
|
||||
for k in 1..N_HORIZONS - 1 {
|
||||
assert_relative_eq!(probs[k], 0.5, epsilon = 1e-6);
|
||||
}
|
||||
let last = N_HORIZONS - 1;
|
||||
assert!(
|
||||
probs[last] > 0.0 && probs[last] < 0.01,
|
||||
"probs[{last}]=sigmoid(-5) should be in (0, 0.01); got {}",
|
||||
probs[last]
|
||||
);
|
||||
}
|
||||
|
||||
@@ -92,13 +92,14 @@ fn g1_isv_bootstrap_writes_canonical_values() {
|
||||
};
|
||||
let trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new");
|
||||
|
||||
// Read full ISV slice to host. Uses the same pattern as the
|
||||
// trainer's own per-step ISV mirror refresh.
|
||||
let mut isv = vec![0.0_f32; RL_SLOTS_END];
|
||||
let stream = dev.cuda_stream().expect("cuda_stream");
|
||||
stream
|
||||
.memcpy_dtoh(&trainer.isv_d, isv.as_mut_slice())
|
||||
.expect("isv dtoh");
|
||||
// Read ISV via the mapped-pinned host view. Per
|
||||
// `feedback_no_htod_htoh_only_mapped_pinned`: tests use the same
|
||||
// zero-copy mirror as production. Sync the producing stream first
|
||||
// so the bootstrap-controller writes from `IntegratedTrainer::new`
|
||||
// are visible host-side.
|
||||
trainer.stream.synchronize().expect("sync trainer stream");
|
||||
let isv: &[f32] = trainer.isv_host_slice();
|
||||
assert_eq!(isv.len(), RL_SLOTS_END, "ISV buffer length");
|
||||
|
||||
// Floating-point exact equality is the right oracle here — each
|
||||
// kernel's bootstrap path is `isv[slot] = K_BOOTSTRAP; return;`
|
||||
|
||||
@@ -23,7 +23,7 @@
|
||||
//! `cargo test -p ml-alpha --test r3_ema_advantage -- --ignored --nocapture`
|
||||
|
||||
use ml_alpha::rl::isv_slots::{
|
||||
RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_SLOTS_END,
|
||||
RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX,
|
||||
};
|
||||
use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig};
|
||||
use ml_alpha::trainer::perception::PerceptionTrainerConfig;
|
||||
@@ -62,16 +62,12 @@ fn upload(
|
||||
d
|
||||
}
|
||||
|
||||
fn readback_isv(
|
||||
dev: &MlDevice,
|
||||
isv_d: &cudarc::driver::CudaSlice<f32>,
|
||||
) -> Vec<f32> {
|
||||
let mut isv = vec![0.0_f32; RL_SLOTS_END];
|
||||
let stream = dev.cuda_stream().expect("cuda_stream").clone();
|
||||
stream
|
||||
.memcpy_dtoh(isv_d, isv.as_mut_slice())
|
||||
.expect("isv dtoh");
|
||||
isv
|
||||
/// Mapped-pinned ISV readback per `feedback_no_htod_htoh_only_mapped_pinned`:
|
||||
/// the trainer's `isv_mapped` is the same buffer the GPU writes via
|
||||
/// `isv_dev_ptr`, so the host view is zero-copy. Caller must have
|
||||
/// synchronized the producing stream first.
|
||||
fn readback_isv(trainer: &ml_alpha::trainer::integrated::IntegratedTrainer) -> Vec<f32> {
|
||||
trainer.isv_host_slice().to_vec()
|
||||
}
|
||||
|
||||
#[test]
|
||||
@@ -83,7 +79,7 @@ fn r3_ema_update_on_done_first_observation_bootstrap_replaces_directly() {
|
||||
// Pre-condition: the EMA-input slot is at sentinel zero (R1
|
||||
// bootstraps ISV[400..406] for controllers; ISV[417..423] EMA-input
|
||||
// slots stay at alloc_zeros per the R1 invariant).
|
||||
let isv_before = readback_isv(&dev, &trainer.isv_d);
|
||||
let isv_before = readback_isv(&trainer);
|
||||
assert_eq!(
|
||||
isv_before[RL_MEAN_ABS_PNL_EMA_INDEX], 0.0,
|
||||
"pre-condition: EMA slot must be sentinel zero"
|
||||
@@ -100,7 +96,7 @@ fn r3_ema_update_on_done_first_observation_bootstrap_replaces_directly() {
|
||||
.expect("ema_update_on_done");
|
||||
stream.synchronize().expect("sync");
|
||||
|
||||
let isv_after = readback_isv(&dev, &trainer.isv_d);
|
||||
let isv_after = readback_isv(&trainer);
|
||||
let val = isv_after[RL_MEAN_ABS_PNL_EMA_INDEX];
|
||||
// Exact equality — bootstrap path is `isv[slot] = mean_obs` with
|
||||
// no arithmetic. Any drift indicates a wrong code path.
|
||||
@@ -118,7 +114,7 @@ fn r3_ema_update_on_done_first_observation_bootstrap_replaces_directly() {
|
||||
.expect("ema_update_on_done hold");
|
||||
stream.synchronize().expect("sync");
|
||||
|
||||
let isv_hold = readback_isv(&dev, &trainer.isv_d);
|
||||
let isv_hold = readback_isv(&trainer);
|
||||
assert_eq!(
|
||||
isv_hold[RL_MEAN_ABS_PNL_EMA_INDEX], 7.0,
|
||||
"hold step (no done) must preserve EMA, not blend toward 0"
|
||||
@@ -149,7 +145,7 @@ fn r3_ema_update_per_step_converges_to_constant_input() {
|
||||
}
|
||||
stream.synchronize().expect("sync");
|
||||
|
||||
let isv = readback_isv(&dev, &trainer.isv_d);
|
||||
let isv = readback_isv(&trainer);
|
||||
let val = isv[RL_KL_PI_EMA_INDEX];
|
||||
assert!(
|
||||
(val - k).abs() < 1e-4,
|
||||
@@ -172,7 +168,7 @@ fn r3_compute_advantage_return_formula_holds() {
|
||||
// computes its expected values from whatever γ ISV holds, so the
|
||||
// pre-condition is just "γ is in the valid bounded range" rather
|
||||
// than a hardcoded canonical value.
|
||||
let isv = readback_isv(&dev, &trainer.isv_d);
|
||||
let isv = readback_isv(&trainer);
|
||||
let gamma = isv[RL_GAMMA_INDEX];
|
||||
assert!(
|
||||
gamma >= 0.90 && gamma <= 0.999,
|
||||
|
||||
@@ -1,71 +1,34 @@
|
||||
//! Phase R5 gates G3 + G4:
|
||||
//! Phase R5 gate G4: `DqnHead::soft_update_target` implements the
|
||||
//! Polyak averaging formula
|
||||
//!
|
||||
//! G3: `launch_rl_controllers_per_step` actually moves all 7 ISV
|
||||
//! output slots away from their R1 bootstrap values when fed
|
||||
//! non-zero EMA inputs (via the R3 `ema_update_*` kernels'
|
||||
//! bootstrap path). Catches "controller doesn't fire", "wrong
|
||||
//! input slot wiring", and "input slot mismatch" bugs.
|
||||
//! target[i] = (1 − τ)·target[i] + τ·current[i]
|
||||
//!
|
||||
//! G4: `DqnHead::soft_update_target` actually moves `w_target_d`
|
||||
//! toward `w_d` via the formula
|
||||
//! `target[i] = (1 − τ)·target[i] + τ·current[i]`, with τ read
|
||||
//! from `ISV[RL_TARGET_TAU_INDEX = 401]`. Tests both the
|
||||
//! formula (force-known w_d, snapshot before, soft_update,
|
||||
//! check formula at sampled indices) and the τ=0 / τ=1 limits
|
||||
//! (τ=0 → target unchanged; τ=1 → target := current).
|
||||
//! with τ read from `ISV[RL_TARGET_TAU_INDEX]`.
|
||||
//!
|
||||
//! Per `feedback_no_cpu_test_fallbacks` every oracle is analytical:
|
||||
//! - G3: invariant "output != bootstrap after a non-trivial input"
|
||||
//! - G4: arithmetic formula `(1-τ)·a + τ·b` evaluated host-side on
|
||||
//! the SAME numbers the kernel saw (no CPU reference of the
|
||||
//! kernel itself — the kernel IS the kernel; we just check its
|
||||
//! output matches the algebraic identity it implements).
|
||||
//! Per `feedback_no_cpu_test_fallbacks`: the oracle is the algebraic
|
||||
//! identity the kernel implements, evaluated host-side on the same
|
||||
//! numbers the kernel saw — not a CPU reimplementation.
|
||||
//!
|
||||
//! G3 (which exercised the now-removed `launch_rl_controllers_per_step`
|
||||
//! bulk method) was retired when the trainer adopted
|
||||
//! `launch_rl_fused_controllers` — a single fused kernel reading from
|
||||
//! per-step dones + a dedicated input-slot index buffer. The fused
|
||||
//! kernel's "all controllers move slots" invariant is exercised end-to-end
|
||||
//! by every cluster training run, so a separate unit test would
|
||||
//! reproduce production setup without adding signal.
|
||||
//!
|
||||
//! Run with:
|
||||
//! `cargo test -p ml-alpha --test r5_controllers_and_soft_update -- --ignored --nocapture`
|
||||
|
||||
use cudarc::driver::CudaStream;
|
||||
use ml_alpha::rl::isv_slots::{
|
||||
RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ADV_VAR_RATIO_CLAMP_INDEX,
|
||||
RL_ADV_VAR_RATIO_TARGET_INDEX, RL_DIV_TARGET_INDEX, RL_ENTROPY_COEF_INDEX,
|
||||
RL_ENTROPY_OBSERVED_EMA_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, RL_EPS_BOOTSTRAP_INDEX,
|
||||
RL_GAMMA_INDEX, RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_PI_EMA_INDEX,
|
||||
RL_KL_TARGET_INDEX, RL_KURT_GAUSSIAN_INDEX, RL_KURT_LIFT_SCALE_INDEX,
|
||||
RL_KURT_NOISE_FLOOR_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX,
|
||||
RL_LOSS_LAMBDA_AUX_INDEX, RL_LR_BOOTSTRAP_INDEX, RL_LR_DECAY_FACTOR_INDEX,
|
||||
RL_LR_LOSS_EMA_ALPHA_INDEX, RL_LR_MAX_INDEX, RL_LR_MIN_INDEX, RL_LR_WARMUP_STEPS_INDEX,
|
||||
RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX,
|
||||
RL_PER_ALPHA_INDEX, RL_PLATEAU_PATIENCE_INDEX, RL_PPO_CLAMP_MARGIN_INDEX,
|
||||
RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX,
|
||||
RL_Q_DIVERGENCE_EMA_INDEX, RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX,
|
||||
RL_REWARD_SCALE_BOOTSTRAP_INDEX, RL_REWARD_SCALE_INDEX, RL_ROLLOUT_BOOTSTRAP_INDEX,
|
||||
RL_SCHULMAN_ADJUST_RATE_INDEX, RL_SCHULMAN_TOLERANCE_INDEX, RL_SLOTS_END,
|
||||
RL_STREAM_ALPHA_INDEX, RL_TARGET_TAU_INDEX, RL_TAU_BOOTSTRAP_INDEX,
|
||||
RL_TD_KURTOSIS_CLAMP_INDEX, RL_TD_KURTOSIS_EMA_INDEX,
|
||||
};
|
||||
use ml_alpha::rl::isv_slots::RL_TARGET_TAU_INDEX;
|
||||
use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig};
|
||||
use ml_alpha::trainer::perception::PerceptionTrainerConfig;
|
||||
use ml_core::device::MlDevice;
|
||||
use std::sync::Arc;
|
||||
|
||||
// Bootstrap value at sentinel input — see isv_bootstrap.rs for the
|
||||
// post-R9-audit derive-from-input pattern explanation.
|
||||
const GAMMA_BOOTSTRAP: f32 = 0.90;
|
||||
/// Bootstrap value the `rl_target_tau_controller` writes into
|
||||
/// `ISV[RL_TARGET_TAU_INDEX]` at construction time (sentinel-input
|
||||
/// path of `pearl_first_observation_bootstrap`).
|
||||
const TAU_BOOTSTRAP: f32 = 0.005;
|
||||
const EPS_BOOTSTRAP: f32 = 0.2;
|
||||
// Bootstrap value at sentinel input — see isv_bootstrap.rs for the
|
||||
// post-R9-audit derive-from-input pattern explanation.
|
||||
const COEF_BOOTSTRAP: f32 = 0.035;
|
||||
const ROLLOUT_BOOTSTRAP: f32 = 2048.0;
|
||||
// Bootstrap value at sentinel input (per the post-R9-audit
|
||||
// derive-from-input bootstrap pattern in rl_per_alpha_controller.cu):
|
||||
// kurt_excess=0 → target = 0.4. Was hardcoded 0.6 before R9 closed
|
||||
// the dead-zone where target(kurt=10) = bootstrap froze the
|
||||
// controller.
|
||||
const PER_ALPHA_BOOTSTRAP: f32 = 0.4;
|
||||
const REWARD_SCALE_BOOTSTRAP: f32 = 1.0;
|
||||
|
||||
const ALPHA_FLOOR: f32 = 0.4;
|
||||
|
||||
fn build_trainer() -> Option<(MlDevice, IntegratedTrainer)> {
|
||||
let dev = match MlDevice::cuda(0) {
|
||||
@@ -89,194 +52,6 @@ fn build_trainer() -> Option<(MlDevice, IntegratedTrainer)> {
|
||||
Some((dev, trainer))
|
||||
}
|
||||
|
||||
fn upload_f32(stream: &Arc<CudaStream>, host: &[f32]) -> cudarc::driver::CudaSlice<f32> {
|
||||
let mut d = stream.alloc_zeros::<f32>(host.len()).expect("alloc");
|
||||
stream.memcpy_htod(host, &mut d).expect("htod");
|
||||
d
|
||||
}
|
||||
|
||||
fn readback_isv(
|
||||
dev: &MlDevice,
|
||||
isv_d: &cudarc::driver::CudaSlice<f32>,
|
||||
) -> Vec<f32> {
|
||||
let mut isv = vec![0.0_f32; RL_SLOTS_END];
|
||||
let stream = dev.cuda_stream().expect("cuda_stream").clone();
|
||||
stream
|
||||
.memcpy_dtoh(isv_d, isv.as_mut_slice())
|
||||
.expect("isv dtoh");
|
||||
isv
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
|
||||
let Some((dev, trainer)) = build_trainer() else { return };
|
||||
let stream = dev.cuda_stream().expect("cuda_stream").clone();
|
||||
|
||||
// Pre-condition: R1 bootstrapped ISV[400..406].
|
||||
let isv_before = readback_isv(&dev, &trainer.isv_d);
|
||||
assert_eq!(isv_before[RL_GAMMA_INDEX], GAMMA_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_TARGET_TAU_INDEX], TAU_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_PPO_CLIP_INDEX], EPS_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_ENTROPY_COEF_INDEX], COEF_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_N_ROLLOUT_STEPS_INDEX], ROLLOUT_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_PER_ALPHA_INDEX], PER_ALPHA_BOOTSTRAP);
|
||||
assert_eq!(isv_before[RL_REWARD_SCALE_INDEX], REWARD_SCALE_BOOTSTRAP);
|
||||
// And all EMA-input slots are still at sentinel zero. Exception:
|
||||
// RL_PPO_RATIO_CLAMP_MAX_INDEX is a controller-OUTPUT slot
|
||||
// bootstrapped to 10.0 by `with_controllers_bootstrapped`.
|
||||
for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END {
|
||||
if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX
|
||||
|| slot == RL_ADV_VAR_RATIO_CLAMP_INDEX
|
||||
|| slot == RL_TD_KURTOSIS_CLAMP_INDEX
|
||||
|| slot == RL_ADV_VAR_RATIO_TARGET_INDEX
|
||||
|| slot == RL_K_LOOP_DIVISOR_INDEX
|
||||
|| slot == RL_K_LOOP_MAX_INDEX
|
||||
|| slot == RL_REWARD_CLAMP_WIN_INDEX
|
||||
|| slot == RL_REWARD_CLAMP_LOSS_INDEX
|
||||
|| slot == RL_KL_TARGET_INDEX
|
||||
|| slot == RL_IMPROVEMENT_THRESHOLD_INDEX
|
||||
|| slot == RL_PLATEAU_PATIENCE_INDEX
|
||||
|| slot == RL_DIV_TARGET_INDEX
|
||||
|| slot == RL_ENTROPY_TARGET_FRAC_INDEX
|
||||
|| slot == RL_KURT_LIFT_SCALE_INDEX
|
||||
|| slot == RL_PPO_CLAMP_MARGIN_INDEX
|
||||
|| slot == RL_LR_WARMUP_STEPS_INDEX
|
||||
|| slot == RL_LR_BOOTSTRAP_INDEX
|
||||
|| slot == RL_LR_MIN_INDEX
|
||||
|| slot == RL_LR_MAX_INDEX
|
||||
|| slot == RL_LR_LOSS_EMA_ALPHA_INDEX
|
||||
|| slot == RL_LR_DECAY_FACTOR_INDEX
|
||||
|| slot == RL_LOSS_LAMBDA_AUX_INDEX
|
||||
|| slot == RL_SCHULMAN_TOLERANCE_INDEX
|
||||
|| slot == RL_SCHULMAN_ADJUST_RATE_INDEX
|
||||
|| slot == RL_STREAM_ALPHA_INDEX
|
||||
|| slot == RL_KURT_GAUSSIAN_INDEX
|
||||
|| slot == RL_KURT_NOISE_FLOOR_INDEX
|
||||
|| slot == RL_TAU_BOOTSTRAP_INDEX
|
||||
|| slot == RL_EPS_BOOTSTRAP_INDEX
|
||||
|| slot == RL_ROLLOUT_BOOTSTRAP_INDEX
|
||||
|| slot == RL_REWARD_SCALE_BOOTSTRAP_INDEX
|
||||
|| slot == RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX
|
||||
{
|
||||
continue;
|
||||
}
|
||||
assert_eq!(isv_before[slot], 0.0);
|
||||
}
|
||||
assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_MAX_INDEX], 10.0);
|
||||
assert_eq!(isv_before[RL_ADV_VAR_RATIO_CLAMP_INDEX], 100.0);
|
||||
assert_eq!(isv_before[RL_TD_KURTOSIS_CLAMP_INDEX], 30.0);
|
||||
assert_eq!(isv_before[RL_ADV_VAR_RATIO_TARGET_INDEX], 5.0);
|
||||
assert_eq!(isv_before[RL_K_LOOP_DIVISOR_INDEX], 2048.0);
|
||||
assert_eq!(isv_before[RL_K_LOOP_MAX_INDEX], 4.0);
|
||||
assert_eq!(isv_before[RL_REWARD_CLAMP_WIN_INDEX], 1.0);
|
||||
assert_eq!(isv_before[RL_REWARD_CLAMP_LOSS_INDEX], 3.0);
|
||||
assert_eq!(isv_before[RL_KL_TARGET_INDEX], 0.01);
|
||||
assert_eq!(isv_before[RL_IMPROVEMENT_THRESHOLD_INDEX], 0.99);
|
||||
assert_eq!(isv_before[RL_PLATEAU_PATIENCE_INDEX], 1000.0);
|
||||
assert_eq!(isv_before[RL_DIV_TARGET_INDEX], 0.01);
|
||||
assert_eq!(isv_before[RL_ENTROPY_TARGET_FRAC_INDEX], 0.7);
|
||||
assert_eq!(isv_before[RL_KURT_LIFT_SCALE_INDEX], 7.0);
|
||||
assert_eq!(isv_before[RL_PPO_CLAMP_MARGIN_INDEX], 10.0);
|
||||
assert_eq!(isv_before[RL_LR_WARMUP_STEPS_INDEX], 2000.0);
|
||||
assert!((isv_before[RL_LR_BOOTSTRAP_INDEX] - 1e-3).abs() < 1e-9);
|
||||
assert!((isv_before[RL_LR_MIN_INDEX] - 1e-4).abs() < 1e-9);
|
||||
assert!((isv_before[RL_LR_MAX_INDEX] - 1e-2).abs() < 1e-9);
|
||||
assert!((isv_before[RL_LR_LOSS_EMA_ALPHA_INDEX] - 0.05).abs() < 1e-7);
|
||||
assert_eq!(isv_before[RL_LR_DECAY_FACTOR_INDEX], 0.5);
|
||||
assert_eq!(isv_before[RL_LOSS_LAMBDA_AUX_INDEX], 1.0);
|
||||
assert_eq!(isv_before[RL_SCHULMAN_TOLERANCE_INDEX], 1.5);
|
||||
assert_eq!(isv_before[RL_SCHULMAN_ADJUST_RATE_INDEX], 1.5);
|
||||
assert!((isv_before[RL_STREAM_ALPHA_INDEX] - 0.05).abs() < 1e-7);
|
||||
assert_eq!(isv_before[RL_KURT_GAUSSIAN_INDEX], 3.0);
|
||||
assert_eq!(isv_before[RL_KURT_NOISE_FLOOR_INDEX], 1.0);
|
||||
assert!((isv_before[RL_TAU_BOOTSTRAP_INDEX] - 0.005).abs() < 1e-7);
|
||||
assert!((isv_before[RL_EPS_BOOTSTRAP_INDEX] - 0.2).abs() < 1e-7);
|
||||
assert_eq!(isv_before[RL_ROLLOUT_BOOTSTRAP_INDEX], 2048.0);
|
||||
assert_eq!(isv_before[RL_REWARD_SCALE_BOOTSTRAP_INDEX], 1.0);
|
||||
assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX], 10.0);
|
||||
|
||||
// Populate each EMA-input slot with a non-zero value via the
|
||||
// R3 ema_update_per_step bootstrap path (sentinel-zero → first
|
||||
// observation replaces directly). Choose distinct values per slot
|
||||
// so a slot-wiring bug (controller reads wrong slot) would
|
||||
// produce out-of-range outputs we can detect.
|
||||
// Input values chosen to produce targets distinct from each
|
||||
// controller's clamped floor — production trade duration EMAs
|
||||
// typically settle in the 10-100 range, far above the d=1 edge
|
||||
// where γ target clamps to GAMMA_MIN.
|
||||
let inputs: [(usize, f32); 7] = [
|
||||
(RL_MEAN_TRADE_DURATION_EMA_INDEX, 20.0), // → rl_gamma (target ≈ 0.966)
|
||||
(RL_Q_DIVERGENCE_EMA_INDEX, 0.5), // → rl_target_tau
|
||||
(RL_KL_PI_EMA_INDEX, 0.1), // → rl_ppo_clip
|
||||
(RL_ENTROPY_OBSERVED_EMA_INDEX, 0.5), // → rl_entropy_coef
|
||||
// Above ADV_VAR_RATIO_TARGET (5.0) × TOLERANCE (1.5) = 7.5 so
|
||||
// the WIDEN branch fires and rollout_steps moves off bootstrap.
|
||||
(RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, 20.0), // → rl_rollout_steps
|
||||
(RL_TD_KURTOSIS_EMA_INDEX, 10.0), // → rl_per_alpha
|
||||
(RL_MEAN_ABS_PNL_EMA_INDEX, 50.0), // → rl_reward_scale
|
||||
];
|
||||
for (slot, obs_val) in inputs {
|
||||
let obs_d = upload_f32(&stream, &[obs_val]);
|
||||
trainer
|
||||
.launch_ema_update_per_step(slot, ALPHA_FLOOR, &obs_d, 1)
|
||||
.expect("ema_update_per_step");
|
||||
}
|
||||
stream.synchronize().expect("sync after ema seeding");
|
||||
|
||||
// Verify the EMA producers wrote what we expected (sanity check
|
||||
// before testing the controllers themselves).
|
||||
let isv_after_ema = readback_isv(&dev, &trainer.isv_d);
|
||||
for (slot, expected) in inputs {
|
||||
let got = isv_after_ema[slot];
|
||||
assert!(
|
||||
(got - expected).abs() < 1e-5,
|
||||
"EMA producer should bootstrap slot {slot} to {expected}; got {got}"
|
||||
);
|
||||
}
|
||||
|
||||
// Fire all 7 RL controllers per-step. Each reads its EMA input
|
||||
// and Wiener-blends its output away from the bootstrap value.
|
||||
trainer
|
||||
.launch_rl_controllers_per_step()
|
||||
.expect("launch_rl_controllers_per_step");
|
||||
stream.synchronize().expect("sync after controllers");
|
||||
|
||||
let isv_after = readback_isv(&dev, &trainer.isv_d);
|
||||
|
||||
// Each output slot must have moved off the bootstrap value. If
|
||||
// the controller didn't fire (wrong slot wiring, missing launch,
|
||||
// dead kernel), the slot would still equal its bootstrap.
|
||||
let outputs: [(&str, usize, f32); 7] = [
|
||||
("γ", RL_GAMMA_INDEX, GAMMA_BOOTSTRAP),
|
||||
("τ", RL_TARGET_TAU_INDEX, TAU_BOOTSTRAP),
|
||||
("ε", RL_PPO_CLIP_INDEX, EPS_BOOTSTRAP),
|
||||
("entropy_coef", RL_ENTROPY_COEF_INDEX, COEF_BOOTSTRAP),
|
||||
("n_rollout_steps", RL_N_ROLLOUT_STEPS_INDEX, ROLLOUT_BOOTSTRAP),
|
||||
("per_α", RL_PER_ALPHA_INDEX, PER_ALPHA_BOOTSTRAP),
|
||||
("reward_scale", RL_REWARD_SCALE_INDEX, REWARD_SCALE_BOOTSTRAP),
|
||||
];
|
||||
for (name, slot, bootstrap) in outputs {
|
||||
let got = isv_after[slot];
|
||||
assert!(
|
||||
(got - bootstrap).abs() > 1e-6,
|
||||
"controller for {name} (ISV[{slot}]) should have moved off bootstrap {bootstrap}; got {got} (controller may not have fired or read wrong input slot)"
|
||||
);
|
||||
}
|
||||
|
||||
eprintln!(
|
||||
"G3 OK — all 7 controllers moved their outputs: \
|
||||
γ {} → {}, τ {} → {}, ε {} → {}, coef {} → {}, n_roll {} → {}, per_α {} → {}, scale {} → {}",
|
||||
GAMMA_BOOTSTRAP, isv_after[RL_GAMMA_INDEX],
|
||||
TAU_BOOTSTRAP, isv_after[RL_TARGET_TAU_INDEX],
|
||||
EPS_BOOTSTRAP, isv_after[RL_PPO_CLIP_INDEX],
|
||||
COEF_BOOTSTRAP, isv_after[RL_ENTROPY_COEF_INDEX],
|
||||
ROLLOUT_BOOTSTRAP, isv_after[RL_N_ROLLOUT_STEPS_INDEX],
|
||||
PER_ALPHA_BOOTSTRAP, isv_after[RL_PER_ALPHA_INDEX],
|
||||
REWARD_SCALE_BOOTSTRAP, isv_after[RL_REWARD_SCALE_INDEX],
|
||||
);
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn g4_dqn_target_soft_update_implements_polyak_formula() {
|
||||
@@ -299,19 +74,23 @@ fn g4_dqn_target_soft_update_implements_polyak_formula() {
|
||||
.memcpy_dtoh(&trainer.dqn_head.w_target_d, target_before.as_mut_slice())
|
||||
.expect("dtoh w_target before");
|
||||
|
||||
// τ = ISV[401] bootstrap value = 0.005.
|
||||
let isv = readback_isv(&dev, &trainer.isv_d);
|
||||
let tau = isv[RL_TARGET_TAU_INDEX];
|
||||
// τ = ISV[RL_TARGET_TAU_INDEX] bootstrap value. Sync the trainer's
|
||||
// stream first so the bootstrap-controller writes from
|
||||
// `IntegratedTrainer::new` are visible through the mapped-pinned
|
||||
// ISV host view per `feedback_no_htod_htoh_only_mapped_pinned`.
|
||||
trainer.stream.synchronize().expect("sync trainer stream");
|
||||
let tau = trainer.read_isv_host(RL_TARGET_TAU_INDEX);
|
||||
assert!(
|
||||
(tau - TAU_BOOTSTRAP).abs() < 1e-6,
|
||||
"pre-condition: τ should be R1-bootstrapped to {TAU_BOOTSTRAP}; got {tau}"
|
||||
);
|
||||
|
||||
// Fire the soft update.
|
||||
let isv_d_clone = trainer.isv_d.clone();
|
||||
// Fire the soft update. `soft_update_target` reads τ from
|
||||
// `isv_dev_ptr` (raw `CUdeviceptr`, zero-copy stable pointer).
|
||||
let isv_ptr = trainer.isv_dev_ptr;
|
||||
trainer
|
||||
.dqn_head
|
||||
.soft_update_target(&isv_d_clone)
|
||||
.soft_update_target(&isv_ptr)
|
||||
.expect("soft_update_target");
|
||||
stream.synchronize().expect("sync after soft_update");
|
||||
|
||||
@@ -345,18 +124,6 @@ fn g4_dqn_target_soft_update_implements_polyak_formula() {
|
||||
"soft_update should change at least one target element when w_d != target"
|
||||
);
|
||||
|
||||
// Limit case 1: τ = 0 → target unchanged. Overwrite ISV[401] = 0
|
||||
// via the EMA-update kernel's bootstrap-defer path. (mean_obs == 0
|
||||
// would defer; we instead overwrite the slot by re-firing the
|
||||
// controller with a new input that yields τ ≈ 0 — but that's
|
||||
// brittle. Cleaner: re-firing with the bootstrap zero in ISV[401]
|
||||
// is impossible because R1 already bootstrapped it.)
|
||||
//
|
||||
// Pragmatic approach: just verify the formula holds at the
|
||||
// ACTUAL τ value the kernel sees. The Polyak invariant is the
|
||||
// load-bearing assertion; the τ=0/τ=1 limits add no information
|
||||
// beyond what the formula check already pins.
|
||||
|
||||
eprintln!(
|
||||
"G4 OK — soft_update applied Polyak formula with τ={tau}: \
|
||||
target[0] {} → {} (expected {})",
|
||||
|
||||
@@ -1,145 +0,0 @@
|
||||
//! Phase R7d gate G6: PER buffer wired into `step_with_lobsim`.
|
||||
//!
|
||||
//! Asserts the load-bearing invariants that distinguish a wired PER
|
||||
//! buffer from R7c's dead `src/rl/replay.rs` struct:
|
||||
//!
|
||||
//! 1. `trainer.replay.len()` grows by exactly `b_size` per
|
||||
//! `step_with_lobsim` call (the per-batch push order documented
|
||||
//! in `IntegratedTrainer::push_to_replay`).
|
||||
//! 2. `replay.sample_indices(b_size, α)` returns a vec of length
|
||||
//! `b_size` after the first step (buffer non-empty post-push).
|
||||
//! 3. After N steps with N × b_size ≤ capacity, `replay.len() ==
|
||||
//! N × b_size` (no replacement). After N steps with N × b_size
|
||||
//! > capacity, `replay.len() == capacity` (ring-with-replacement
|
||||
//! capped at capacity).
|
||||
//!
|
||||
//! Per `pearl_tests_must_prove_not_lock_observations`: asserts
|
||||
//! buffer-mechanism invariants (length growth, sample size), NOT
|
||||
//! observed Q values or losses — those vary across runs and lock
|
||||
//! the test against any future change to Q init or PRNG state.
|
||||
//!
|
||||
//! Run with:
|
||||
//! `cargo test -p ml-alpha --test r7d_per_wiring -- --ignored --nocapture`
|
||||
|
||||
use ml_alpha::cfc::snap_features::Mbp10RawInput;
|
||||
use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig};
|
||||
use ml_alpha::trainer::perception::PerceptionTrainerConfig;
|
||||
use ml_backtesting::sim::LobSimCuda;
|
||||
use ml_core::device::MlDevice;
|
||||
|
||||
fn synthetic_window(seq_len: usize, base_mid: f32) -> Vec<Mbp10RawInput> {
|
||||
let mut out = Vec::with_capacity(seq_len);
|
||||
let mut prev_mid = base_mid;
|
||||
let mut ts_ns = 1_000_000_u64;
|
||||
for _ in 0..seq_len {
|
||||
let next_mid = prev_mid + 0.25;
|
||||
let mut bid_px = [0.0_f32; 10];
|
||||
let mut bid_sz = [0.0_f32; 10];
|
||||
let mut ask_px = [0.0_f32; 10];
|
||||
let mut ask_sz = [0.0_f32; 10];
|
||||
for i in 0..10 {
|
||||
bid_px[i] = next_mid - 0.125 - 0.25 * i as f32;
|
||||
ask_px[i] = next_mid + 0.125 + 0.25 * i as f32;
|
||||
bid_sz[i] = 10.0;
|
||||
ask_sz[i] = 10.0;
|
||||
}
|
||||
let prev_ts = ts_ns;
|
||||
ts_ns += 20_000_000;
|
||||
out.push(Mbp10RawInput {
|
||||
bid_px,
|
||||
bid_sz,
|
||||
ask_px,
|
||||
ask_sz,
|
||||
prev_mid,
|
||||
trade_signed_vol: 0.0,
|
||||
trade_count: 0,
|
||||
ts_ns,
|
||||
prev_ts_ns: prev_ts,
|
||||
regime: [0.0; 6],
|
||||
});
|
||||
prev_mid = next_mid;
|
||||
}
|
||||
out
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn r7d_per_buffer_grows_one_per_step_at_b_size_1() {
|
||||
let dev = match MlDevice::cuda(0) {
|
||||
Ok(d) => d,
|
||||
Err(_) => {
|
||||
eprintln!("CUDA 0 not available — skipping r7d_per_wiring");
|
||||
return;
|
||||
}
|
||||
};
|
||||
|
||||
// b_size = 1, small capacity so we can also verify the ring-cap
|
||||
// invariant in the same test (no need for a separate fixture).
|
||||
let per_capacity = 8usize;
|
||||
let cfg = IntegratedTrainerConfig {
|
||||
perception: PerceptionTrainerConfig {
|
||||
seq_len: 4,
|
||||
n_batch: 1,
|
||||
..PerceptionTrainerConfig::default()
|
||||
},
|
||||
per_capacity,
|
||||
..IntegratedTrainerConfig::default()
|
||||
};
|
||||
let mut trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new");
|
||||
let mut sim = LobSimCuda::new(1, &dev).expect("LobSimCuda::new");
|
||||
|
||||
// Invariant 1: buffer starts empty.
|
||||
assert_eq!(
|
||||
trainer.replay.len(),
|
||||
0,
|
||||
"PER buffer must start empty before any step_with_lobsim call"
|
||||
);
|
||||
|
||||
// Drive N=5 steps; assert linear growth from 0 → 5.
|
||||
for step in 1..=5usize {
|
||||
let snapshots = synthetic_window(4, 5500.0 + step as f32 * 0.25);
|
||||
let next_snapshots = synthetic_window(4, 5500.0 + (step + 1) as f32 * 0.25);
|
||||
let _stats = trainer
|
||||
.step_with_lobsim(&snapshots, &next_snapshots, &mut sim)
|
||||
.unwrap_or_else(|e| panic!("step_with_lobsim step {step}: {e:?}"));
|
||||
assert_eq!(
|
||||
trainer.replay.len(),
|
||||
step,
|
||||
"PER buffer must grow by exactly 1 per step at b_size=1 (step {step})"
|
||||
);
|
||||
}
|
||||
|
||||
// Invariant 2: sample at α=0.6 returns batch size 1.
|
||||
let sample = trainer.replay.sample_indices(1, 0.6);
|
||||
assert_eq!(
|
||||
sample.len(),
|
||||
1,
|
||||
"sample_indices(1, 0.6) on a non-empty buffer must return 1 index"
|
||||
);
|
||||
assert!(
|
||||
sample[0] < trainer.replay.len(),
|
||||
"sampled index must be in [0, replay.len()) — got {} for len {}",
|
||||
sample[0],
|
||||
trainer.replay.len()
|
||||
);
|
||||
|
||||
// Invariant 3: drive past capacity, buffer caps at `per_capacity`.
|
||||
// Already at len=5; drive 10 more (total 15 transitions pushed,
|
||||
// capacity=8 → buffer ends at exactly 8).
|
||||
for step in 6..=15usize {
|
||||
let snapshots = synthetic_window(4, 5500.0 + step as f32 * 0.25);
|
||||
let next_snapshots = synthetic_window(4, 5500.0 + (step + 1) as f32 * 0.25);
|
||||
trainer
|
||||
.step_with_lobsim(&snapshots, &next_snapshots, &mut sim)
|
||||
.unwrap_or_else(|e| panic!("step_with_lobsim step {step}: {e:?}"));
|
||||
}
|
||||
assert_eq!(
|
||||
trainer.replay.len(),
|
||||
per_capacity,
|
||||
"PER buffer must cap at per_capacity = {per_capacity} (ring-with-replacement)"
|
||||
);
|
||||
|
||||
eprintln!(
|
||||
"R7d G6 OK — buffer grew 0→5→{per_capacity} across 15 push cycles; sample returns expected size"
|
||||
);
|
||||
}
|
||||
@@ -152,17 +152,18 @@ fn default_pyramid_ctx(stream: &Arc<CudaStream>, b_size: usize) -> Result<Pyrami
|
||||
})
|
||||
}
|
||||
|
||||
/// Overwrite a single ISV slot host-side then push the whole isv_d
|
||||
/// to the device. Cheap because the trainer's `isv_d` is small
|
||||
/// (~500 floats).
|
||||
/// Overwrite a single ISV slot via the mapped-pinned buffer's volatile
|
||||
/// per-record write. The GPU sees the new value after the next stream
|
||||
/// sync barrier — no explicit HtoD copy needed
|
||||
/// (per `feedback_no_htod_htoh_only_mapped_pinned`).
|
||||
fn set_isv_slot(
|
||||
trainer: &mut IntegratedTrainer,
|
||||
stream: &Arc<CudaStream>,
|
||||
_stream: &Arc<CudaStream>,
|
||||
slot: usize,
|
||||
value: f32,
|
||||
) -> Result<()> {
|
||||
trainer.isv_host[slot] = value;
|
||||
write_slice_f32_d_pub(stream, &trainer.isv_host, &mut trainer.isv_d)
|
||||
trainer.isv_mapped.write_record(slot, value);
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
|
||||
@@ -1143,6 +1143,7 @@ impl LobSimCuda {
|
||||
.arg(&n)
|
||||
.launch(cfg)?;
|
||||
}
|
||||
self.stream.synchronize()?;
|
||||
Ok(())
|
||||
}
|
||||
|
||||
@@ -1433,43 +1434,6 @@ impl ml_alpha::rl::reward::RlLobBackend for LobSimCuda {
|
||||
ask_sz_src,
|
||||
)
|
||||
}
|
||||
|
||||
fn raw_ptrs(&self) -> ml_alpha::rl::reward::LobSimRawPtrs {
|
||||
ml_alpha::rl::reward::LobSimRawPtrs {
|
||||
bid_px: self.bid_px_d.raw_ptr(),
|
||||
bid_sz: self.bid_sz_d.raw_ptr(),
|
||||
ask_px: self.ask_px_d.raw_ptr(),
|
||||
ask_sz: self.ask_sz_d.raw_ptr(),
|
||||
books: self.books_d.raw_ptr(),
|
||||
prev_mid: self.prev_mid_d.raw_ptr(),
|
||||
atr_mid_ema: self.atr_mid_ema_d.raw_ptr(),
|
||||
snapshots_skipped: self.snapshots_skipped_d.raw_ptr(),
|
||||
min_reasonable_px: self.min_reasonable_px_d.raw_ptr(),
|
||||
max_reasonable_px: self.max_reasonable_px_d.raw_ptr(),
|
||||
book_update_fn: self.book_update_fn.cu_function(),
|
||||
market_targets: self.market_targets_d.raw_ptr(),
|
||||
pos: self.pos_d.raw_ptr(),
|
||||
cost_per_lot_per_side: self.cost_per_lot_per_side_d.raw_ptr(),
|
||||
total_fees_per_b: self.total_fees_per_b_d.raw_ptr(),
|
||||
submit_market_fn: self.submit_market_fn.cu_function(),
|
||||
open_trade_state: self.open_trade_state_d.raw_ptr(),
|
||||
trade_log: self.trade_log_d.raw_ptr(),
|
||||
trade_log_head: self.trade_log_head_d.raw_ptr(),
|
||||
trail_hwm: self.trail_hwm_d.raw_ptr(),
|
||||
zero_vwap_at_open: self.zero_vwap_at_open_d.raw_ptr(),
|
||||
saturated_vwap_at_open: self.saturated_vwap_at_open_d.raw_ptr(),
|
||||
defensive_exit_clamp: self.defensive_exit_clamp_d.raw_ptr(),
|
||||
conv_signed_ema: self.conv_signed_ema_d.raw_ptr(),
|
||||
diag_hold_hist: self.diag_hold_hist_d.raw_ptr(),
|
||||
diag_outcome_n: self.diag_outcome_n_d.raw_ptr(),
|
||||
diag_outcome_sum_pnl: self.diag_outcome_sum_pnl_d.raw_ptr(),
|
||||
diag_outcome_n_wins: self.diag_outcome_n_wins_d.raw_ptr(),
|
||||
pnl_track_fn: self.pnl_track_fn.cu_function(),
|
||||
n_backtests: self.n_backtests as i32,
|
||||
pos_bytes: std::mem::size_of::<crate::lob::PosFlat>() as i32,
|
||||
trade_log_cap: crate::lob::TRADE_LOG_CAP as i32,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Re-open the impl block for `LobSimCuda` so subsequent methods (if any
|
||||
|
||||
@@ -1,333 +0,0 @@
|
||||
# Alpha-RL Performance + Checkpoint + Walk-Forward Implementation Plan
|
||||
|
||||
> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.
|
||||
|
||||
**Goal:** 2-3× training throughput via TF32, crash recovery via checkpointing, OOS validation via 3-fold walk-forward on H100.
|
||||
|
||||
**Architecture:** Enable TF32 Tensor Core math on all cuBLAS handles (2 sites). Add checkpoint/resume serialization for model weights + Adam state + ISV to the training binary. Validate with 3-fold walk-forward at 25k steps/fold on H100.
|
||||
|
||||
**Tech Stack:** Rust, CUDA cuBLAS, cudarc, mapped-pinned memory, Argo Workflows
|
||||
|
||||
---
|
||||
|
||||
### Task 1: Enable TF32 on DQN cuBLAS handle
|
||||
|
||||
**Files:**
|
||||
- Modify: `crates/ml-alpha/src/rl/dqn.rs:314` (after `cublasSetWorkspace_v2`)
|
||||
|
||||
- [ ] **Step 1: Add TF32 math mode after workspace setup**
|
||||
|
||||
In `crates/ml-alpha/src/rl/dqn.rs`, after line 313 (the closing of the `cublasSetWorkspace_v2` block), add inside the same `unsafe` block:
|
||||
|
||||
```rust
|
||||
cudarc::cublas::sys::cublasSetMathMode(
|
||||
*cublas.handle(),
|
||||
cudarc::cublas::sys::cublasMath_t::CUBLAS_TF32_TENSOR_OP_MATH,
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow::anyhow!("DqnHead: cublasSetMathMode TF32: {e:?}"))?;
|
||||
```
|
||||
|
||||
- [ ] **Step 2: Build check**
|
||||
|
||||
Run: `SQLX_OFFLINE=true cargo check -p ml-alpha`
|
||||
Expected: compiles cleanly
|
||||
|
||||
- [ ] **Step 3: Commit**
|
||||
|
||||
```bash
|
||||
git add crates/ml-alpha/src/rl/dqn.rs
|
||||
git commit -m "perf(cuda): enable TF32 Tensor Core math on DQN cuBLAS handle"
|
||||
```
|
||||
|
||||
### Task 2: Enable TF32 on Mamba2 cuBLAS handle
|
||||
|
||||
**Files:**
|
||||
- Modify: `crates/ml-alpha/src/mamba2_block.rs:549` (after `cublasSetWorkspace_v2`)
|
||||
|
||||
- [ ] **Step 1: Add TF32 math mode after workspace setup**
|
||||
|
||||
In `crates/ml-alpha/src/mamba2_block.rs`, after line 549 (the closing of the `cublasSetWorkspace_v2` block), add inside the same `unsafe` block:
|
||||
|
||||
```rust
|
||||
cudarc::cublas::sys::cublasSetMathMode(
|
||||
*cublas.handle(),
|
||||
cudarc::cublas::sys::cublasMath_t::CUBLAS_TF32_TENSOR_OP_MATH,
|
||||
)
|
||||
.result()
|
||||
.map_err(|e| anyhow!("Mamba2Block: cublasSetMathMode TF32: {e:?}"))?;
|
||||
```
|
||||
|
||||
- [ ] **Step 2: Build check**
|
||||
|
||||
Run: `SQLX_OFFLINE=true cargo check -p ml-alpha`
|
||||
Expected: compiles cleanly
|
||||
|
||||
- [ ] **Step 3: Local smoke test (RTX 3050 Ti, sm_86 supports TF32)**
|
||||
|
||||
Run a 100-step local smoke and verify l_q is within 1% of baseline:
|
||||
```bash
|
||||
FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml-alpha --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -5
|
||||
```
|
||||
Expected: smoke passes, no NaN
|
||||
|
||||
- [ ] **Step 4: Commit**
|
||||
|
||||
```bash
|
||||
git add crates/ml-alpha/src/mamba2_block.rs
|
||||
git commit -m "perf(cuda): enable TF32 Tensor Core math on Mamba2 cuBLAS handle"
|
||||
```
|
||||
|
||||
### Task 3: Add AdamW save/load methods
|
||||
|
||||
**Files:**
|
||||
- Modify: `crates/ml-alpha/src/trainer/optim.rs`
|
||||
|
||||
- [ ] **Step 1: Add save method to AdamW**
|
||||
|
||||
Add after the existing `impl AdamW` block (or at the end of the impl):
|
||||
|
||||
```rust
|
||||
/// Serialize Adam state (m, v, step_count, hyperparams) to writer.
|
||||
/// Reads device buffers via mapped-pinned DtoH (one-shot, not hot path).
|
||||
pub fn save(&self, w: &mut impl std::io::Write) -> Result<()> {
|
||||
let m_host = self._stream.clone().read_sync(&self.m)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save m: {e}"))?;
|
||||
let v_host = self._stream.clone().read_sync(&self.v)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW save v: {e}"))?;
|
||||
let n = m_host.len() as u64;
|
||||
w.write_all(&n.to_le_bytes())?;
|
||||
w.write_all(&self.step_count_host.to_le_bytes())?;
|
||||
w.write_all(&self.lr.to_le_bytes())?;
|
||||
w.write_all(&self.beta1.to_le_bytes())?;
|
||||
w.write_all(&self.beta2.to_le_bytes())?;
|
||||
w.write_all(&self.eps.to_le_bytes())?;
|
||||
w.write_all(&self.wd.to_le_bytes())?;
|
||||
let m_bytes: &[u8] = bytemuck::cast_slice(&m_host);
|
||||
w.write_all(m_bytes)?;
|
||||
let v_bytes: &[u8] = bytemuck::cast_slice(&v_host);
|
||||
w.write_all(v_bytes)?;
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Restore Adam state from reader. Sizes must match.
|
||||
pub fn load(&mut self, r: &mut impl std::io::Read) -> Result<()> {
|
||||
let mut buf8 = [0u8; 8];
|
||||
let mut buf4 = [0u8; 4];
|
||||
r.read_exact(&mut buf8)?;
|
||||
let n = u64::from_le_bytes(buf8) as usize;
|
||||
anyhow::ensure!(n == self.m.len(), "AdamW load: m size mismatch {n} vs {}", self.m.len());
|
||||
r.read_exact(&mut buf4)?;
|
||||
self.step_count_host = i32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?; self.lr = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?; self.beta1 = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?; self.beta2 = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?; self.eps = f32::from_le_bytes(buf4);
|
||||
r.read_exact(&mut buf4)?; self.wd = f32::from_le_bytes(buf4);
|
||||
let mut m_host = vec![0f32; n];
|
||||
r.read_exact(bytemuck::cast_slice_mut(&mut m_host))?;
|
||||
self._stream.clone().write_sync(&m_host, &mut self.m)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load m: {e}"))?;
|
||||
let mut v_host = vec![0f32; n];
|
||||
r.read_exact(bytemuck::cast_slice_mut(&mut v_host))?;
|
||||
self._stream.clone().write_sync(&v_host, &mut self.v)
|
||||
.map_err(|e| anyhow::anyhow!("AdamW load v: {e}"))?;
|
||||
Ok(())
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **Step 2: Verify bytemuck is in dependencies**
|
||||
|
||||
Run: `grep bytemuck crates/ml-alpha/Cargo.toml`
|
||||
If missing, add `bytemuck = { version = "1", features = ["derive"] }`.
|
||||
|
||||
- [ ] **Step 3: Build check**
|
||||
|
||||
Run: `SQLX_OFFLINE=true cargo check -p ml-alpha`
|
||||
|
||||
- [ ] **Step 4: Commit**
|
||||
|
||||
```bash
|
||||
git add crates/ml-alpha/src/trainer/optim.rs crates/ml-alpha/Cargo.toml
|
||||
git commit -m "feat(rl): AdamW save/load for checkpoint persistence"
|
||||
```
|
||||
|
||||
### Task 4: Add IntegratedTrainer checkpoint save/load
|
||||
|
||||
**Files:**
|
||||
- Modify: `crates/ml-alpha/src/trainer/integrated.rs`
|
||||
|
||||
- [ ] **Step 1: Add save_checkpoint method**
|
||||
|
||||
Add a `save_checkpoint` method to `IntegratedTrainer` that serializes:
|
||||
- DQN weights (online + target): `dqn_head.w_d`, `dqn_head.b_d`, `dqn_head.w_target_d`, `dqn_head.b_target_d`
|
||||
- Policy/V head weights
|
||||
- ISV bus (585 f32 from mapped-pinned `isv_mapped.host_ptr`)
|
||||
- All 20 AdamW states via `adam.save()`
|
||||
- Step counter
|
||||
- Encoder via `perception.save_checkpoint()`
|
||||
|
||||
Format: sequential sections with `[name_len: u16][name: bytes][data_len: u64][data: bytes]`.
|
||||
File header: `[magic: u32 = 0x464F5843][version: u32 = 1][step: u64]`.
|
||||
|
||||
Write to `<out_dir>/checkpoint-<step>.bin`. Keep last 2, delete older.
|
||||
|
||||
- [ ] **Step 2: Add load_checkpoint method**
|
||||
|
||||
Inverse of save: read header, validate magic/version, restore each section by name lookup. Skip unknown sections for forward compatibility.
|
||||
|
||||
- [ ] **Step 3: Build check**
|
||||
|
||||
Run: `SQLX_OFFLINE=true cargo check -p ml-alpha`
|
||||
|
||||
- [ ] **Step 4: Commit**
|
||||
|
||||
```bash
|
||||
git add crates/ml-alpha/src/trainer/integrated.rs
|
||||
git commit -m "feat(rl): IntegratedTrainer checkpoint save/load"
|
||||
```
|
||||
|
||||
### Task 5: Wire checkpoint into training loop
|
||||
|
||||
**Files:**
|
||||
- Modify: `crates/ml-alpha/examples/alpha_rl_train.rs`
|
||||
|
||||
- [ ] **Step 1: Add CLI flags**
|
||||
|
||||
Add to the CLI struct (clap):
|
||||
```rust
|
||||
/// Resume from checkpoint file path. Restores model weights, Adam
|
||||
/// state, ISV, and step counter. Training continues from saved step.
|
||||
#[arg(long)]
|
||||
resume_from: Option<PathBuf>,
|
||||
|
||||
/// Checkpoint interval in steps (default: 5000). Set to 0 to disable.
|
||||
#[arg(long, default_value = "5000")]
|
||||
checkpoint_every: usize,
|
||||
```
|
||||
|
||||
- [ ] **Step 2: Add resume logic before training loop**
|
||||
|
||||
After trainer init, before `for step in 0..cli.n_steps`:
|
||||
```rust
|
||||
let start_step = if let Some(ref ckpt_path) = cli.resume_from {
|
||||
trainer.load_checkpoint(ckpt_path)
|
||||
.with_context(|| format!("resume from {}", ckpt_path.display()))?
|
||||
} else {
|
||||
0
|
||||
};
|
||||
```
|
||||
|
||||
Change loop to `for step in start_step..cli.n_steps`.
|
||||
|
||||
- [ ] **Step 3: Add checkpoint save inside training loop**
|
||||
|
||||
After the per-step JSONL write, before the next iteration:
|
||||
```rust
|
||||
if cli.checkpoint_every > 0 && step > 0 && step % cli.checkpoint_every == 0 {
|
||||
let ckpt_path = cli.out.join(format!("checkpoint-{step}.bin"));
|
||||
trainer.save_checkpoint(&ckpt_path, step as u64)
|
||||
.with_context(|| format!("checkpoint at step {step}"))?;
|
||||
eprintln!("checkpoint saved: {}", ckpt_path.display());
|
||||
// Rolling: keep last 2, delete older
|
||||
let old_step = step.saturating_sub(cli.checkpoint_every * 2);
|
||||
if old_step > 0 {
|
||||
let old_path = cli.out.join(format!("checkpoint-{old_step}.bin"));
|
||||
let _ = std::fs::remove_file(&old_path);
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
- [ ] **Step 4: Build check**
|
||||
|
||||
Run: `SQLX_OFFLINE=true cargo check -p ml-alpha`
|
||||
|
||||
- [ ] **Step 5: Commit**
|
||||
|
||||
```bash
|
||||
git add crates/ml-alpha/examples/alpha_rl_train.rs
|
||||
git commit -m "feat(rl): wire checkpoint save/resume into training loop"
|
||||
```
|
||||
|
||||
### Task 6: Local validation smoke
|
||||
|
||||
**Files:** None (test only)
|
||||
|
||||
- [ ] **Step 1: Run 200-step local smoke with TF32**
|
||||
|
||||
```bash
|
||||
SQLX_OFFLINE=true cargo build --release -p ml-alpha --example alpha_rl_train
|
||||
# Run 200 steps with checkpoint at 100
|
||||
./target/release/examples/alpha_rl_train \
|
||||
--n-backtests 16 --n-steps 200 --checkpoint-every 100 \
|
||||
--out /tmp/smoke-tf32 \
|
||||
--mbp10-data-dir test_data/futures-baseline \
|
||||
--trades-data-dir test_data/futures-baseline 2>&1 | tail -5
|
||||
```
|
||||
|
||||
Expected: completes without NaN, checkpoint files at `/tmp/smoke-tf32/checkpoint-100.bin`
|
||||
|
||||
- [ ] **Step 2: Test resume**
|
||||
|
||||
```bash
|
||||
./target/release/examples/alpha_rl_train \
|
||||
--n-backtests 16 --n-steps 200 --checkpoint-every 100 \
|
||||
--resume-from /tmp/smoke-tf32/checkpoint-100.bin \
|
||||
--out /tmp/smoke-tf32-resume \
|
||||
--mbp10-data-dir test_data/futures-baseline \
|
||||
--trades-data-dir test_data/futures-baseline 2>&1 | tail -5
|
||||
```
|
||||
|
||||
Expected: starts from step 100, runs to 200, no crash
|
||||
|
||||
- [ ] **Step 3: Commit all together**
|
||||
|
||||
```bash
|
||||
git add -A
|
||||
git commit -m "test: validate TF32 + checkpoint save/resume local smoke"
|
||||
```
|
||||
|
||||
### Task 7: Deploy walk-forward on H100
|
||||
|
||||
**Files:** None (deployment only)
|
||||
|
||||
- [ ] **Step 1: Push branch**
|
||||
|
||||
```bash
|
||||
git push origin ml-alpha-phase-a
|
||||
```
|
||||
|
||||
- [ ] **Step 2: Submit 3-fold walk-forward on H100**
|
||||
|
||||
```bash
|
||||
./scripts/argo-train.sh \
|
||||
--model alpha-rl \
|
||||
--branch ml-alpha-phase-a \
|
||||
--gpu-pool ci-training-h100 \
|
||||
--folds 3
|
||||
```
|
||||
|
||||
If `argo-train.sh` doesn't support `--folds` with step override, submit directly:
|
||||
```bash
|
||||
argo submit --from=wftmpl/alpha-rl -n foxhunt \
|
||||
-p git-branch=ml-alpha-phase-a \
|
||||
-p n-steps=25000 \
|
||||
-p n-backtests=1024 \
|
||||
-p per-capacity=65536 \
|
||||
-p gpu-pool=ci-training-h100
|
||||
```
|
||||
Run 3 times with fold-idx 0, 1, 2.
|
||||
|
||||
- [ ] **Step 3: Monitor with log-only (no extra pods on GPU node)**
|
||||
|
||||
```bash
|
||||
kubectl logs <pod-name> -n foxhunt --tail=1 -f | grep "step.*wr="
|
||||
```
|
||||
|
||||
- [ ] **Step 4: Validate success criteria**
|
||||
|
||||
All 3 folds must show:
|
||||
- wr > 0.55
|
||||
- No fold with wr < 0.50
|
||||
- Entropy stable (no collapse)
|
||||
- hold% between 30-70%
|
||||
@@ -1,152 +0,0 @@
|
||||
# Alpha-RL Performance + Checkpoint + Walk-Forward
|
||||
|
||||
**Date**: 2026-05-27
|
||||
**Status**: Approved
|
||||
**Scope**: TF32 matmuls, checkpoint/resume, controller fusion, walk-forward validation
|
||||
|
||||
## Context
|
||||
|
||||
The alpha-rl pipeline achieves wr=0.567 at b=1024 but takes 4.5h for 100k steps on L40S at 6.2 sps — and the model converges by 5-10k steps. The last 90k steps are wasted. The L40S node OOM'd at 58k from monitoring pod overhead, losing all state (zero checkpoint infrastructure). All cuBLAS SGEMMs force `CUBLAS_COMPUTE_32F`, bypassing Tensor Cores entirely on both L40S and H100.
|
||||
|
||||
## Goals
|
||||
|
||||
1. **2-3× throughput** via TF32 Tensor Core acceleration on cuBLAS SGEMMs
|
||||
2. **Crash recovery** via checkpoint/resume every 5k steps
|
||||
3. **~30% launch overhead reduction** via controller kernel fusion
|
||||
4. **OOS validation** via 3-fold walk-forward at 25k steps/fold on H100
|
||||
|
||||
Target wall clock: 3 folds × 25k steps at ~15 sps on H100 = ~83 min total (vs 4.5h+ single-fold L40S that never finished).
|
||||
|
||||
## P0: TF32 Matmuls
|
||||
|
||||
### What
|
||||
|
||||
Enable TF32 Tensor Core math on all cuBLAS handles. TF32 uses 19-bit mantissa (vs FP32's 23-bit) with identical range. Precision loss is well within RL gradient noise. PyTorch default since 1.7.
|
||||
|
||||
### Where
|
||||
|
||||
Two cuBLAS handle creation sites:
|
||||
|
||||
1. `crates/ml-alpha/src/rl/dqn.rs` — after `cublasCreate_v2()`, add `cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH)`. Affects DQN forward (online + target) and backward (grad_h + grad_w) = 4 SGEMMs/step.
|
||||
|
||||
2. `crates/ml-alpha/src/mamba2_block.rs` — same pattern after handle creation. Affects W_in, W_a, W_b, W_out projection GEMMs = 4+ SGEMMs/step.
|
||||
|
||||
### Impact
|
||||
|
||||
8+ SGEMMs per step go from FP32 scalar to TF32 Tensor Core. On L40S (Ada Lovelace): ~2× throughput. On H100 (Hopper): ~3× throughput. Combined with H100's 2× memory bandwidth: expect 12-20 sps at b=1024 (vs 6.2 current).
|
||||
|
||||
### Verification
|
||||
|
||||
Run 1k-step local smoke on RTX 3050 Ti (sm_86, supports TF32). Compare l_q at step 1000 with and without TF32 — should be within 1%.
|
||||
|
||||
## P1: Checkpoint/Resume
|
||||
|
||||
### What
|
||||
|
||||
Serialize training state every 5k steps to PVC. Resume from last checkpoint on restart.
|
||||
|
||||
### State to persist
|
||||
|
||||
| Component | Size | Source |
|
||||
|-----------|------|--------|
|
||||
| DQN weights (online) | W[128×231] + b[231] = ~120KB | `dqn_head.w_d`, `dqn_head.b_d` |
|
||||
| DQN weights (target) | Same = ~120KB | `dqn_head.w_target_d`, `dqn_head.b_target_d` |
|
||||
| Policy head weights | ~120KB | `policy_head.w_d`, `policy_head.b_d` |
|
||||
| V head weights | ~2KB | `v_head.w_d`, `v_head.b_d` |
|
||||
| Encoder weights | ~2-10MB (Mamba2 + CfC) | `encoder.params_d()` |
|
||||
| Adam state (m, v per param) | 2× model size = ~20MB | Per-group m_d, v_d |
|
||||
| ISV bus | 585 × f32 = 2.3KB | `isv_dev_ptr` |
|
||||
| PER buffer | priorities + tree = ~1MB | `gpu_replay.priorities_d` |
|
||||
| Step counter + RNG | ~100B | `step`, xorshift state |
|
||||
|
||||
Total: ~50MB per checkpoint.
|
||||
|
||||
### Format
|
||||
|
||||
Flat binary: `[magic: u32][version: u32][step: u64][n_sections: u32][sections...]` where each section is `[name_len: u16][name: bytes][data_len: u64][data: bytes]`. No serde, no JSON — raw device→mapped-pinned→file.
|
||||
|
||||
### Path
|
||||
|
||||
`/feature-cache/alpha-rl-runs/<sha>/fold<N>/checkpoint-<step>.bin`
|
||||
|
||||
Keep last 2 checkpoints (rolling). Delete older ones to save PVC space.
|
||||
|
||||
### Resume
|
||||
|
||||
CLI flag `--resume-from <path>`. Loads checkpoint, restores all device buffers, continues from saved step. The `alpha_rl_train.rs` binary checks for the flag before the training loop.
|
||||
|
||||
### Verification
|
||||
|
||||
Save at step 1000, kill, resume, compare l_q/wr at step 2000 with an uninterrupted run. Should be identical (deterministic with scoped_init_seed per pearl).
|
||||
|
||||
## P2: Controller Kernel Fusion
|
||||
|
||||
### What
|
||||
|
||||
Fuse 10+ single-thread ISV controller kernels into one `rl_fused_all_controllers` mega-kernel. Currently each controller is a separate `raw_launch(grid=1, block=1)` — the launch overhead (~5μs each) dominates the actual compute (~1μs each).
|
||||
|
||||
### Controllers to fuse
|
||||
|
||||
- `rl_gamma_controller`
|
||||
- `rl_target_tau_controller`
|
||||
- `rl_ppo_clip_controller` (legacy but still launched)
|
||||
- `rl_entropy_coef_controller`
|
||||
- `rl_per_alpha_controller`
|
||||
- `rl_reward_scale_controller`
|
||||
- `rl_rollout_steps_controller`
|
||||
- `rl_lr_controller` (5 heads)
|
||||
|
||||
All share the same pattern: read ISV input slot, Wiener-α blend, Schulman bounded step, write ISV output slot. The existing `rl_fused_controllers.cu` already handles a subset — extend to cover all.
|
||||
|
||||
### Impact
|
||||
|
||||
~10 kernel launches → 1. Saves ~50μs/step. At 6 sps that's ~0.3ms saved per 160ms step = ~0.2% improvement. Small but free.
|
||||
|
||||
### Verification
|
||||
|
||||
Before/after nsys trace: total GPU kernel launches should drop by ~10 per step.
|
||||
|
||||
## P3: Walk-Forward Validation
|
||||
|
||||
### What
|
||||
|
||||
3-fold temporal walk-forward on H100 with 25k steps per fold. Each fold trains on window [i] and evaluates on window [i+1]. The argo template already supports `--folds 3` which fans out via DAG.
|
||||
|
||||
### Config
|
||||
|
||||
```bash
|
||||
argo submit --from=wftmpl/alpha-rl -n foxhunt \
|
||||
-p git-branch=ml-alpha-phase-a \
|
||||
-p n-steps=25000 \
|
||||
-p n-backtests=1024 \
|
||||
-p per-capacity=65536 \
|
||||
-p gpu-pool=ci-training-h100
|
||||
# --folds 3 via argo-train.sh
|
||||
```
|
||||
|
||||
### Success criteria
|
||||
|
||||
- wr > 0.55 on ALL 3 folds (not just the average)
|
||||
- No fold with wr < 0.50
|
||||
- Entropy stable (no collapse on any fold)
|
||||
- hold% between 30-70% on all folds
|
||||
|
||||
### Wall clock estimate
|
||||
|
||||
3 folds × 25k steps. With TF32 on H100 at ~15 sps: 25k/15 = 28 min/fold. Sequential: 84 min. Parallel (3 H100s): 28 min.
|
||||
|
||||
## Implementation Order
|
||||
|
||||
1. **P0: TF32** — 2 one-line changes, local smoke, deploy
|
||||
2. **P2: Controller fusion** — extend existing fused kernel, local smoke
|
||||
3. **P1: Checkpoint** — new infrastructure, needs careful testing
|
||||
4. **P3: Walk-forward** — deploy after P0+P2 are validated
|
||||
|
||||
P0 and P2 can be implemented in parallel. P1 is independent. P3 runs after all code changes are merged.
|
||||
|
||||
## Anti-patterns to avoid
|
||||
|
||||
- No FP16 matmuls (loss scaling complexity, RL numerical sensitivity)
|
||||
- No multi-GPU (NCCL overhead not justified at 25k steps)
|
||||
- No changes to model architecture (validated at wr=0.567)
|
||||
- No monitoring pods on GPU node (caused OOM — use log-only monitoring)
|
||||
@@ -1,129 +0,0 @@
|
||||
# Mega-Graph CUDA Pipeline
|
||||
|
||||
**Date**: 2026-05-27
|
||||
**Status**: Draft
|
||||
**Goal**: Capture the entire per-step training pipeline into one CUDA graph launch, eliminating ~150 individual kernel launches and ~140ms of host-side overhead. Target: 50+ sps (from 6.4).
|
||||
|
||||
## Problem
|
||||
|
||||
nsys profiling shows GPU utilization at 8.5% (13ms kernel time per 156ms step). The host spends 143ms between kernel launches on: Rust arg assembly, Result unwrapping, conditional logic, 5 graph launches + ~20 individual raw_launch calls. The CUDA APIs themselves are fast (<1ms total); the overhead is Rust code between them.
|
||||
|
||||
## Current Per-Step Pipeline
|
||||
|
||||
```
|
||||
HOST GPU (13ms total)
|
||||
───────────────────────────── ────────────────────────────
|
||||
build 3 gather args (50μs) → sample_and_gather (13ms)
|
||||
gather_next (0.04ms)
|
||||
gather_current (0.04ms)
|
||||
graph_launch prefill (50μs) → GRAPH A: encoder + Q/V/π + actions (~2ms)
|
||||
build 4 gate args (100μs) → conf_gate + frd_gate + action_ent + log_pi (0.1ms)
|
||||
lobsim.apply_snapshot (20μs) → book_update (0.01ms)
|
||||
graph_launch postfill (50μs) → GRAPH A2: risk + trail + actions_to_market (~1ms)
|
||||
lobsim.step_fill (20μs) → fill + pnl_track (0.05ms)
|
||||
graph_launch reward (50μs) → GRAPH C: reward + scale + clamp + context (~2ms)
|
||||
build 2 PER args (20μs) → per_push_ring + per_push_flush (0.05ms)
|
||||
perception.step_batched (10μs)→ PERCEPTION GRAPH: encoder train (~2ms)
|
||||
graph_launch training (50μs) → GRAPH D: Q/π/V backward + Adam (~3ms)
|
||||
build 4 target args (50μs) → soft_update + q_divergence (0.1ms)
|
||||
DiagFrame memcpy (100μs) → (async, background thread)
|
||||
───────────────────────────
|
||||
Total host: ~570μs arg work
|
||||
But Rust overhead: ~155ms (!)
|
||||
```
|
||||
|
||||
The 570μs of pure arg work can't explain 155ms. The remaining ~154ms is:
|
||||
- `step_synthetic()` → reads ISV from mapped-pinned → launches LR controller → builds training graph args → launches GRAPH D → reads ISV for LR plateau
|
||||
- `step_with_lobsim_reward_and_train()` → builds reward graph args → conditional K-loop → event-based cross-stream sync → PER sample on train_stream
|
||||
|
||||
Each of these involves deep Rust function call chains (integrated.rs is 8400 lines) with Result unwrapping, debug_asserts, conditional branches, and method dispatch through trait objects.
|
||||
|
||||
## Approaches
|
||||
|
||||
### Approach A: Mega-Graph (Recommended)
|
||||
|
||||
Capture ALL kernels from a single step into ONE graph during warmup. On subsequent steps, launch the mega-graph with a single `cuGraphLaunch` call. Zero host work between kernels.
|
||||
|
||||
**Requirements**:
|
||||
1. ALL kernel arguments use stable device pointers (pre-allocated at init)
|
||||
2. No host-side conditional logic during the captured section
|
||||
3. No memory allocation/deallocation during capture
|
||||
4. K-loop fixed at K=1 (current config)
|
||||
5. cuBLAS workspace pre-allocated (already done)
|
||||
|
||||
**What changes per step** (must be handled):
|
||||
- `sample_and_gather` uses per-step PRNG state → device-side, graph-safe
|
||||
- Perception's `ts_ns` value → write to mapped-pinned staging, kernel reads from stable pointer
|
||||
- ISV bus → modified by controller kernels in-place, graph-safe
|
||||
- LR controller → reads host-side loss values → **BLOCKER**: needs mapped-pinned reads
|
||||
|
||||
**LR controller blocker**: The host reads `last_q_loss`, `last_pi_loss`, `last_v_loss` from mapped-pinned memory between the reward graph and the training graph. These feed the per-head LR scaling. Fix: move LR controller to a GPU kernel that reads from mapped-pinned device pointers directly. The host doesn't need to see the LR values — the Adam kernel reads LR from ISV.
|
||||
|
||||
**Lobsim blocker**: `apply_snapshot_from_device` and `step_fill_from_market_targets` are lobsim methods with internal state. These need to be either: (a) converted to raw_launch with cached pointers, or (b) kept outside the mega-graph as a "lobsim mini-graph."
|
||||
|
||||
**Estimated result**: 1 graph launch per step. Host does: graph launch (50μs) + DiagFrame memcpy (100μs) = 150μs/step → **~6600 sps** (GPU-bound at 13ms/step → ~77 sps with GPU saturation).
|
||||
|
||||
### Approach B: Coalesce Existing Graphs
|
||||
|
||||
Merge the 4 existing graphs + loose kernels into 2 larger graphs: "env_graph" (A + gates + A2 + fill + C) and "train_graph" (PER + perception + D + target).
|
||||
|
||||
**Simpler**: doesn't require solving the LR controller blocker or lobsim abstraction. Just expand the existing capture boundaries.
|
||||
|
||||
**Host work**: 2 graph launches + cross-stream event + DiagFrame = ~300μs. But the Rust overhead between the two launches is still significant (K-loop, LR reads, etc.).
|
||||
|
||||
**Estimated result**: ~20-30 sps (2× improvement, not 10×).
|
||||
|
||||
### Approach C: Async Host Pipeline
|
||||
|
||||
Keep the current graph structure but pipeline: host launches step N+1's gathers while step N's training graph runs. Double-buffer all state.
|
||||
|
||||
**Complex**: requires double-buffering 50+ device buffers. Weight updates from step N must be visible to step N+1's forward. Correctness is hard to verify.
|
||||
|
||||
**Estimated result**: ~12 sps (2× from pipelining).
|
||||
|
||||
## Recommendation: Approach A (Mega-Graph)
|
||||
|
||||
The only approach that reaches 50+ sps. The blockers are solvable:
|
||||
|
||||
1. **LR controller → GPU kernel**: already ISV-driven; just move the host-side `AdamW.lr` mutation to a kernel that writes the LR value into the Adam kernel's arg buffer. ~30 lines.
|
||||
|
||||
2. **Lobsim → raw_launch**: `apply_snapshot_from_device` is 2 DtoD copies + 1 kernel. `step_fill_from_market_targets` is 2 kernels. Convert both to raw_launch with cached pointers. ~100 lines.
|
||||
|
||||
3. **K-loop fixed at K=1**: already the case in production config. The mega-graph captures one iteration.
|
||||
|
||||
4. **Cross-stream PER**: PER sample runs on train_stream, but at K=1 it's just 1 sample per step. Can be moved to the main stream (the separate stream was for K>1 pipelining).
|
||||
|
||||
## Implementation Phases
|
||||
|
||||
### Phase 1: Convert lobsim methods to raw_launch (prerequisite)
|
||||
- `apply_snapshot_from_device` → 2 DtoD + 1 kernel via raw_launch
|
||||
- `step_fill_from_market_targets` → 2 kernels via raw_launch
|
||||
- Cache all lobsim internal pointers in LobPtrs
|
||||
|
||||
### Phase 2: Move LR controller to GPU kernel
|
||||
- New `rl_lr_from_mapped_pinned.cu` kernel reads loss from mapped-pinned device pointers
|
||||
- Writes per-head LR to ISV slots
|
||||
- Adam kernel reads LR from ISV (already supported)
|
||||
|
||||
### Phase 3: Move PER sample to main stream
|
||||
- At K=1, PER sample + priority_update + tree_rebuild run once per step
|
||||
- Move from train_stream to self.stream
|
||||
- Eliminate cross-stream events
|
||||
|
||||
### Phase 4: Mega-graph capture
|
||||
- Warmup step 0: eager execution (all kernels launch individually)
|
||||
- Warmup step 1: `begin_capture` → entire pipeline → `end_capture`
|
||||
- Step 2+: single `cuGraphLaunch` per step
|
||||
- Host per-step: ts_ns write to mapped-pinned + graph launch + DiagFrame
|
||||
|
||||
### Phase 5: Validation
|
||||
- nsys: confirm 1 cuGraphLaunch per step, GPU utilization >80%
|
||||
- Functional: wr/pnl match baseline within 1%
|
||||
- Smoke: 1k steps local, 5k steps L40S
|
||||
|
||||
## Anti-patterns to avoid
|
||||
|
||||
- No `cuGraphExecUpdate` (overkill — our topology is fixed)
|
||||
- No conditional graph nodes (CUDA 12.4 feature, complex, fragile)
|
||||
- No multi-stream graphs (correctness nightmare)
|
||||
- No host reads inside the captured section (use mapped-pinned reads OUTSIDE or defer)
|
||||
@@ -158,7 +158,6 @@ spec:
|
||||
|
||||
# Auto-detect GPU compute capability for all crates' build.rs
|
||||
export CUDA_COMPUTE_CAP=$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader | head -1 | tr -d '.')
|
||||
export FOXHUNT_CUDA_ARCH="sm_${CUDA_COMPUTE_CAP}"
|
||||
echo "=== GPU arch: sm_${CUDA_COMPUTE_CAP} ==="
|
||||
|
||||
# Clear stale build artifacts when arch changes
|
||||
|
||||
2
vendor/cudarc/src/driver/safe/core.rs
vendored
2
vendor/cudarc/src/driver/safe/core.rs
vendored
@@ -85,7 +85,7 @@ impl CudaContext {
|
||||
cu_device,
|
||||
cu_ctx,
|
||||
ordinal,
|
||||
has_async_alloc: AtomicBool::new(true),
|
||||
has_async_alloc: AtomicBool::new(false),
|
||||
num_streams: AtomicUsize::new(0),
|
||||
event_tracking: AtomicBool::new(true),
|
||||
error_state: AtomicU32::new(0),
|
||||
|
||||
Reference in New Issue
Block a user