gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

perf(rl): device-resident step counter + fused controllers (-9 launches)

Browse Source 
Two CUDA Graph prerequisites implemented:

1. Device-resident step counter (ISV[548]):
   - New rl_increment_step.cu kernel (single thread, ISV += 1)
   - All kernels that took current_step as scalar now read from ISV
   - Updated: confidence_gate, frd_gate, unit_state_update,
     trade_context_update, gate_threshold_controller
   - Enables CUDA Graph capture (no scalar arg changes between replays)

2. Fused controllers (rl_fused_controllers.cu):
   - Combines 10 single-thread controllers into 1 kernel launch:
     gamma, tau, ppo_clip, entropy, rollout_steps, per_alpha,
     reward_scale (with ±2% clamp), ppo_ratio_clamp,
     gate_threshold, q_distill_lambda
   - Saves 9 kernel launches per step (~40-80μs)
   - Individual .cu files retained for testing/documentation

ISV slot 548 (step counter). Local smoke: 100 steps, no crash.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-25 21:35:50 +02:00
parent 993201250a
commit a7c1d763d8
11 changed files with 661 additions and 166 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
crates/ml-alpha/build.rs
View File

@@ -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

5
crates/ml-alpha/cuda/rl_confidence_gate.cu
View File

@@ -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;

5
crates/ml-alpha/cuda/rl_frd_gate.cu
View File

@@ -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;

464
crates/ml-alpha/cuda/rl_fused_controllers.cu Normal file
View File

@@ -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 <stdint.h>
// ═══════════════════════════════════════════════════════════════════════
// 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;
}
}
}

5
crates/ml-alpha/cuda/rl_gate_threshold_controller.cu
View File

@@ -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;

20
crates/ml-alpha/cuda/rl_increment_step.cu Normal file
View File

@@ -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;
}

6
crates/ml-alpha/cuda/rl_trade_context_update.cu
View File

@@ -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;

6
crates/ml-alpha/cuda/rl_unit_state_update.cu
View File

@@ -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];

12
crates/ml-alpha/src/rl/isv_slots.rs
View File

@@ -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;

270
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -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<CudaModule>,
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<CudaModule>,
#[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<CudaModule>,
@@ -510,7 +524,16 @@ pub struct IntegratedTrainer {
_rl_recent_outcome_update_module: Arc<CudaModule>,
rl_recent_outcome_update_fn: CudaFunction,
_rl_gate_threshold_controller_module: Arc<CudaModule>,
// 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<CudaModule>,
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<i32>,
_rl_reward_shaping_module: Arc<CudaModule>,
rl_reward_shaping_fn: CudaFunction,
_rl_asymmetric_trail_decay_module: Arc<CudaModule>,
@@ -523,6 +546,9 @@ pub struct IntegratedTrainer {
rl_trade_context_update_fn: CudaFunction,
_rl_multires_features_update_module: Arc<CudaModule>,
rl_multires_features_update_fn: CudaFunction,
// Device-resident step counter (ISV[548] += 1.0 per step).
_rl_increment_step_module: Arc<CudaModule>,
rl_increment_step_fn: CudaFunction,
pub outcome_ema_d: CudaSlice<f32>,
pub trade_context_d: CudaSlice<f32>,
pub multires_output_d: CudaSlice<f32>,
@@ -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::<i32>(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<u8>,
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(&current_step);
.arg(&pos_bytes_i);
unsafe {
launch
.launch(cfg)
@@ -2579,7 +2609,6 @@ impl IntegratedTrainer {
pos_state_d: &CudaSlice<u8>,
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(&current_step);
.arg(&pos_bytes_i);
unsafe {
launch
.launch(cfg)
@@ -2630,7 +2658,6 @@ impl IntegratedTrainer {
pyramid_units_count_d: &mut CudaSlice<i32>,
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(&current_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(&current_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(&current_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<u8> = 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(&current_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(&current_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(&current_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

32
crates/ml-alpha/tests/trade_management_kernels.rs
View File

@@ -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!(