diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 823f037f8..2b4ba06e7 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -97,6 +97,8 @@ const KERNELS: &[&str] = &[ "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 + "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) ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu index 521c0c66d..44f893a66 100644 --- a/crates/ml-alpha/cuda/rl_confidence_gate.cu +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -30,6 +30,7 @@ #define RL_CONF_GATE_SIGMA_NORM_INDEX 514 #define RL_CONF_GATE_FIRED_COUNT_INDEX 515 #define RL_GATE_WARMUP_STEPS_INDEX 524 +#define RL_STEP_COUNTER_ISV_INDEX 548 extern "C" __global__ void rl_confidence_gate( int* __restrict__ actions, // [B] IN/OUT @@ -38,12 +39,12 @@ extern "C" __global__ void rl_confidence_gate( const unsigned char* __restrict__ pos_state, // [B × pos_bytes] float* __restrict__ isv, int b_size, - int pos_bytes, - int current_step + int pos_bytes ) { const int b = blockIdx.x; if (b >= b_size) 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; diff --git a/crates/ml-alpha/cuda/rl_frd_gate.cu b/crates/ml-alpha/cuda/rl_frd_gate.cu index 39680999e..bb18c8b3e 100644 --- a/crates/ml-alpha/cuda/rl_frd_gate.cu +++ b/crates/ml-alpha/cuda/rl_frd_gate.cu @@ -34,6 +34,7 @@ #define RL_FRD_GATE_THR_SHORT_INDEX 517 #define RL_FRD_GATE_FIRED_COUNT_INDEX 518 #define RL_GATE_WARMUP_STEPS_INDEX 524 +#define RL_STEP_COUNTER_ISV_INDEX 548 extern "C" __global__ void rl_frd_gate( int* __restrict__ actions, // [B] IN/OUT @@ -41,12 +42,12 @@ extern "C" __global__ void rl_frd_gate( const unsigned char* __restrict__ pos_state, // [B × pos_bytes] float* __restrict__ isv, int b_size, - int pos_bytes, - int current_step + int pos_bytes ) { const int b = blockIdx.x; if (b >= b_size) 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; diff --git a/crates/ml-alpha/cuda/rl_fused_controllers.cu b/crates/ml-alpha/cuda/rl_fused_controllers.cu new file mode 100644 index 000000000..9bbb46a44 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_fused_controllers.cu @@ -0,0 +1,464 @@ +// rl_fused_controllers.cu — single-kernel fusion of 10 RL ISV controllers. +// +// Eliminates 9 kernel-launch overheads (~40-80μs/step) by running all +// controllers sequentially in one thread. Each controller's logic is +// copied verbatim from its standalone .cu file. The standalone files +// are retained for documentation and individual testing. +// +// Controllers fused (in execution order): +// 1. rl_gamma_controller ISV[400] ← ISV[input_slots[0]] +// 2. rl_target_tau_controller ISV[401] ← ISV[input_slots[1]] +// 3. rl_ppo_clip_controller ISV[402] ← ISV[input_slots[2]] +// 4. rl_entropy_coef_controller ISV[403] ← ISV[input_slots[3]] +// 5. rl_rollout_steps_controller ISV[404] ← ISV[input_slots[4]] +// 6. rl_per_alpha_controller ISV[405] ← ISV[input_slots[5]] +// 7. rl_reward_scale_controller ISV[406] ← ISV[input_slots[6]] +// 8. rl_ppo_ratio_clamp_controller ISV[440] ← ISV[402] (reads ε just written) +// 9. rl_gate_threshold_controller ISV[512,516,517] ← dones[B] +// 10. rl_q_distill_lambda_controller ISV[486] ← ISV[488] +// +// Per `feedback_no_atomicadd`: single thread, no atomics. +// Per `feedback_cpu_is_read_only`: all state in ISV, no host roundtrip. + +#include + +// ═══════════════════════════════════════════════════════════════════════ +// ISV slot indices — mirrored from individual controller headers. +// ═══════════════════════════════════════════════════════════════════════ + +// --- 1. Gamma controller --- +#define RL_GAMMA_INDEX 400 +#define GAMMA_MIN 0.995f +#define GAMMA_MAX 0.999f + +// --- 2. Target tau controller --- +#define RL_TARGET_TAU_INDEX 401 +#define TAU_MIN 0.001f +#define TAU_MAX 0.05f +#define RL_TAU_BOOTSTRAP_INDEX 473 +#define RL_DIV_TARGET_INDEX 457 +#define DIV_NOISE_FLOOR_FRAC 0.01f + +// --- 3. PPO clip controller --- +#define RL_PPO_CLIP_INDEX 402 +#define EPS_MIN 0.05f +#define EPS_MAX 0.5f +#define RL_EPS_BOOTSTRAP_INDEX 474 +#define RL_KL_TARGET_INDEX 454 +#define KL_NOISE_FLOOR_FRAC 0.01f + +// --- 4. Entropy coef controller --- +#define RL_ENTROPY_COEF_INDEX 403 +#define N_ACTIONS 11 +#define COEF_MIN 0.0f +#define COEF_MAX 0.05f +#define RL_ENTROPY_TARGET_FRAC_INDEX 458 + +// --- 5. Rollout steps controller --- +#define RL_N_ROLLOUT_STEPS_INDEX 404 +#define ROLLOUT_MIN 256.0f +#define ROLLOUT_MAX 8192.0f +#define RL_ROLLOUT_BOOTSTRAP_INDEX 475 +#define RL_ADV_VAR_RATIO_TARGET_INDEX 449 +#define ADV_VAR_RATIO_NOISE_FLOOR_FRAC 0.01f + +// --- 6. PER alpha controller --- +#define RL_PER_ALPHA_INDEX 405 +#define PER_ALPHA_MIN 0.3f +#define PER_ALPHA_MAX 1.0f +#define RL_KURT_GAUSSIAN_INDEX 471 +#define RL_KURT_NOISE_FLOOR_INDEX 472 +#define RL_KURT_LIFT_SCALE_INDEX 459 + +// --- 7. Reward scale controller --- +#define RL_REWARD_SCALE_INDEX 406 +#define RL_REWARD_SCALE_MIN_INDEX 492 +#define REWARD_SCALE_MAX 1e3f +#define RL_REWARD_SCALE_BOOTSTRAP_INDEX 476 +#define EPS_PNL 1e-3f + +// --- 8. PPO ratio clamp controller --- +#define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 +#define RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX 477 +#define RL_PPO_CLAMP_MARGIN_INDEX 460 +#define PPO_RATIO_CLAMP_MIN_OUT 2.0f +#define PPO_RATIO_CLAMP_MAX_OUT 1000.0f + +// --- 9. Gate threshold controller --- +#define RL_CONF_GATE_THRESHOLD_INDEX 512 +#define RL_FRD_GATE_THR_LONG_INDEX 516 +#define RL_FRD_GATE_THR_SHORT_INDEX 517 +#define RL_GATE_WARMUP_STEPS_INDEX 524 +#define RL_GATE_DONES_EMA_INDEX 525 +#define RL_GATE_DONES_TARGET_INDEX 526 +#define RL_GATE_CONF_MIN_INDEX 527 +#define RL_GATE_CONF_MAX_INDEX 528 +#define RL_GATE_FRD_MIN_INDEX 529 +#define RL_GATE_FRD_MAX_INDEX 530 +#define RL_GATE_ADJUST_RATE_INDEX 531 + +// --- 10. Q distill lambda controller --- +#define RL_Q_DISTILL_LAMBDA_INDEX 486 +#define RL_Q_DISTILL_KL_EMA_INDEX 488 +#define RL_Q_DISTILL_KL_TARGET_INDEX 491 +#define MIN_LAMBDA 0.05f +#define MAX_LAMBDA 1.0f +#define KL_TOLERANCE 3.0f +#define LAMBDA_RAMP_RATE 1.2f +#define LAMBDA_DECAY_RATE 0.998f + +// --- Shared Schulman params --- +#define RL_SCHULMAN_TOLERANCE_INDEX 468 +#define RL_SCHULMAN_ADJUST_RATE_INDEX 469 + +// --- Shared Wiener --- +#define WIENER_ALPHA_FLOOR 0.4f + +// --- Step counter (device-resident) --- +#define RL_STEP_COUNTER_ISV_INDEX 548 + + +// ═══════════════════════════════════════════════════════════════════════ +// Fused kernel: all 10 controllers in one launch. +// +// input_slots[7]: ISV indices for the 7 base controllers' EMA inputs. +// [0] = RL_MEAN_TRADE_DURATION_EMA_INDEX (417) — gamma +// [1] = RL_Q_DIVERGENCE_EMA_INDEX (418) — target_tau +// [2] = RL_KL_PI_EMA_INDEX (419) — ppo_clip +// [3] = RL_ENTROPY_OBSERVED_EMA_INDEX (420) — entropy_coef +// [4] = RL_ADVANTAGE_VAR_RATIO_EMA_INDEX (421) — rollout_steps +// [5] = RL_TD_KURTOSIS_EMA_INDEX (422) — per_alpha +// [6] = RL_MEAN_ABS_PNL_EMA_INDEX (423) — reward_scale +// ═══════════════════════════════════════════════════════════════════════ +extern "C" __global__ void rl_fused_controllers( + float* __restrict__ isv, + const float* __restrict__ dones, // [b_size] for gate threshold + float alpha, // Wiener-α for the 7 base + ratio clamp + int b_size, // batch size for gate threshold dones scan + const int* __restrict__ input_slots // [7] ISV indices for base controllers +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX]; + + // ═══════════════════════════════════════════════════════════════════ + // 1. Gamma controller → ISV[400] + // ═══════════════════════════════════════════════════════════════════ + { + const float gamma_prev = isv[RL_GAMMA_INDEX]; + const float mean_trade_duration_events = isv[input_slots[0]]; + const float d = fmaxf(mean_trade_duration_events, 1.0f); + float gamma_target = powf(0.5f, 1.0f / d); + gamma_target = fmaxf(GAMMA_MIN, fminf(gamma_target, GAMMA_MAX)); + + if (gamma_prev == 0.0f) { + isv[RL_GAMMA_INDEX] = gamma_target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float gamma_new = (1.0f - a) * gamma_prev + a * gamma_target; + gamma_new = fmaxf(GAMMA_MIN, fminf(gamma_new, GAMMA_MAX)); + isv[RL_GAMMA_INDEX] = gamma_new; + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 2. Target tau controller → ISV[401] + // ═══════════════════════════════════════════════════════════════════ + { + const float tau_prev = isv[RL_TARGET_TAU_INDEX]; + if (tau_prev == 0.0f) { + isv[RL_TARGET_TAU_INDEX] = isv[RL_TAU_BOOTSTRAP_INDEX]; + } else { + const float q_divergence_norm = isv[input_slots[1]]; + if (q_divergence_norm != 0.0f) { + const float div_target = isv[RL_DIV_TARGET_INDEX]; + const float div_noise_floor = div_target * DIV_NOISE_FLOOR_FRAC; + if (q_divergence_norm >= div_noise_floor) { + float ratio; + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + if (q_divergence_norm > div_target * tolerance) { + ratio = adjust_rate; + } else if (q_divergence_norm < div_target / tolerance) { + ratio = 1.0f / adjust_rate; + } else { + ratio = 1.0f; + } + float tau_target = tau_prev * ratio; + tau_target = fmaxf(TAU_MIN, fminf(tau_target, TAU_MAX)); + + if (tau_prev == isv[RL_TAU_BOOTSTRAP_INDEX]) { + isv[RL_TARGET_TAU_INDEX] = tau_target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float tau_new = (1.0f - a) * tau_prev + a * tau_target; + tau_new = fmaxf(TAU_MIN, fminf(tau_new, TAU_MAX)); + isv[RL_TARGET_TAU_INDEX] = tau_new; + } + } + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 3. PPO clip controller → ISV[402] + // ═══════════════════════════════════════════════════════════════════ + { + const float eps_prev = isv[RL_PPO_CLIP_INDEX]; + if (eps_prev == 0.0f) { + isv[RL_PPO_CLIP_INDEX] = isv[RL_EPS_BOOTSTRAP_INDEX]; + } else { + const float kl_ema = isv[input_slots[2]]; + if (kl_ema != 0.0f) { + const float kl_target = isv[RL_KL_TARGET_INDEX]; + const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC; + if (kl_ema >= kl_noise_floor) { + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + float ratio; + if (kl_ema > kl_target * tolerance) { + ratio = 1.0f / adjust_rate; + } else if (kl_ema < kl_target / tolerance) { + ratio = adjust_rate; + } else { + ratio = 1.0f; + } + float eps_target = eps_prev * ratio; + eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX)); + + if (eps_prev == isv[RL_EPS_BOOTSTRAP_INDEX]) { + isv[RL_PPO_CLIP_INDEX] = eps_target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float eps_new = (1.0f - a) * eps_prev + a * eps_target; + eps_new = fmaxf(EPS_MIN, fminf(eps_new, EPS_MAX)); + isv[RL_PPO_CLIP_INDEX] = eps_new; + } + } + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 4. Entropy coef controller → ISV[403] + // ═══════════════════════════════════════════════════════════════════ + { + const float coef_prev = isv[RL_ENTROPY_COEF_INDEX]; + const float entropy_observed_ema = isv[input_slots[3]]; + const float h_max = logf((float)N_ACTIONS); + const float target_frac = isv[RL_ENTROPY_TARGET_FRAC_INDEX]; + const float h_target = target_frac * h_max; + const float deficit = fmaxf(0.0f, h_target - entropy_observed_ema); + float coef_target = (deficit / h_max) * COEF_MAX; + coef_target = fmaxf(COEF_MIN, fminf(coef_target, COEF_MAX)); + + if (coef_prev == 0.0f) { + isv[RL_ENTROPY_COEF_INDEX] = coef_target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float coef_new = (1.0f - a) * coef_prev + a * coef_target; + coef_new = fmaxf(COEF_MIN, fminf(coef_new, COEF_MAX)); + isv[RL_ENTROPY_COEF_INDEX] = coef_new; + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 5. Rollout steps controller → ISV[404] + // ═══════════════════════════════════════════════════════════════════ + { + const float prev = isv[RL_N_ROLLOUT_STEPS_INDEX]; + if (prev == 0.0f) { + isv[RL_N_ROLLOUT_STEPS_INDEX] = isv[RL_ROLLOUT_BOOTSTRAP_INDEX]; + } else { + const float advantage_var_over_abs_mean = isv[input_slots[4]]; + if (advantage_var_over_abs_mean != 0.0f) { + const float adv_var_target = isv[RL_ADV_VAR_RATIO_TARGET_INDEX]; + const float adv_var_noise_floor = + adv_var_target * ADV_VAR_RATIO_NOISE_FLOOR_FRAC; + if (advantage_var_over_abs_mean >= adv_var_noise_floor) { + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + float scale; + if (advantage_var_over_abs_mean > adv_var_target * tolerance) { + scale = adjust_rate; + } else if (advantage_var_over_abs_mean < adv_var_target / tolerance) { + scale = 1.0f / adjust_rate; + } else { + scale = 1.0f; + } + float target = prev * scale; + target = fmaxf(ROLLOUT_MIN, fminf(target, ROLLOUT_MAX)); + + if (prev == isv[RL_ROLLOUT_BOOTSTRAP_INDEX]) { + isv[RL_N_ROLLOUT_STEPS_INDEX] = target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + out = fmaxf(ROLLOUT_MIN, fminf(out, ROLLOUT_MAX)); + isv[RL_N_ROLLOUT_STEPS_INDEX] = out; + } + } + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 6. PER alpha controller → ISV[405] + // ═══════════════════════════════════════════════════════════════════ + { + const float prev = isv[RL_PER_ALPHA_INDEX]; + const float td_kurtosis_ema = isv[input_slots[5]]; + + if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < isv[RL_KURT_NOISE_FLOOR_INDEX]) { + if (prev != 0.0f) goto per_alpha_done; + } + + { + const float kurt_excess = fmaxf(0.0f, td_kurtosis_ema - isv[RL_KURT_GAUSSIAN_INDEX]); + const float kurt_lift_scale = isv[RL_KURT_LIFT_SCALE_INDEX]; + float target = 0.4f + 0.2f * (kurt_excess / kurt_lift_scale); + target = fmaxf(PER_ALPHA_MIN, fminf(target, PER_ALPHA_MAX)); + + if (prev == 0.0f) { + isv[RL_PER_ALPHA_INDEX] = target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + out = fmaxf(PER_ALPHA_MIN, fminf(out, PER_ALPHA_MAX)); + isv[RL_PER_ALPHA_INDEX] = out; + } + } + per_alpha_done:; + } + + // ═══════════════════════════════════════════════════════════════════ + // 7. Reward scale controller → ISV[406] + // ═══════════════════════════════════════════════════════════════════ + { + const float prev = isv[RL_REWARD_SCALE_INDEX]; + if (prev == 0.0f) { + isv[RL_REWARD_SCALE_INDEX] = isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]; + } else { + const float mean_abs_pnl_ema = isv[input_slots[6]]; + if (mean_abs_pnl_ema != 0.0f) { + const float denom = fmaxf(mean_abs_pnl_ema, EPS_PNL); + float target = 1.0f / denom; + const float scale_min = isv[RL_REWARD_SCALE_MIN_INDEX]; + target = fmaxf(scale_min, fminf(target, REWARD_SCALE_MAX)); + + if (prev == isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]) { + isv[RL_REWARD_SCALE_INDEX] = target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + + // Per-step ±2% movement clamp. + const float max_move = prev * 1.02f; + const float min_move = prev * 0.98f; + out = fmaxf(min_move, fminf(out, max_move)); + + out = fmaxf(scale_min, fminf(out, REWARD_SCALE_MAX)); + isv[RL_REWARD_SCALE_INDEX] = out; + } + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 8. PPO ratio clamp controller → ISV[440] + // Reads ε from ISV[402] which was just written above. + // ═══════════════════════════════════════════════════════════════════ + { + const float prev = isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]; + if (prev == 0.0f) { + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX]; + } else { + const float eps = isv[RL_PPO_CLIP_INDEX]; + const float clamp_margin = isv[RL_PPO_CLAMP_MARGIN_INDEX]; + float target = (1.0f + eps) * clamp_margin; + target = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, + fminf(target, PPO_RATIO_CLAMP_MAX_OUT)); + + if (prev == isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX]) { + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = target; + } else { + const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + out = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, + fminf(out, PPO_RATIO_CLAMP_MAX_OUT)); + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = out; + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 9. Gate threshold controller → ISV[512,516,517] + // Scans dones[B] to compute done_rate EMA → ISV[525]. + // Adjusts conf + FRD thresholds via Schulman bounded step. + // ═══════════════════════════════════════════════════════════════════ + { + const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX]; + if (current_step >= warmup) { + float done_count = 0.0f; + for (int b = 0; b < b_size; ++b) { + if (dones[b] > 0.5f) done_count += 1.0f; + } + const float done_rate = done_count / (float)b_size; + const float gate_alpha = 0.01f; + const float prev_ema = isv[RL_GATE_DONES_EMA_INDEX]; + const float dones_ema = (prev_ema == 0.0f && done_rate > 0.0f) + ? done_rate + : (1.0f - gate_alpha) * prev_ema + gate_alpha * done_rate; + isv[RL_GATE_DONES_EMA_INDEX] = dones_ema; + + const float target = isv[RL_GATE_DONES_TARGET_INDEX]; + const float adjust = isv[RL_GATE_ADJUST_RATE_INDEX]; + + float conf_thr = isv[RL_CONF_GATE_THRESHOLD_INDEX]; + float frd_long = isv[RL_FRD_GATE_THR_LONG_INDEX]; + float frd_short = isv[RL_FRD_GATE_THR_SHORT_INDEX]; + + if (dones_ema < target * 0.2f) { + conf_thr *= adjust; + frd_long *= adjust; + frd_short *= adjust; + } else if (dones_ema > target * 5.0f) { + conf_thr /= adjust; + frd_long /= adjust; + frd_short /= adjust; + } + + const float conf_min = isv[RL_GATE_CONF_MIN_INDEX]; + const float conf_max = isv[RL_GATE_CONF_MAX_INDEX]; + const float frd_min = isv[RL_GATE_FRD_MIN_INDEX]; + const float frd_max = isv[RL_GATE_FRD_MAX_INDEX]; + + isv[RL_CONF_GATE_THRESHOLD_INDEX] = fmaxf(conf_min, fminf(conf_thr, conf_max)); + isv[RL_FRD_GATE_THR_LONG_INDEX] = fmaxf(frd_min, fminf(frd_long, frd_max)); + isv[RL_FRD_GATE_THR_SHORT_INDEX] = fmaxf(frd_min, fminf(frd_short, frd_max)); + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 10. Q→π distill lambda controller → ISV[486] + // Asymmetric bounded step on KL_EMA vs target. + // ═══════════════════════════════════════════════════════════════════ + { + const float kl_observed = isv[RL_Q_DISTILL_KL_EMA_INDEX]; + const float kl_target = isv[RL_Q_DISTILL_KL_TARGET_INDEX]; + float lambda = isv[RL_Q_DISTILL_LAMBDA_INDEX]; + + if (kl_observed > 0.0f && kl_target > 0.0f && lambda > 0.0f) { + const float upper = kl_target * KL_TOLERANCE; + const float lower = kl_target / KL_TOLERANCE; + + if (kl_observed > upper) { + lambda = fminf(MAX_LAMBDA, lambda * LAMBDA_RAMP_RATE); + } else if (kl_observed < lower) { + lambda = fmaxf(MIN_LAMBDA, lambda * LAMBDA_DECAY_RATE); + } + + isv[RL_Q_DISTILL_LAMBDA_INDEX] = lambda; + } + } +} diff --git a/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu index 440be8d0f..76b4c5a36 100644 --- a/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu +++ b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu @@ -22,15 +22,16 @@ #define RL_GATE_FRD_MIN_INDEX 529 #define RL_GATE_FRD_MAX_INDEX 530 #define RL_GATE_ADJUST_RATE_INDEX 531 +#define RL_STEP_COUNTER_ISV_INDEX 548 extern "C" __global__ void rl_gate_threshold_controller( float* __restrict__ isv, const float* __restrict__ dones, // [B] - int b_size, - int current_step + int b_size ) { if (threadIdx.x != 0 || blockIdx.x != 0) 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; diff --git a/crates/ml-alpha/cuda/rl_increment_step.cu b/crates/ml-alpha/cuda/rl_increment_step.cu new file mode 100644 index 000000000..02c68b939 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_increment_step.cu @@ -0,0 +1,20 @@ +// rl_increment_step.cu — bump the device-resident step counter. +// +// Single thread (grid=1, block=1), runs once per step at the very start +// of the step_with_lobsim pipeline, BEFORE any kernel that reads +// current_step. Kernels read `(int)isv[RL_STEP_COUNTER_ISV_INDEX]` +// instead of receiving `int current_step` as a scalar argument. This +// makes ALL kernel argument lists constant across calls, enabling CUDA +// Graph capture (scalar args can't change between replays). +// +// Per `feedback_no_atomicadd`: single thread, no atomics. +// Per `feedback_cpu_is_read_only`: pure device-side increment. + +#define RL_STEP_COUNTER_ISV_INDEX 548 + +extern "C" __global__ void rl_increment_step( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + isv[RL_STEP_COUNTER_ISV_INDEX] += 1.0f; +} diff --git a/crates/ml-alpha/cuda/rl_trade_context_update.cu b/crates/ml-alpha/cuda/rl_trade_context_update.cu index 8c880eafe..b89978580 100644 --- a/crates/ml-alpha/cuda/rl_trade_context_update.cu +++ b/crates/ml-alpha/cuda/rl_trade_context_update.cu @@ -19,6 +19,7 @@ #define MAX_UNITS 4 #define RL_MEAN_ABS_PNL_EMA_INDEX 423 #define TRADE_CONTEXT_DIM 4 +#define RL_STEP_COUNTER_ISV_INDEX 548 extern "C" __global__ void rl_trade_context_update( float* __restrict__ trade_context, // [B × TRADE_CONTEXT_DIM] OUT @@ -32,12 +33,13 @@ extern "C" __global__ void rl_trade_context_update( const float* __restrict__ ask_px, // [BOOK_LEVELS] const float* __restrict__ isv, int b_size, - int pos_bytes, - int current_step + int pos_bytes ) { const int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= b_size) return; + const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX]; + const int base = b * MAX_UNITS; const int out_base = b * TRADE_CONTEXT_DIM; diff --git a/crates/ml-alpha/cuda/rl_unit_state_update.cu b/crates/ml-alpha/cuda/rl_unit_state_update.cu index d7a025e26..2f19bd818 100644 --- a/crates/ml-alpha/cuda/rl_unit_state_update.cu +++ b/crates/ml-alpha/cuda/rl_unit_state_update.cu @@ -27,6 +27,7 @@ #define RL_TRAIL_K_INIT_INDEX 496 #define RL_MEAN_ABS_PNL_EMA_INDEX 423 #define RL_PYRAMID_ADD_COUNT_INDEX 507 +#define RL_STEP_COUNTER_ISV_INDEX 548 // PosFlat layout (matches crates/ml-backtesting/src/lob/mod.rs::PosFlat): // offset 0: position_lots: i32 @@ -46,12 +47,13 @@ extern "C" __global__ void rl_unit_state_update( unsigned char* __restrict__ unit_active, // [B * MAX_UNITS] int* __restrict__ pyramid_units_count, // [B] int b_size, - int pos_bytes, - int current_step + int pos_bytes ) { const int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= b_size) return; + const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX]; + const int curr_pos = *(const int*)(pos_state + b * pos_bytes + 0); const float vwap = *(const float*)(pos_state + b * pos_bytes + 4); const int prev_pos = prev_pos_lots[b]; diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 4e11cff6d..2896384b9 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -51,8 +51,9 @@ //! //! | 543-545 | 3 IQN head slots (N_TAU + ensemble α + LR) | IqnHead + ensemble action selector | //! | 547 | Noisy linear σ_init (NoisyNet exploration) | NoisyLinear layers | +//! | 548 | Device-resident step counter (graph-safe) | rl_increment_step | //! -//! Total: 148 slots; `RL_SLOTS_END = 548`. +//! Total: 149 slots; `RL_SLOTS_END = 549`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -990,6 +991,13 @@ pub const RL_ASYM_WIN_THRESHOLD_INDEX: usize = 546; /// action selection. pub const RL_NOISY_SIGMA_INIT_INDEX: usize = 547; +/// Device-resident step counter. Incremented by `rl_increment_step` +/// kernel each step (single thread, ISV[548] += 1.0). Kernels read +/// this instead of a scalar `current_step` argument, enabling CUDA +/// Graph capture (scalar args can't change between replays). Reset to +/// 0.0 at ISV allocation (alloc_zeros); first step sets it to 1.0. +pub const RL_STEP_COUNTER_ISV_INDEX: usize = 548; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 548; +pub const RL_SLOTS_END: usize = 549; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index b7c8a5fd5..be81a4d75 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -209,6 +209,12 @@ const RL_RECENT_OUTCOME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_recent_outcome_update.cubin")); const RL_GATE_THRESHOLD_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_gate_threshold_controller.cubin")); +// Fused kernel: 10 RL ISV controllers in one launch — saves 9 kernel +// launch overheads (~40-80μs/step). Replaces individual launches of +// gamma, tau, ppo_clip, entropy_coef, rollout_steps, per_alpha, +// reward_scale, ppo_ratio_clamp, gate_threshold, q_distill_lambda. +const RL_FUSED_CONTROLLERS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_fused_controllers.cubin")); const RL_REWARD_SHAPING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_shaping.cubin")); const RL_ASYMMETRIC_TRAIL_DECAY_CUBIN: &[u8] = @@ -225,6 +231,11 @@ const RL_TRAIL_STOP_CHECK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_stop_check.cubin")); const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/actions_to_market_targets.cubin")); +// Device-resident step counter bump — ISV[548] += 1.0 each step. +// Prerequisite for CUDA Graph capture: removes scalar current_step +// from all downstream kernel argument lists. +const RL_INCREMENT_STEP_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_increment_step.cubin")); // Element-wise in-place add: dst[i] += src[i]. Used by noisy // exploration to fold NoisyLinear output into ensemble_q_d before @@ -492,7 +503,10 @@ pub struct IntegratedTrainer { _rl_ensemble_action_value_module: Arc, rl_ensemble_action_value_fn: CudaFunction, // λ_distill adaptive controller (rljzl followup 2026-05-24). + // Retained for bootstrap / testing; per-step launch fused into + // rl_fused_controllers. _rl_q_distill_lambda_controller_module: Arc, + #[allow(dead_code)] rl_q_distill_lambda_controller_fn: CudaFunction, // SP20 P1+P5 (audit-wiring fix for dead a7/a8). _rl_unit_state_update_module: Arc, @@ -510,7 +524,16 @@ pub struct IntegratedTrainer { _rl_recent_outcome_update_module: Arc, rl_recent_outcome_update_fn: CudaFunction, _rl_gate_threshold_controller_module: Arc, + // Retained for testing; per-step launch fused into + // rl_fused_controllers. + #[allow(dead_code)] rl_gate_threshold_controller_fn: CudaFunction, + // Fused kernel: 10 RL ISV controllers in one launch. + _rl_fused_controllers_module: Arc, + rl_fused_controllers_fn: CudaFunction, + /// Device buffer holding the 7 ISV input-slot indices for the fused + /// controller kernel. Allocated once at init, never mutated. + fused_ctrl_input_slots_d: CudaSlice, _rl_reward_shaping_module: Arc, rl_reward_shaping_fn: CudaFunction, _rl_asymmetric_trail_decay_module: Arc, @@ -523,6 +546,9 @@ pub struct IntegratedTrainer { rl_trade_context_update_fn: CudaFunction, _rl_multires_features_update_module: Arc, rl_multires_features_update_fn: CudaFunction, + // Device-resident step counter (ISV[548] += 1.0 per step). + _rl_increment_step_module: Arc, + rl_increment_step_fn: CudaFunction, pub outcome_ema_d: CudaSlice, pub trade_context_d: CudaSlice, pub multires_output_d: CudaSlice, @@ -1103,6 +1129,27 @@ impl IntegratedTrainer { let rl_gate_threshold_controller_fn = rl_gate_threshold_controller_module .load_function("rl_gate_threshold_controller") .context("load rl_gate_threshold_controller")?; + let rl_fused_controllers_module = ctx + .load_cubin(RL_FUSED_CONTROLLERS_CUBIN.to_vec()) + .context("load rl_fused_controllers cubin")?; + let rl_fused_controllers_fn = rl_fused_controllers_module + .load_function("rl_fused_controllers") + .context("load rl_fused_controllers")?; + // Device buffer for the 7 ISV input-slot indices consumed by + // the fused controller kernel. Allocated once, never mutated. + let fused_ctrl_input_slots_h: [i32; 7] = [ + crate::rl::isv_slots::RL_MEAN_TRADE_DURATION_EMA_INDEX as i32, + crate::rl::isv_slots::RL_Q_DIVERGENCE_EMA_INDEX as i32, + crate::rl::isv_slots::RL_KL_PI_EMA_INDEX as i32, + crate::rl::isv_slots::RL_ENTROPY_OBSERVED_EMA_INDEX as i32, + crate::rl::isv_slots::RL_ADVANTAGE_VAR_RATIO_EMA_INDEX as i32, + crate::rl::isv_slots::RL_TD_KURTOSIS_EMA_INDEX as i32, + crate::rl::isv_slots::RL_MEAN_ABS_PNL_EMA_INDEX as i32, + ]; + let mut fused_ctrl_input_slots_d = stream + .alloc_zeros::(7) + .context("alloc fused_ctrl_input_slots_d")?; + write_slice_i32_d_pub(&stream, &fused_ctrl_input_slots_h, &mut fused_ctrl_input_slots_d)?; let rl_reward_shaping_module = ctx .load_cubin(RL_REWARD_SHAPING_CUBIN.to_vec()) .context("load rl_reward_shaping cubin")?; @@ -1139,6 +1186,12 @@ impl IntegratedTrainer { let rl_multires_features_update_fn = rl_multires_features_update_module .load_function("rl_multires_features_update") .context("load rl_multires_features_update")?; + let rl_increment_step_module = ctx + .load_cubin(RL_INCREMENT_STEP_CUBIN.to_vec()) + .context("load rl_increment_step cubin")?; + let rl_increment_step_fn = rl_increment_step_module + .load_function("rl_increment_step") + .context("load rl_increment_step")?; let actions_to_market_targets_module = ctx .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) .context("load actions_to_market_targets cubin")?; @@ -1624,6 +1677,9 @@ impl IntegratedTrainer { rl_recent_outcome_update_fn, _rl_gate_threshold_controller_module: rl_gate_threshold_controller_module, rl_gate_threshold_controller_fn, + _rl_fused_controllers_module: rl_fused_controllers_module, + rl_fused_controllers_fn, + fused_ctrl_input_slots_d, _rl_reward_shaping_module: rl_reward_shaping_module, rl_reward_shaping_fn, _rl_asymmetric_trail_decay_module: rl_asymmetric_trail_decay_module, @@ -1636,6 +1692,8 @@ impl IntegratedTrainer { rl_trade_context_update_fn, _rl_multires_features_update_module: rl_multires_features_update_module, rl_multires_features_update_fn, + _rl_increment_step_module: rl_increment_step_module, + rl_increment_step_fn, outcome_ema_d, trade_context_d, multires_output_d, @@ -2097,74 +2155,48 @@ impl IntegratedTrainer { Ok(()) } - /// Phase R5: fire all 7 RL adaptive controllers in sequence, - /// each reading its EMA input from the corresponding ISV slot - /// populated by the Phase R3 `ema_update_*` kernels. + /// Fire all 10 RL ISV controllers in a single fused kernel launch. /// - /// Slot pairing (output ← input): - /// ISV[400] γ ← ISV[417] MEAN_TRADE_DURATION_EMA - /// ISV[401] τ ← ISV[418] Q_DIVERGENCE_EMA - /// ISV[402] ε ← ISV[419] KL_PI_EMA - /// ISV[403] entropy_coef ← ISV[420] ENTROPY_OBSERVED_EMA - /// ISV[404] n_rollout_steps← ISV[421] ADVANTAGE_VAR_RATIO_EMA - /// ISV[405] per_α ← ISV[422] TD_KURTOSIS_EMA - /// ISV[406] reward_scale ← ISV[423] MEAN_ABS_PNL_EMA + /// Replaces 10 individual kernel launches (gamma, tau, ppo_clip, + /// entropy_coef, rollout_steps, per_alpha, reward_scale, + /// ppo_ratio_clamp, gate_threshold, q_distill_lambda) with one + /// launch of `rl_fused_controllers`, saving ~40-80us/step of + /// kernel-launch overhead. /// - /// R6's `step_with_lobsim` will call this AFTER the EMA producers - /// have updated the input slots. Idempotent on a single trainer - /// step but should fire exactly once. - pub fn launch_rl_controllers_per_step(&self) -> Result<()> { + /// Slot pairing (output <- input): + /// ISV[400] gamma <- ISV[417] MEAN_TRADE_DURATION_EMA + /// ISV[401] tau <- ISV[418] Q_DIVERGENCE_EMA + /// ISV[402] eps <- ISV[419] KL_PI_EMA + /// ISV[403] entropy_coef <- ISV[420] ENTROPY_OBSERVED_EMA + /// ISV[404] n_rollout_steps <- ISV[421] ADVANTAGE_VAR_RATIO_EMA + /// ISV[405] per_alpha <- ISV[422] TD_KURTOSIS_EMA + /// ISV[406] reward_scale <- ISV[423] MEAN_ABS_PNL_EMA + /// ISV[440] ratio_clamp <- ISV[402] (reads eps just written) + /// ISV[512,516,517] gates <- dones[B] + /// ISV[486] lambda_distill <- ISV[488] KL_EMA + pub fn launch_rl_fused_controllers( + &self, + b_size: usize, + ) -> Result<()> { + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; let alpha = RL_LR_CONTROLLER_ALPHA; - self.launch_isv_controller_3arg( - &self.rl_gamma_controller_fn, - alpha, - crate::rl::isv_slots::RL_MEAN_TRADE_DURATION_EMA_INDEX as i32, - ) - .context("R5 launch rl_gamma_controller")?; - self.launch_isv_controller_3arg( - &self.rl_target_tau_controller_fn, - alpha, - crate::rl::isv_slots::RL_Q_DIVERGENCE_EMA_INDEX as i32, - ) - .context("R5 launch rl_target_tau_controller")?; - self.launch_isv_controller_3arg( - &self.rl_ppo_clip_controller_fn, - alpha, - crate::rl::isv_slots::RL_KL_PI_EMA_INDEX as i32, - ) - .context("R5 launch rl_ppo_clip_controller")?; - self.launch_isv_controller_3arg( - &self.rl_entropy_coef_controller_fn, - alpha, - crate::rl::isv_slots::RL_ENTROPY_OBSERVED_EMA_INDEX as i32, - ) - .context("R5 launch rl_entropy_coef_controller")?; - self.launch_isv_controller_3arg( - &self.rl_rollout_steps_controller_fn, - alpha, - crate::rl::isv_slots::RL_ADVANTAGE_VAR_RATIO_EMA_INDEX as i32, - ) - .context("R5 launch rl_rollout_steps_controller")?; - self.launch_isv_controller_3arg( - &self.rl_per_alpha_controller_fn, - alpha, - crate::rl::isv_slots::RL_TD_KURTOSIS_EMA_INDEX as i32, - ) - .context("R5 launch rl_per_alpha_controller")?; - self.launch_isv_controller_3arg( - &self.rl_reward_scale_controller_fn, - alpha, - crate::rl::isv_slots::RL_MEAN_ABS_PNL_EMA_INDEX as i32, - ) - .context("R5 launch rl_reward_scale_controller")?; - // audit — PPO ratio clamp ceiling, anchored on ε at RL_PPO_CLIP_INDEX - // (which the rl_ppo_clip_controller just refreshed two calls up). - self.launch_isv_controller_3arg( - &self.rl_ppo_ratio_clamp_controller_fn, - alpha, - crate::rl::isv_slots::RL_PPO_CLIP_INDEX as i32, - ) - .context("R9 launch rl_ppo_ratio_clamp_controller")?; + let b_size_i = b_size as i32; + let mut launch = self.stream.launch_builder(&self.rl_fused_controllers_fn); + launch + .arg(&self.isv_d) + .arg(&self.dones_d) + .arg(&alpha) + .arg(&b_size_i) + .arg(&self.fused_ctrl_input_slots_d); + unsafe { + launch + .launch(cfg) + .context("rl_fused_controllers launch")?; + } Ok(()) } @@ -2539,7 +2571,6 @@ impl IntegratedTrainer { pos_state_d: &CudaSlice, b_size: usize, pos_bytes: usize, - current_step: i32, ) -> Result<()> { debug_assert_eq!(actions_d.len(), b_size); debug_assert_eq!(q_logits_d.len(), b_size * N_ACTIONS * Q_N_ATOMS); @@ -2559,8 +2590,7 @@ impl IntegratedTrainer { .arg(pos_state_d) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -2579,7 +2609,6 @@ impl IntegratedTrainer { pos_state_d: &CudaSlice, b_size: usize, pos_bytes: usize, - current_step: i32, ) -> Result<()> { debug_assert_eq!(actions_d.len(), b_size); let frd_out_dim = crate::rl::frd::FRD_OUT_DIM; @@ -2599,8 +2628,7 @@ impl IntegratedTrainer { .arg(pos_state_d) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -2630,7 +2658,6 @@ impl IntegratedTrainer { pyramid_units_count_d: &mut CudaSlice, b_size: usize, pos_bytes: usize, - current_step: i32, ) -> Result<()> { debug_assert_eq!(pos_state_d.len(), b_size * pos_bytes); debug_assert_eq!(prev_pos_lots_d.len(), b_size); @@ -2661,8 +2688,7 @@ impl IntegratedTrainer { .arg(unit_active_d) .arg(pyramid_units_count_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -3180,27 +3206,10 @@ impl IntegratedTrainer { } } - // λ_distill adaptive controller — runs AFTER distill kernel - // (which writes KL_EMA to slot 488). Adjusts λ slot 486 via - // Schulman bounded step on KL_EMA vs KL_TARGET slot 491. The - // new λ takes effect on the NEXT step's distill kernel. - // Single-thread, single-block. - { - let cfg_lambda_ctrl = LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }; - let mut launch = self - .stream - .launch_builder(&self.rl_q_distill_lambda_controller_fn); - launch.arg(&self.isv_d); - unsafe { - launch - .launch(cfg_lambda_ctrl) - .context("rl_q_distill_lambda_controller launch")?; - } - } + // NOTE: q_distill_lambda controller is now fused into + // rl_fused_controllers (fired in the per-step block). Reads + // KL_EMA from the previous step's backward pass — one-step + // delay is negligible for a 0.998/1.2 step-size controller. self.policy_head .backward_to_w_b_h( @@ -3994,6 +4003,24 @@ impl IntegratedTrainer { ); } + // ── Step 0: bump device-resident step counter (ISV[548]). + // Must run BEFORE any kernel that reads current_step from ISV. + // Single thread, single block — graph-safe (no scalar args change). + { + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&self.rl_increment_step_fn); + launch.arg(&self.isv_d); + unsafe { + launch + .launch(cfg) + .context("rl_increment_step launch")?; + } + } + // ── Step 1a (R7c): encoder forward on NEXT snapshots to land // h_{t+1} at slot K-1. Done FIRST so the subsequent // forward_encoder(snapshots) call leaves perception's internal @@ -4310,7 +4337,6 @@ impl IntegratedTrainer { }; let b_size_i = b_size as i32; let pos_bytes_i = lobsim.pos_bytes() as i32; - let current_step_i = (self.step_counter as i32).max(0); let mut launch = self.stream.launch_builder(&self.rl_confidence_gate_fn); launch .arg(&mut self.actions_d) @@ -4319,8 +4345,7 @@ impl IntegratedTrainer { .arg(pos_d_ref) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step_i); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -4338,7 +4363,6 @@ impl IntegratedTrainer { }; let b_size_i = b_size as i32; let pos_bytes_i = lobsim.pos_bytes() as i32; - let current_step_i = (self.step_counter as i32).max(0); let mut launch = self.stream.launch_builder(&self.rl_frd_gate_fn); launch .arg(&mut self.actions_d) @@ -4346,8 +4370,7 @@ impl IntegratedTrainer { .arg(pos_d_ref) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step_i); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -4616,7 +4639,6 @@ impl IntegratedTrainer { // Maintains its own prev_pos tracker (unit_prev_pos_lots_d) // separate from extract_realized_pnl_delta's prev_position_lots_d // so kernels are independently composable. - let current_step_i = (self.step_counter as i32).max(0); { let pos_d_ref: &CudaSlice = lobsim.pos_d(); let cfg = LaunchConfig { @@ -4637,8 +4659,7 @@ impl IntegratedTrainer { .arg(&mut self.unit_active_d) .arg(&mut self.pyramid_units_count_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step_i); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -4658,7 +4679,6 @@ impl IntegratedTrainer { }; let b_size_i = b_size as i32; let pos_bytes_i = lobsim.pos_bytes() as i32; - let current_step_i = (self.step_counter as i32).max(0); let mut launch = self.stream.launch_builder(&self.rl_trade_context_update_fn); launch @@ -4673,8 +4693,7 @@ impl IntegratedTrainer { .arg(ask_px_d) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i) - .arg(¤t_step_i); + .arg(&pos_bytes_i); unsafe { launch .launch(cfg) @@ -4881,35 +4900,14 @@ impl IntegratedTrainer { } } - // Fire all 7 RL controllers — each reads its ISV EMA-input - // slot and Wiener-blends its output. Critical: this happens - // BEFORE apply_reward_scale so the scale ISV[406] reflects - // this step's mean_abs_pnl EMA update. - self.launch_rl_controllers_per_step() - .context("launch_rl_controllers_per_step")?; - - // Gate threshold controller — adapts conf + FRD thresholds from - // trade frequency. Runs after other controllers, before reward scale. - { - let cfg = LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), - shared_mem_bytes: 0, - }; - let b_size_i = b_size as i32; - let current_step_i = (self.step_counter as i32).max(0); - let mut launch = self.stream.launch_builder(&self.rl_gate_threshold_controller_fn); - launch - .arg(&self.isv_d) - .arg(&self.dones_d) - .arg(&b_size_i) - .arg(¤t_step_i); - unsafe { - launch - .launch(cfg) - .context("rl_gate_threshold_controller launch")?; - } - } + // Fire all 10 RL ISV controllers in a single fused kernel + // launch (gamma, tau, ppo_clip, entropy_coef, rollout_steps, + // per_alpha, reward_scale, ppo_ratio_clamp, gate_threshold, + // q_distill_lambda). Critical: this happens BEFORE + // apply_reward_scale so the scale ISV[406] reflects this + // step's mean_abs_pnl EMA update. + self.launch_rl_fused_controllers(b_size) + .context("launch_rl_fused_controllers")?; // Apply the reward scale on device (in place on rewards_d) + // adaptive clamp + per-step pre-clamp diag + per-step diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 8f2a63536..ffe0c3664 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -43,8 +43,8 @@ use ml_alpha::rl::isv_slots::{ RL_ANTIMARTINGALE_KAPPA_INDEX, RL_ANTIMARTINGALE_MAX_INDEX, RL_ANTIMARTINGALE_MIN_INDEX, RL_CONF_GATE_LAMBDA_INDEX, RL_CONF_GATE_SIGMA_NORM_INDEX, RL_CONF_GATE_THRESHOLD_INDEX, RL_FRD_GATE_THR_LONG_INDEX, RL_FRD_GATE_THR_SHORT_INDEX, RL_HEAT_CAP_MAX_LOTS_INDEX, - RL_PYRAMID_THRESHOLD_INDEX, RL_TRAIL_ADJUST_RATE_INDEX, RL_TRAIL_K_INIT_INDEX, - RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX, + RL_PYRAMID_THRESHOLD_INDEX, RL_STEP_COUNTER_ISV_INDEX, RL_TRAIL_ADJUST_RATE_INDEX, + RL_TRAIL_K_INIT_INDEX, RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX, }; use ml_alpha::trainer::integrated::{ read_slice_d_pub, read_slice_i32_d_pub, read_slice_u8_d_pub, write_slice_f32_d_pub, @@ -323,6 +323,7 @@ fn unit_state_transitions() -> Result<()> { // ── Step 0: OPEN (was-flat → long 2 @ 100.0). let pos_open_d = upload_u8(&stream, &pos_buf(2, 100.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 0.0)?; trainer.launch_rl_unit_state_update( &pos_open_d, &mut prev_pos_d, @@ -335,7 +336,6 @@ fn unit_state_transitions() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 0, )?; let active = read_slice_u8_d_pub(&stream, &unit_active_d, b_size * MAX_UNITS)?; @@ -368,6 +368,7 @@ fn unit_state_transitions() -> Result<()> { // ── Step 1: CLOSE (long 2 → flat 0). let pos_close_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 1.0)?; trainer.launch_rl_unit_state_update( &pos_close_d, &mut prev_pos_d, @@ -380,7 +381,6 @@ fn unit_state_transitions() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 1, )?; let active = read_slice_u8_d_pub(&stream, &unit_active_d, b_size * MAX_UNITS)?; let lots = read_slice_i32_d_pub(&stream, &unit_lots_d, b_size * MAX_UNITS)?; @@ -660,6 +660,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { // ── Step 0: OPEN — flat → long 1 @ 100.0. Slot 0 active. let pos0_d = upload_u8(&stream, &pos_buf(1, 100.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 0.0)?; trainer.launch_rl_unit_state_update( &pos0_d, &mut prev_pos_d, @@ -672,7 +673,6 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 0, )?; let pyramid = read_slice_i32_d_pub(&stream, &pyramid_count_d, b_size)?; assert_eq!(pyramid[0], 1, "after OPEN: pyramid_count == 1"); @@ -680,6 +680,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { // ── Step 1: Pyramid ADD — long 1 → long 2 @ 101.0 (price moved). // |curr|=2 > |prev|=1 → allocates slot 1. let pos1_d = upload_u8(&stream, &pos_buf(2, 101.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 1.0)?; trainer.launch_rl_unit_state_update( &pos1_d, &mut prev_pos_d, @@ -692,7 +693,6 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 1, )?; let active = read_slice_u8_d_pub(&stream, &unit_active_d, b_size * MAX_UNITS)?; @@ -711,6 +711,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { // ── Step 2: Partial FLAT — long 2 → long 1 @ 101.0. // |curr|=1 < |prev|=2 → deactivates oldest active (slot 0). let pos2_d = upload_u8(&stream, &pos_buf(1, 101.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 2.0)?; trainer.launch_rl_unit_state_update( &pos2_d, &mut prev_pos_d, @@ -723,7 +724,6 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 2, )?; let active = read_slice_u8_d_pub(&stream, &unit_active_d, b_size * MAX_UNITS)?; @@ -737,6 +737,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { // ── Step 3: Full CLOSE — long 1 → flat 0. let pos3_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 3.0)?; trainer.launch_rl_unit_state_update( &pos3_d, &mut prev_pos_d, @@ -749,7 +750,6 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut pyramid_count_d, b_size, pos_bytes, - 3, )?; let active = read_slice_u8_d_pub(&stream, &unit_active_d, b_size * MAX_UNITS)?; @@ -901,6 +901,8 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_THRESHOLD_INDEX, 0.5)?; set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_LAMBDA_INDEX, 1.0)?; set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0)?; + // Step counter past warmup so gates are active. + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 99999.0)?; // Flat position (opening action should be gated). let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; @@ -918,7 +920,6 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -943,7 +944,6 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?; assert_eq!( @@ -961,7 +961,6 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_long_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?; assert_eq!( @@ -978,7 +977,6 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?; assert_eq!( @@ -1001,6 +999,8 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_LONG_INDEX, 0.35)?; set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_SHORT_INDEX, 0.35)?; + // Step counter past warmup so gates are active. + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 99999.0)?; let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; let frd_dim = FRD_N_HORIZONS * FRD_N_ATOMS; @@ -1017,7 +1017,6 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -1041,7 +1040,6 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?; assert_eq!( @@ -1062,7 +1060,6 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?; assert_eq!( @@ -1080,7 +1077,6 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_long_d, b_size, pos_bytes, - 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?; assert_eq!( @@ -1375,6 +1371,8 @@ fn both_gates_block_same_action() -> Result<()> { set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0)?; set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_LONG_INDEX, 0.99)?; set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_SHORT_INDEX, 0.99)?; + // Step counter past warmup so gates are active. + set_isv_slot(&mut trainer, &stream, RL_STEP_COUNTER_ISV_INDEX, 99999.0)?; let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; @@ -1394,7 +1392,6 @@ fn both_gates_block_same_action() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions_after_conf = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -1409,7 +1406,6 @@ fn both_gates_block_same_action() -> Result<()> { &pos_flat_d, b_size, pos_bytes, - 99999, )?; let actions_after_frd = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!(