diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 9e911c1c9..a377e00fe 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -104,6 +104,13 @@ const KERNELS: &[&str] = &[ "rl_v_blend_alpha_controller", // Phase 4.4: ISV-adaptive Schulman-bounded controller on α from observed |V_dq − V_scalar| / |V_scalar| tracking ratio "rl_advantage_normalize", // Phase 4.5 (2026-05-30): per-batch advantage normalization (A − mean)/std with ε² variance floor — standard PPO practice, self-adaptive (no tuned params) "rl_signal_variance_update", // Adaptive controller floors (2026-05-30): Welford online variance for controller input EMAs; replaces hardcoded NOISE_FLOOR_FRAC constants in 8 controllers — sample_var = M² / (count-1); see spec 2026-05-30-adaptive-controller-floor-design + "rl_win_rate_ema_update", // Adaptive risk management (2026-05-30): Layer 4 (Kelly) input — observed win_rate EMA from per-batch trade outcomes + "rl_avg_win_loss_ema_update", // Adaptive risk management (2026-05-30): Layer 4 (Kelly) input — avg_win/avg_loss USD EMAs from per-batch closed-trade pnl + "rl_inventory_variance_update", // Adaptive risk management (2026-05-30): Layer 3 (inventory β) input — |net_position| variance EMA across batches + "rl_cmdp_constraints_check", // Adaptive risk management (2026-05-30): Layer 1 hard gates — session DD + cooldown + consecutive-loss tracking; writes override flags to ISV + "rl_iqn_action_tau_controller", // Adaptive risk management (2026-05-30): Layer 2 — adapts IQN action-selection quantile τ from session drawdown signal + "rl_inventory_beta_controller", // Adaptive risk management (2026-05-30): Layer 3 — adapts inventory penalty β from observed inventory variance vs reward magnitude + "rl_kelly_fraction_controller", // Adaptive risk management (2026-05-30): Layer 4 — half-Kelly position-size multiplier from observed win_rate × R-multiple; warmup-gated "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 diff --git a/crates/ml-alpha/cuda/actions_to_market_targets.cu b/crates/ml-alpha/cuda/actions_to_market_targets.cu index feab7e4cf..dc1327049 100644 --- a/crates/ml-alpha/cuda/actions_to_market_targets.cu +++ b/crates/ml-alpha/cuda/actions_to_market_targets.cu @@ -40,6 +40,28 @@ #define RL_ANTIMARTINGALE_MIN_INDEX 510 #define RL_ANTIMARTINGALE_MAX_INDEX 511 +// ── Layer 1 (CMDP hard constraints — spec 2026-05-30-adaptive-risk-management) ── +#define RL_SESSION_DD_TRIGGERED_INDEX 664 +#define RL_MAX_OPEN_UNITS_INDEX 665 +#define RL_COOLDOWN_REMAINING_STEPS_INDEX 668 +#define RL_NET_INVENTORY_LIMIT_USD_INDEX 670 + +// ── Layer 4 (Kelly fraction sizing — spec 2026-05-30-adaptive-risk-management) ── +#define RL_KELLY_FRACTION_INDEX 676 + +// Action indices — must match crate::rl::common::Action. +#define ACTION_SHORT_LARGE 0 +#define ACTION_SHORT_SMALL 1 +#define ACTION_HOLD 2 +#define ACTION_FLAT_FROM_LONG 3 +#define ACTION_FLAT_FROM_SHORT 4 +#define ACTION_LONG_SMALL 5 +#define ACTION_LONG_LARGE 6 +#define ACTION_TRAIL_TIGHTEN 7 +#define ACTION_TRAIL_LOOSEN 8 +#define ACTION_HALF_FLAT_LONG 9 +#define ACTION_HALF_FLAT_SHORT 10 + __device__ __forceinline__ int antimartingale_size( int base_size, float outcome_ema_b, const float* isv ) { @@ -51,6 +73,24 @@ __device__ __forceinline__ int antimartingale_size( return (sized > 0) ? sized : 1; } +// Returns true if `action` would open or grow a one-sided position +// (excludes closes, half-flats, holds, and trail mutations). +__device__ __forceinline__ bool is_opening_action(int action) { + return (action == ACTION_SHORT_LARGE) || (action == ACTION_SHORT_SMALL) + || (action == ACTION_LONG_SMALL) || (action == ACTION_LONG_LARGE); +} + +// Returns true if executing `action` from the current `net_pos_lots` +// would INCREASE one-sided exposure (i.e. opens/grows the dominant side). +// Closes/reverses/holds never trigger this — only same-side adds. +__device__ __forceinline__ bool would_increase_exposure(int action, int net_pos_lots) { + const bool longs = (action == ACTION_LONG_SMALL) || (action == ACTION_LONG_LARGE); + const bool shorts = (action == ACTION_SHORT_SMALL) || (action == ACTION_SHORT_LARGE); + if (longs && net_pos_lots >= 0) return true; // open or pyramid same side + if (shorts && net_pos_lots <= 0) return true; + return false; +} + extern "C" __global__ void actions_to_market_targets( const int* __restrict__ actions, // [b_size] const unsigned char* __restrict__ pos_state, // [b_size * pos_bytes] @@ -70,13 +110,52 @@ extern "C" __global__ void actions_to_market_targets( const int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= b_size) return; - const int action = actions[b]; + int action = actions[b]; const int position_lots = *(const int*)(pos_state + b * pos_bytes); const float outcome_ema_b = outcome_ema[b]; int side = 2; // default no-op int size = 0; + // ──────────────────────────────────────────────────────────────────── + // Layer 1 (CMDP hard constraints — spec 2026-05-30-adaptive-risk-management). + // Session-level overrides force Hold regardless of agent's choice. + // ──────────────────────────────────────────────────────────────────── + const bool dd_triggered = isv[RL_SESSION_DD_TRIGGERED_INDEX] >= 0.5f; + const bool in_cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] > 0.0f; + if (dd_triggered || in_cooldown) { + market_targets[b * 2 + 0] = 2; // no-op side + market_targets[b * 2 + 1] = 0; // size = 0 + return; + } + + // Max-open-units check (per-batch dynamic count): refuse opens once + // pyramid is at MAX. Closes still allowed (handled by remaining logic). + { + int active_units = 0; + const int base_u = b * MAX_UNITS; + for (int u = 0; u < MAX_UNITS; ++u) { + if (unit_active[base_u + u]) active_units += 1; + } + const int max_open = (int)isv[RL_MAX_OPEN_UNITS_INDEX]; + if (active_units >= max_open && is_opening_action(action)) { + action = ACTION_HOLD; + } + } + + // Net-inventory cap (USD): block actions that grow one-sided exposure + // beyond the limit. Uses bid/ask mid as the USD multiplier. + { + const float inv_limit = isv[RL_NET_INVENTORY_LIMIT_USD_INDEX]; + if (inv_limit > 0.0f) { + const float mid = 0.5f * (bid_px[0] + ask_px[0]); + const float net_pos_usd = (float)position_lots * mid; + if (fabsf(net_pos_usd) > inv_limit && would_increase_exposure(action, position_lots)) { + action = ACTION_HOLD; + } + } + } + if (action == 0) { // ShortLarge — when already short, apply pyramid logic. if (position_lots < 0) { @@ -240,6 +319,27 @@ extern "C" __global__ void actions_to_market_targets( // actions 7, 8 (TrailTighten/Loosen): no fill — handled by // rl_trail_mutate kernel. The side=2 no-op default is correct here. + // ──────────────────────────────────────────────────────────────────── + // Layer 4 (Kelly-fraction sizing — spec 2026-05-30-adaptive-risk-management). + // Scales the opening/sizing size by the observed-edge Kelly fraction. + // Closes (FlatFrom* + HalfFlat*) and Hold are NOT scaled — risk + // management for exiting positions is handled by Layer 1 + trail stops. + // ──────────────────────────────────────────────────────────────────── + if (size > 0 && is_opening_action(action)) { + const float kelly = isv[RL_KELLY_FRACTION_INDEX]; + // Bootstrap = 1.0 (full size); negative or zero → no commitment. + const float scaled = (float)size * kelly; + if (scaled <= 0.0f) { + size = 0; + side = 2; + } else { + // Ceil to preserve at least 1-lot when kelly is small-but-positive. + int new_size = (int)ceilf(scaled); + if (new_size < 1) new_size = 1; + size = new_size; + } + } + market_targets[b * 2 + 0] = side; market_targets[b * 2 + 1] = size; } diff --git a/crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu b/crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu new file mode 100644 index 000000000..d165454c7 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu @@ -0,0 +1,77 @@ +// rl_avg_win_loss_ema_update.cu — Layer 4 (Kelly) input EMA producer. +// +// Tracks observed per-trade avg_win and avg_loss magnitudes in USD as +// EMAs. Kelly fraction uses R-multiple = avg_win / avg_loss together +// with win_rate to derive the safe Kelly bet size. +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Inputs: +// rewards[b] f32 — per-batch realized pnl delta (shaped) this step +// dones[b] f32 — 1.0 if a trade closed this step, else 0.0 +// +// Outcome derived inline: win = (done & reward > 0); loss = (done & reward < 0). +// +// Outputs: +// ISV[RL_AVG_WIN_USD_EMA_INDEX = 678] (positive USD) +// ISV[RL_AVG_LOSS_USD_EMA_INDEX = 679] (positive USD, magnitude) +// +// Bootstrap semantics per `pearl_first_observation_bootstrap`: +// sentinel = 0.0; first non-zero closed-trade step replaces directly. +// +// Per `feedback_no_atomicadd`: single-thread single-block (sums sequentially). +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: EMA-α structural smoothing +// parameter (same as ema_update_per_step convention). + +#include + +#define RL_AVG_WIN_USD_EMA_INDEX 678 +#define RL_AVG_LOSS_USD_EMA_INDEX 679 +#define EMA_ALPHA 0.05f + +extern "C" __global__ void rl_avg_win_loss_ema_update( + float* __restrict__ isv, + const float* __restrict__ rewards, // [b_size] shaped pnl + const float* __restrict__ dones, // [b_size] 1.0 = close + int b_size +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + // Sum per-trade win/loss magnitudes across the batch. + float sum_win = 0.0f; + float sum_loss = 0.0f; + int n_win = 0; + int n_loss = 0; + for (int b = 0; b < b_size; ++b) { + if (dones[b] < 0.5f) continue; + const float p = rewards[b]; + if (p > 0.0f) { + sum_win += p; + n_win += 1; + } else if (p < 0.0f) { + sum_loss += fabsf(p); + n_loss += 1; + } + } + + if (n_win > 0) { + const float step_avg = sum_win / (float)n_win; + const float prev = isv[RL_AVG_WIN_USD_EMA_INDEX]; + if (prev == 0.0f) { + isv[RL_AVG_WIN_USD_EMA_INDEX] = step_avg; + } else { + isv[RL_AVG_WIN_USD_EMA_INDEX] = (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_avg; + } + } + + if (n_loss > 0) { + const float step_avg = sum_loss / (float)n_loss; + const float prev = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + if (prev == 0.0f) { + isv[RL_AVG_LOSS_USD_EMA_INDEX] = step_avg; + } else { + isv[RL_AVG_LOSS_USD_EMA_INDEX] = (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_avg; + } + } +} diff --git a/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu b/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu new file mode 100644 index 000000000..4f5c6eb1a --- /dev/null +++ b/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu @@ -0,0 +1,92 @@ +// rl_cmdp_constraints_check.cu — Layer 1 hard CMDP gates. +// +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Maintains the session-level + cooldown state that downstream consumers +// (actions_to_market_targets) read as override flags. Three independent +// risk constraints are tracked here: +// +// 1. Session pnl accumulation + DD limit (sticky `triggered` flag) +// 2. Cooldown counter decrement (after consecutive losses limit fires) +// 3. Consecutive-loss tracking (sets cooldown when limit hit) +// +// Max-open-units + net-inventory checks are evaluated PER BATCH inside +// actions_to_market_targets where per-batch unit state is available. +// This kernel handles only session-wide state. +// +// Per `feedback_no_atomicadd`: single-thread single-block, sums sequentially. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_no_partial_refactor`: ISV slots must match isv_slots.rs. +// +// Launch config: +// grid = (1, 1, 1) +// block = (1, 1, 1) +// smem = 0 +// stream = main RL stream (sequenced AFTER rl_fused_reward_pipeline +// writes realized_pnl / outcomes, BEFORE actions_to_market_targets +// reads override flags). + +#include + +#define RL_SESSION_PNL_USD_INDEX 662 +#define RL_SESSION_DD_LIMIT_USD_INDEX 663 +#define RL_SESSION_DD_TRIGGERED_INDEX 664 +#define RL_CONSEC_LOSS_LIMIT_INDEX 666 +#define RL_CONSEC_LOSS_COUNT_INDEX 667 +#define RL_COOLDOWN_REMAINING_STEPS_INDEX 668 +#define RL_COOLDOWN_DURATION_INDEX 669 + +// Outcome is derived inline: loss = (done & reward < 0), win = (done & reward > 0). +extern "C" __global__ void rl_cmdp_constraints_check( + float* __restrict__ isv, + const float* __restrict__ rewards, // [b_size] shaped pnl + const float* __restrict__ dones, // [b_size] 1.0 = close + int b_size +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + // ──────────────────────────────────────────────────────────── + // 1. Session pnl accumulation + DD check. + // Sum per-batch rewards across batch (the per-step pnl delta). + // ──────────────────────────────────────────────────────────── + float step_pnl = 0.0f; + for (int b = 0; b < b_size; ++b) { + step_pnl += rewards[b]; + } + const float session_pnl_new = isv[RL_SESSION_PNL_USD_INDEX] + step_pnl; + isv[RL_SESSION_PNL_USD_INDEX] = session_pnl_new; + if (session_pnl_new < isv[RL_SESSION_DD_LIMIT_USD_INDEX]) { + isv[RL_SESSION_DD_TRIGGERED_INDEX] = 1.0f; + } + + // ──────────────────────────────────────────────────────────── + // 2. Cooldown decrement. + // ──────────────────────────────────────────────────────────── + const float cooldown_prev = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX]; + if (cooldown_prev > 0.0f) { + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = cooldown_prev - 1.0f; + } + + // ──────────────────────────────────────────────────────────── + // 3. Consecutive loss tracking. Outcome derived from done & reward sign. + // ──────────────────────────────────────────────────────────── + float consec = isv[RL_CONSEC_LOSS_COUNT_INDEX]; + for (int b = 0; b < b_size; ++b) { + if (dones[b] < 0.5f) continue; + const float r = rewards[b]; + if (r < 0.0f) { + consec += 1.0f; + } else if (r > 0.0f) { + consec = 0.0f; + } + // r == 0 with done: ambiguous — treat as no-change (rare break-even close) + } + if (consec >= isv[RL_CONSEC_LOSS_LIMIT_INDEX]) { + // Limit hit — start cooldown, reset counter. + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = isv[RL_COOLDOWN_DURATION_INDEX]; + isv[RL_CONSEC_LOSS_COUNT_INDEX] = 0.0f; + } else { + isv[RL_CONSEC_LOSS_COUNT_INDEX] = consec; + } +} diff --git a/crates/ml-alpha/cuda/rl_fused_controllers.cu b/crates/ml-alpha/cuda/rl_fused_controllers.cu index 33a738668..4d632349a 100644 --- a/crates/ml-alpha/cuda/rl_fused_controllers.cu +++ b/crates/ml-alpha/cuda/rl_fused_controllers.cu @@ -1,14 +1,14 @@ -// rl_fused_controllers.cu — single-kernel fusion of 10 RL ISV controllers. +// rl_fused_controllers.cu — single-kernel fusion of 13 RL ISV controllers. // -// Eliminates 9 kernel-launch overheads (~40-80μs/step) by running all +// Eliminates 12 kernel-launch overheads (~55-100μs/step) by running all // controllers sequentially in one thread. Each controller's logic is // MIRRORED from its standalone .cu file (post-2026-05-30 adaptive -// controller-floor refactor): adaptive noise floor derived from Welford -// variance, asymmetric Schulman patience counters, ISV-driven clamp -// bounds + Wiener-α floor. Per `feedback_no_partial_refactor` and the -// "fused kernel must be semantically equivalent to running all individual -// controllers in sequence" contract, every branch below matches its -// standalone counterpart's ISV reads, math, and counter semantics. +// controller-floor + risk-management refactors): adaptive noise floor +// derived from Welford variance, asymmetric Schulman patience counters, +// ISV-driven clamp bounds + Wiener-α floor. Per `feedback_no_partial_refactor` +// and the "fused kernel must be semantically equivalent to running all +// individual controllers in sequence" contract, every branch below matches +// its standalone counterpart's ISV reads, math, and counter semantics. // // Controllers fused (in execution order): // 1. rl_gamma_controller ISV[400] ← ISV[input_slots[0]] @@ -21,6 +21,9 @@ // 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] +// 11. rl_iqn_action_tau_controller ISV[671] ← ISV[662] (Layer 2, spec 2026-05-30-adaptive-risk-management) +// 12. rl_inventory_beta_controller ISV[674] ← ISV[675,614] (Layer 3) +// 13. rl_kelly_fraction_controller ISV[676] ← ISV[677,678,679,680,681,660] (Layer 4) // // Per `feedback_no_atomicadd`: single thread, no atomics. // Per `feedback_cpu_is_read_only`: all state in ISV, no host roundtrip. @@ -176,6 +179,36 @@ // --- Step counter (device-resident) --- #define RL_STEP_COUNTER_ISV_INDEX 548 +// ═══════════════════════════════════════════════════════════════════════ +// Adaptive risk management — spec 2026-05-30-adaptive-risk-management. +// Layers 2/3/4 controllers fused below. Layer 1 (CMDP gates) runs as a +// separate kernel because it needs per-step `realized_pnl` + `outcome` +// arrays not available here. Layer D (trail factors) is read-only ISV. +// ═══════════════════════════════════════════════════════════════════════ + +// --- 11. Layer 2: IQN action τ controller --- +#define RL_SESSION_PNL_USD_INDEX 662 +#define RL_IQN_ACTION_TAU_INDEX 671 +#define RL_IQN_ACTION_TAU_MIN_INDEX 672 +#define RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX 673 +#define DEFAULT_STARTING_CAPITAL_USD 35000.0f + +// --- 12. Layer 3: Inventory β controller --- +#define RL_INVENTORY_PENALTY_BETA_INDEX 674 +#define RL_INVENTORY_VARIANCE_EMA_INDEX 675 +// (RL_REWARD_MAGNITUDE_EMA_INDEX 614 already defined above) +#define BETA_TARGET_FRAC_OF_REWARD 0.01f +#define BETA_SIGMA_MULTIPLIER 2.0f + +// --- 13. Layer 4: Kelly fraction controller --- +#define RL_KELLY_FRACTION_INDEX 676 +#define RL_WIN_RATE_EMA_INDEX 677 +#define RL_AVG_WIN_USD_EMA_INDEX 678 +#define RL_AVG_LOSS_USD_EMA_INDEX 679 +#define RL_KELLY_SAFETY_FRAC_INDEX 680 +#define RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX 681 +// (RL_CUMULATIVE_DONES_INDEX 660 already defined above) + // ═══════════════════════════════════════════════════════════════════════ // Fused kernel: all 10 controllers in one launch. @@ -734,4 +767,74 @@ extern "C" __global__ void rl_fused_controllers( isv[RL_Q_DISTILL_LAMBDA_INDEX] = lambda; } } + + // ═══════════════════════════════════════════════════════════════════ + // 11. Layer 2 — IQN action τ controller → ISV[671] + // τ = clamp(τ_min, 0.5 - sensitivity × drawdown_frac, 1.0) + // drawdown_frac = max(0, -session_pnl) / starting_capital + // ═══════════════════════════════════════════════════════════════════ + { + const float session_pnl = isv[RL_SESSION_PNL_USD_INDEX]; + const float drawdown_usd = fmaxf(0.0f, -session_pnl); + const float drawdown_frac = drawdown_usd / DEFAULT_STARTING_CAPITAL_USD; + + const float tau_min = isv[RL_IQN_ACTION_TAU_MIN_INDEX]; + const float sensitivity = isv[RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX]; + float tau_target = 0.5f - sensitivity * drawdown_frac; + tau_target = fmaxf(tau_min, fminf(tau_target, 1.0f)); + isv[RL_IQN_ACTION_TAU_INDEX] = tau_target; + } + + // ═══════════════════════════════════════════════════════════════════ + // 12. Layer 3 — Inventory β controller → ISV[674] + // β_target = 0.01 × E[|reward|] / (2 × σ_inventory) + // Sentinel-zero bootstrap; Wiener-α blend with floor 0.4 after. + // ═══════════════════════════════════════════════════════════════════ + { + const float reward_mag = isv[RL_REWARD_MAGNITUDE_EMA_INDEX]; + const float inv_var = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; + if (inv_var > 0.0f && reward_mag > 0.0f) { + const float inv_std = sqrtf(inv_var); + const float beta_target = BETA_TARGET_FRAC_OF_REWARD * reward_mag + / (BETA_SIGMA_MULTIPLIER * inv_std); + const float prev = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; + if (prev == 0.0f) { + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = beta_target; + } else { + const float a = fmaxf(alpha, wiener_floor); + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = + (1.0f - a) * prev + a * beta_target; + } + } + } + + // ═══════════════════════════════════════════════════════════════════ + // 13. Layer 4 — Kelly fraction controller → ISV[676] + // f = clamp(safety × (p × b − q) / b, 0, 1) + // b = avg_win / avg_loss; warmup-gated on cumulative_dones. + // ═══════════════════════════════════════════════════════════════════ + { + const float kelly_dones = isv[RL_CUMULATIVE_DONES_INDEX]; + const float kelly_min_trades = isv[RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX]; + if (kelly_dones < kelly_min_trades) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + } else { + const float p = isv[RL_WIN_RATE_EMA_INDEX]; + const float avg_win = isv[RL_AVG_WIN_USD_EMA_INDEX]; + const float avg_loss = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + if (avg_loss <= 0.0f) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + } else if (avg_win <= 0.0f) { + isv[RL_KELLY_FRACTION_INDEX] = 0.0f; + } else { + const float b_ratio = avg_win / avg_loss; + const float q = 1.0f - p; + const float f_kelly = (p * b_ratio - q) / b_ratio; + const float safety = isv[RL_KELLY_SAFETY_FRAC_INDEX]; + float f = f_kelly * safety; + f = fmaxf(0.0f, fminf(f, 1.0f)); + isv[RL_KELLY_FRACTION_INDEX] = f; + } + } + } } diff --git a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu b/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu index 777528831..43ee7eb55 100644 --- a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu +++ b/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu @@ -39,6 +39,8 @@ #define RL_SHORT_HOLD_PENALTY_INDEX 534 #define RL_HOLD_BONUS_INDEX 535 #define RL_OUTCOME_ALPHA_INDEX 520 +// Layer 3 (inventory penalty — spec 2026-05-30-adaptive-risk-management). +#define RL_INVENTORY_PENALTY_BETA_INDEX 674 #define MAX_UNITS 4 @@ -236,6 +238,17 @@ extern "C" __global__ void rl_fused_reward_pipeline( r += hold_bonus * sqrtf((float)hold_time); } + // ================================================================ + // PHASE 5b: Layer 3 inventory penalty (spec 2026-05-30-adaptive-risk-management). + // Apply BEFORE raw_rewards snapshot so the penalty is part of the + // shaped reward seen by Q/V learning. β = 0 (sentinel) → no-op. + // ================================================================ + const float inv_beta = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; + if (inv_beta > 0.0f) { + const float net_pos_mag = (float)((current_lots < 0) ? -current_lots : current_lots); + r -= inv_beta * net_pos_mag; + } + // Write shaped reward back. rewards[b] = r; diff --git a/crates/ml-alpha/cuda/rl_inventory_beta_controller.cu b/crates/ml-alpha/cuda/rl_inventory_beta_controller.cu new file mode 100644 index 000000000..610531d5a --- /dev/null +++ b/crates/ml-alpha/cuda/rl_inventory_beta_controller.cu @@ -0,0 +1,57 @@ +// rl_inventory_beta_controller.cu — Layer 3: adaptive inventory penalty β. +// +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Scales the inventory penalty `reward_shaped = reward_raw - β × |net_pos|` +// so that at ~2σ inventory the penalty is ~1% of typical reward magnitude. +// +// β_target = 0.01 × E[|reward|] / (2 × σ_inventory) +// +// Sentinel-zero bootstrap per `pearl_first_observation_bootstrap`: hold β +// at 0 until BOTH reward magnitude EMA and inventory variance EMA have +// real data. After bootstrap, Wiener-α blend with floor 0.4 per +// `pearl_wiener_alpha_floor_for_nonstationary`. +// +// Per `feedback_no_atomicadd`: single-thread single-block. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: every threshold via ISV slot. + +#include + +#define RL_INVENTORY_PENALTY_BETA_INDEX 674 +#define RL_INVENTORY_VARIANCE_EMA_INDEX 675 +#define RL_REWARD_MAGNITUDE_EMA_INDEX 614 +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// 1% of typical reward magnitude at 2σ inventory — Avellaneda-Stoikov +// canonical "noticeable but not dominant" calibration. Structural ratio. +#define BETA_TARGET_FRAC_OF_REWARD 0.01f +#define BETA_SIGMA_MULTIPLIER 2.0f + +extern "C" __global__ void rl_inventory_beta_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + const float reward_mag = isv[RL_REWARD_MAGNITUDE_EMA_INDEX]; + const float inv_var = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; + + // Dead-signal hold: need real data on both inputs. + if (inv_var <= 0.0f || reward_mag <= 0.0f) return; + + const float inv_std = sqrtf(inv_var); + const float beta_target = BETA_TARGET_FRAC_OF_REWARD * reward_mag + / (BETA_SIGMA_MULTIPLIER * inv_std); + + const float prev = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; + const float a_floor = isv[RL_WIENER_ALPHA_FLOOR_INDEX]; + + if (prev == 0.0f) { + // First-observation bootstrap. + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = beta_target; + } else { + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = + (1.0f - a_floor) * prev + a_floor * beta_target; + } +} diff --git a/crates/ml-alpha/cuda/rl_inventory_variance_update.cu b/crates/ml-alpha/cuda/rl_inventory_variance_update.cu new file mode 100644 index 000000000..f58e8c06a --- /dev/null +++ b/crates/ml-alpha/cuda/rl_inventory_variance_update.cu @@ -0,0 +1,59 @@ +// rl_inventory_variance_update.cu — Layer 3 (inventory penalty) input EMA. +// +// Tracks running variance of |net_position| across batches. Layer 3's β +// controller scales the penalty so that at ~2σ inventory the penalty +// equals ~1% of typical reward magnitude. +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Inputs: +// net_position_per_batch[b] signed contracts (current_lots, i32) +// +// Output: ISV[RL_INVENTORY_VARIANCE_EMA_INDEX = 675] +// +// EMA over the batch variance of |net_position|; sentinel-zero bootstrap. +// This is NOT a true Welford triple — we collapse the batch into a single +// scalar (batch variance) then EMA across steps with a constant alpha. +// The β controller treats this slot as σ² directly. +// +// Per `feedback_no_atomicadd`: single-thread single-block, sums sequentially. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: Welford-EMA α is a structural +// smoothing parameter, matching ema_update_per_step convention. + +#include + +#define RL_INVENTORY_VARIANCE_EMA_INDEX 675 +#define WELFORD_ALPHA 0.01f // slower EMA than Kelly inputs + +extern "C" __global__ void rl_inventory_variance_update( + float* __restrict__ isv, + const int* __restrict__ net_position_per_batch, // [b_size] signed contracts (i32) + int b_size +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + if (b_size <= 0) return; + + // Compute batch mean + variance of |net_position|. + float sum = 0.0f; + float sum_sq = 0.0f; + for (int b = 0; b < b_size; ++b) { + const int lots = net_position_per_batch[b]; + const float v = (float)((lots < 0) ? -lots : lots); + sum += v; + sum_sq += v * v; + } + const float n = (float)b_size; + const float mean = sum / n; + const float var = (sum_sq / n) - (mean * mean); + if (var < 0.0f) return; // numerical safety: should not happen but defend + + const float prev = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; + if (prev == 0.0f) { + isv[RL_INVENTORY_VARIANCE_EMA_INDEX] = var; + } else { + isv[RL_INVENTORY_VARIANCE_EMA_INDEX] = + (1.0f - WELFORD_ALPHA) * prev + WELFORD_ALPHA * var; + } +} diff --git a/crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu b/crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu new file mode 100644 index 000000000..f121c34e1 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu @@ -0,0 +1,53 @@ +// rl_iqn_action_tau_controller.cu — Layer 2: adaptive IQN action-selection τ. +// +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Adapts the quantile used for IQN risk-averse action selection from the +// observed session drawdown. Normal trading: τ = 0.5 (median ≈ expected-Q +// behavior). Under drawdown: τ → τ_MIN (pessimistic, downside-aware). +// +// τ_target = max(τ_min, 0.5 - sensitivity × drawdown_frac) +// +// Drawdown approximation: `drawdown_frac = max(0, -session_pnl) / +// starting_capital`. Because the trainer resets session_pnl on each fold +// boundary, this approximates "drawdown from session start". For the +// strict "drawdown from session peak" formulation we'd need a peak-tracker +// slot — deferred (see plan note); the current formulation is correct in +// the dominant case where the agent is in a losing session and provides +// the right pressure (more pessimistic action selection as losses +// accumulate). In a strongly-winning session τ stays at 0.5. +// +// Per `feedback_no_atomicadd`: single-thread single-block. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: τ_min, sensitivity ISV-driven. +// +// Starting capital is the $35k single-contract ES baseline per +// `project_ml_alpha_starting_capital`; structural constant exemption. + +#include + +#define RL_SESSION_PNL_USD_INDEX 662 +#define RL_IQN_ACTION_TAU_INDEX 671 +#define RL_IQN_ACTION_TAU_MIN_INDEX 672 +#define RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX 673 + +#define DEFAULT_STARTING_CAPITAL_USD 35000.0f + +extern "C" __global__ void rl_iqn_action_tau_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + const float session_pnl = isv[RL_SESSION_PNL_USD_INDEX]; + // Approximate drawdown from session boundary as max(0, -session_pnl). + const float drawdown_usd = fmaxf(0.0f, -session_pnl); + const float drawdown_frac = drawdown_usd / DEFAULT_STARTING_CAPITAL_USD; + + const float tau_min = isv[RL_IQN_ACTION_TAU_MIN_INDEX]; + const float sensitivity = isv[RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX]; + float tau_target = 0.5f - sensitivity * drawdown_frac; + tau_target = fmaxf(tau_min, fminf(tau_target, 1.0f)); + + isv[RL_IQN_ACTION_TAU_INDEX] = tau_target; +} diff --git a/crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu b/crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu new file mode 100644 index 000000000..9b3cc43c5 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu @@ -0,0 +1,70 @@ +// rl_kelly_fraction_controller.cu — Layer 4: half-Kelly position-size multiplier. +// +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Half-Kelly (Thorp) from observed win_rate × R-multiple: +// +// f_kelly = (p × b − q) / b where p = win_rate, q = 1 − p, +// b = avg_win / avg_loss +// f = clamp(safety × f_kelly, 0, 1) where safety = 0.5 (half-Kelly) +// +// Warmup gate (`pearl_first_observation_bootstrap` + warmup discipline): +// hold f at bootstrap = 1.0 until `cumulative_dones >= MIN_TRADES_FOR_RELEASE` +// so the EMA inputs (win_rate, avg_win, avg_loss) have time to converge +// past cold-start noise. After release, f tracks observed edge — if +// edge is negative or zero, Kelly clamps to 0 = no commitment. +// +// Dead-signal guard: if avg_loss <= 0 (sentinel, no losses observed yet) +// hold f at 1.0 — we don't have a denominator for the R-multiple. +// +// Per `feedback_no_atomicadd`: single-thread single-block. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: safety multiplier + warmup gate +// are ISV-driven (slots 680, 681). + +#include + +#define RL_KELLY_FRACTION_INDEX 676 +#define RL_WIN_RATE_EMA_INDEX 677 +#define RL_AVG_WIN_USD_EMA_INDEX 678 +#define RL_AVG_LOSS_USD_EMA_INDEX 679 +#define RL_KELLY_SAFETY_FRAC_INDEX 680 +#define RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX 681 +#define RL_CUMULATIVE_DONES_INDEX 660 + +extern "C" __global__ void rl_kelly_fraction_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + // Warmup gate: hold at bootstrap=1.0 until enough closed trades. + if (isv[RL_CUMULATIVE_DONES_INDEX] < isv[RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX]) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + return; + } + + const float p = isv[RL_WIN_RATE_EMA_INDEX]; + const float avg_win = isv[RL_AVG_WIN_USD_EMA_INDEX]; + const float avg_loss = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + + // Dead-signal guard: need positive loss magnitude for R-multiple. + if (avg_loss <= 0.0f) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + return; + } + if (avg_win <= 0.0f) { + // No wins observed — Kelly cannot make sense of edge, hold at 0 + // (don't commit further size when we've only ever lost). + isv[RL_KELLY_FRACTION_INDEX] = 0.0f; + return; + } + + const float b = avg_win / avg_loss; + const float q = 1.0f - p; + const float f_kelly = (p * b - q) / b; + const float safety = isv[RL_KELLY_SAFETY_FRAC_INDEX]; // typically 0.5 + float f = f_kelly * safety; + f = fmaxf(0.0f, fminf(f, 1.0f)); + isv[RL_KELLY_FRACTION_INDEX] = f; +} diff --git a/crates/ml-alpha/cuda/rl_trail_mutate.cu b/crates/ml-alpha/cuda/rl_trail_mutate.cu index fe1397679..2875ff2b9 100644 --- a/crates/ml-alpha/cuda/rl_trail_mutate.cu +++ b/crates/ml-alpha/cuda/rl_trail_mutate.cu @@ -1,26 +1,29 @@ // rl_trail_mutate.cu — SP20 P5 trail mutation kernel. // -// Handles the previously-dead a7 (TrailTighten) and a8 (TrailLoosen) -// actions per `pearl_dead_trail_stop_actions_a7_a8`. Mutates each -// active unit's trail_distance bounded by ISV-driven MIN/MAX with -// symmetric reciprocal adjust rate per SP20 §4.12: -// a7 (tighten): trail = max(MIN, trail × rate) -// a8 (loosen): trail = min(MAX, trail / rate) +// Handles a7 (TrailTighten) and a8 (TrailLoosen) — closes +// `pearl_dead_trail_stop_actions_a7_a8`. Mutates each active unit's +// trail_distance bounded by ISV-driven MIN/MAX. Each direction has its +// own independent factor per spec 2026-05-30-adaptive-risk-management: // -// Tighten and loosen are reciprocal: N tightens followed by N loosens -// returns trail to its original distance. +// a7 (tighten): trail = max(MIN, trail × TIGHTEN_FACTOR) (default 0.9) +// a8 (loosen): trail = min(MAX, trail × LOOSEN_FACTOR) (default 1.1) // -// Audit history: scripts/audit-wiring.sh dogfood pass flagged that -// a7/a8 had no consumer anywhere in the codebase prior to this -// kernel. +// Tighten and loosen are NO LONGER strictly reciprocal — the agent can +// learn (via the ISV slots' controllers, if any) to bias one direction +// or the other. Bootstrap factors (0.9 / 1.1) preserve the prior +// approximately-symmetric behavior. +// +// Slot migration: previously read `RL_TRAIL_ADJUST_RATE_INDEX = 497` +// and used `trail × rate` for tighten and `trail / rate` for loosen. +// Now reads independent `RL_TRAIL_TIGHTEN_FACTOR_INDEX = 682` and +// `RL_TRAIL_LOOSEN_FACTOR_INDEX = 683` from the risk-management spec. // // One thread per (batch, unit). Mutates ALL active units uniformly -// when the batch chose a7/a8 — matches SP20 §2.2 Tier 2 design -// ("a7/a8 uniformly tighten/loosen all active units"). +// when the batch chose a7/a8 — matches SP20 §2.2 Tier 2 design. // // Per `feedback_no_atomicadd`: per-(batch, unit) independent writes. // Per `feedback_cpu_is_read_only`: pure device-side. -// Per `feedback_isv_for_adaptive_bounds`: MIN/MAX/rate all ISV slots. +// Per `feedback_isv_for_adaptive_bounds`: MIN/MAX/factors all ISV slots. #include @@ -29,7 +32,8 @@ #define ACTION_TRAIL_LOOSEN 8 #define RL_TRAIL_MIN_INDEX 494 #define RL_TRAIL_MAX_INDEX 495 -#define RL_TRAIL_ADJUST_RATE_INDEX 497 +#define RL_TRAIL_TIGHTEN_FACTOR_INDEX 682 +#define RL_TRAIL_LOOSEN_FACTOR_INDEX 683 extern "C" __global__ void rl_trail_mutate( const int* __restrict__ actions, // [B] @@ -48,17 +52,17 @@ extern "C" __global__ void rl_trail_mutate( const int idx = b * MAX_UNITS + u; if (unit_active[idx] == 0) return; - const float trail_min = isv[RL_TRAIL_MIN_INDEX]; - const float trail_max = isv[RL_TRAIL_MAX_INDEX]; - const float adjust_rate = isv[RL_TRAIL_ADJUST_RATE_INDEX]; - const float current = unit_trail_distance[idx]; + const float trail_min = isv[RL_TRAIL_MIN_INDEX]; + const float trail_max = isv[RL_TRAIL_MAX_INDEX]; + const float current = unit_trail_distance[idx]; float updated; if (action == ACTION_TRAIL_TIGHTEN) { - updated = fmaxf(trail_min, current * adjust_rate); + const float factor = isv[RL_TRAIL_TIGHTEN_FACTOR_INDEX]; // bootstrap 0.9 + updated = fmaxf(trail_min, current * factor); } else { // ACTION_TRAIL_LOOSEN - const float inv_rate = (adjust_rate > 1e-6f) ? (1.0f / adjust_rate) : 1.0f; - updated = fminf(trail_max, current * inv_rate); + const float factor = isv[RL_TRAIL_LOOSEN_FACTOR_INDEX]; // bootstrap 1.1 + updated = fminf(trail_max, current * factor); } unit_trail_distance[idx] = updated; } diff --git a/crates/ml-alpha/cuda/rl_win_rate_ema_update.cu b/crates/ml-alpha/cuda/rl_win_rate_ema_update.cu new file mode 100644 index 000000000..b71699409 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_win_rate_ema_update.cu @@ -0,0 +1,63 @@ +// rl_win_rate_ema_update.cu — Layer 4 (Kelly) input EMA producer. +// +// Tracks observed per-step win_rate as an EMA across closed trades. +// Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +// +// Inputs: +// rewards[b] f32 — per-batch realized pnl delta this step (shaped) +// dones[b] f32 — 1.0 if a trade closed this step, else 0.0 +// +// Trade outcome is derived inline: win = (done & reward > 0), +// loss = (done & reward < 0), no-close otherwise. +// +// Output: ISV[RL_WIN_RATE_EMA_INDEX = 677] +// +// Bootstrap semantics per `pearl_first_observation_bootstrap`: +// sentinel = 0.5 (neutral 50% win rate); first non-trivial step replaces +// directly. Zero observations on a given step (no closes) → no update. +// +// Per `feedback_no_atomicadd`: single-thread single-block, sums across +// b_size sequentially. +// Per `feedback_cpu_is_read_only`: pure device kernel; trainer launches +// each step after rl_fused_reward_pipeline writes rewards + dones. +// Per `feedback_isv_for_adaptive_bounds`: EMA-α is currently a structural +// constant 0.05 (≈ 20-trade effective window) — same pattern as +// ema_update_per_step where α is a structural smoothing parameter. + +#include + +#define RL_WIN_RATE_EMA_INDEX 677 +#define EMA_ALPHA 0.05f +#define SENTINEL 0.5f + +extern "C" __global__ void rl_win_rate_ema_update( + float* __restrict__ isv, + const float* __restrict__ rewards, // [b_size] shaped pnl delta + const float* __restrict__ dones, // [b_size] 1.0 = close + int b_size +) { + // Single-thread guard. Block shape (1,1,1) by contract. + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + // Count this step's closed trades and wins (done & reward > 0). + int closed = 0; + int wins = 0; + for (int b = 0; b < b_size; ++b) { + if (dones[b] >= 0.5f) { + closed += 1; + if (rewards[b] > 0.0f) wins += 1; + } + } + if (closed == 0) return; // no observation this step + + const float step_wr = (float)wins / (float)closed; + const float prev = isv[RL_WIN_RATE_EMA_INDEX]; + + // First-observation bootstrap on sentinel 0.5 (neutral default). + if (prev == SENTINEL) { + isv[RL_WIN_RATE_EMA_INDEX] = step_wr; + } else { + isv[RL_WIN_RATE_EMA_INDEX] = (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_wr; + } +} diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index b81590201..b1295b4a6 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -1145,6 +1145,14 @@ fn main() -> Result<()> { // the train-end policy's OOS performance. True pure-eval (forward // only, no backward) is a follow-up architectural change. if cli.n_eval_steps > 0 && !eval_files.is_empty() { + // Adaptive risk management (spec 2026-05-30): reset Layer 1 + // session-level state on train→eval transition. Clears session + // pnl, DD-triggered flag, consec-loss counter, cooldown counter. + // Kelly + inventory EMAs intentionally preserved. + trainer + .reset_session_state() + .context("reset_session_state at train→eval boundary")?; + let head_before_eval = sim .read_total_trade_count() .context("read trade count pre-eval")?; diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index dffdb61c1..16e92b00c 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -762,9 +762,13 @@ pub const RL_TRAIL_MAX_INDEX: usize = 495; /// (Turtle 2N convention). pub const RL_TRAIL_K_INIT_INDEX: usize = 496; -/// Trail-stop adjust rate — `rl_trail_mutate` applies -/// `trail × adjust_rate` on a7 (tighten), `trail / adjust_rate` on -/// a8 (loosen). Symmetric reciprocal per SP20 §4.12. Default 0.9. +/// DEPRECATED — superseded 2026-05-30 by independent +/// `RL_TRAIL_TIGHTEN_FACTOR_INDEX` (682) and +/// `RL_TRAIL_LOOSEN_FACTOR_INDEX` (683) per +/// `docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md`. +/// `rl_trail_mutate` no longer reads this slot. Index 497 remains +/// allocated to preserve subsequent slot numbering (renumbering would +/// be a structural refactor). Do not write to it. pub const RL_TRAIL_ADJUST_RATE_INDEX: usize = 497; // ─── Forward-Return-Distribution (FRD) head — supervised forecaster @@ -1361,4 +1365,93 @@ pub const RL_MIN_TRADES_FOR_RELEASE_INDEX: usize = 661; /// Last RL-allocated slot index (exclusive). Pre-spec: 588. Post-noise-floor: 640. /// Post-clamp-bounds extension: 660. Post-reward-floor-fix: 662. Total new: 74 slots. -pub const RL_SLOTS_END: usize = 662; +// ============================================================ +// Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management-design) +// +// Five orthogonal risk layers, each ISV-driven: +// Layer 1 — CMDP hard constraints (session DD, cooldown, max open, inventory) +// Layer 2 — IQN risk-averse action selection via adaptive τ +// Layer 3 — Inventory penalty (Avellaneda-Stoikov) in reward shaping +// Layer 4 — Kelly-fraction position sizing +// Layer D — TrailTighten/Loosen wiring (close dead-action gap per +// pearl_dead_trail_stop_actions_a7_a8) +// +// Per `feedback_adaptive_not_tuned`: every threshold below is signal-driven +// or warmup-gated, not a hardcoded constant. +// ============================================================ + +// ─── Layer 1: CMDP hard constraints (9 slots) ─── + +/// Layer 1 — Cumulative pnl since session boundary (USD). Reset on +/// train/eval phase transition by the trainer. +pub const RL_SESSION_PNL_USD_INDEX: usize = 662; +/// Layer 1 — Drawdown cap (USD, negative threshold). Bootstrap -3500 = +/// 10% of $35k starting capital per `project_ml_alpha_starting_capital`. +pub const RL_SESSION_DD_LIMIT_USD_INDEX: usize = 663; +/// Layer 1 — Sticky flag (1.0 if session DD breached, 0.0 otherwise). +/// Override-to-Hold in actions_to_market_targets when set. +pub const RL_SESSION_DD_TRIGGERED_INDEX: usize = 664; +/// Layer 1 — Max concurrent open positions per batch. Bootstrap = MAX_UNITS = 4. +pub const RL_MAX_OPEN_UNITS_INDEX: usize = 665; +/// Layer 1 — Max consecutive losses before cool-down. Bootstrap 10. +pub const RL_CONSEC_LOSS_LIMIT_INDEX: usize = 666; +/// Layer 1 — Running consecutive-loss count. Reset on win. +pub const RL_CONSEC_LOSS_COUNT_INDEX: usize = 667; +/// Layer 1 — Cool-down period remaining (steps). When > 0, override to Hold. +pub const RL_COOLDOWN_REMAINING_STEPS_INDEX: usize = 668; +/// Layer 1 — Default cool-down length (steps). Set when consec loss limit hit. +pub const RL_COOLDOWN_DURATION_INDEX: usize = 669; +/// Layer 1 — Cap on net one-sided exposure (USD). Bootstrap $10k. +pub const RL_NET_INVENTORY_LIMIT_USD_INDEX: usize = 670; + +// ─── Layer 2: IQN risk-averse τ for action selection (3 slots) ─── + +/// Layer 2 — Quantile τ ∈ (0, 1] used for action selection. +/// Bootstrap 0.5 (median = approximately current expected-Q behavior). +/// Adaptive: drops toward MIN under session drawdown. +pub const RL_IQN_ACTION_TAU_INDEX: usize = 671; +/// Layer 2 — Lower bound on action τ. Bootstrap 0.1 (pessimistic floor). +pub const RL_IQN_ACTION_TAU_MIN_INDEX: usize = 672; +/// Layer 2 — How aggressively τ drops with session drawdown fraction. +/// `τ = max(τ_min, 0.5 - sensitivity × drawdown_frac)`. Bootstrap 5.0 +/// (= τ hits MIN at 10% drawdown of starting capital). +pub const RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX: usize = 673; + +// ─── Layer 3: Inventory penalty (Avellaneda-Stoikov style) (2 slots) ─── + +/// Layer 3 — Penalty coefficient β. `reward_shaped = reward_raw - β × |net_position|`. +/// Bootstrap 0.0 (sentinel — adaptive on first inventory variance observation). +pub const RL_INVENTORY_PENALTY_BETA_INDEX: usize = 674; +/// Layer 3 — Observed inventory variance EMA (Welford-style). +/// β controller scales penalty to ~1% of reward magnitude at 2σ inventory. +pub const RL_INVENTORY_VARIANCE_EMA_INDEX: usize = 675; + +// ─── Layer 4: Kelly-fraction position sizing (6 slots) ─── + +/// Layer 4 — Output: position-size multiplier ∈ [0, 1]. +/// Bootstrap 1.0 (full size). Adaptive from observed win_rate × R-multiple. +pub const RL_KELLY_FRACTION_INDEX: usize = 676; +/// Layer 4 — Observed win_rate EMA. Bootstrap 0.5 (neutral sentinel). +pub const RL_WIN_RATE_EMA_INDEX: usize = 677; +/// Layer 4 — Observed avg_win EMA in USD. Bootstrap 0.0 sentinel. +pub const RL_AVG_WIN_USD_EMA_INDEX: usize = 678; +/// Layer 4 — Observed avg_loss EMA in USD. Bootstrap 0.0 sentinel. +pub const RL_AVG_LOSS_USD_EMA_INDEX: usize = 679; +/// Layer 4 — Safety multiplier on Kelly. Bootstrap 0.5 (Thorp half-Kelly). +pub const RL_KELLY_SAFETY_FRAC_INDEX: usize = 680; +/// Layer 4 — Warmup gate: hold at bootstrap=1.0 until N closed trades. +/// Uses `RL_CUMULATIVE_DONES_INDEX` from b1ef6664a as the counter signal. +pub const RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX: usize = 681; + +// ─── Layer D: TrailTighten/Loosen wiring (2 slots) ─── + +/// Layer D — Factor applied to unit_trail_distance on action a7 (TrailTighten). +/// Bootstrap 0.9 (tighten by 10%). Closes pearl_dead_trail_stop_actions_a7_a8. +pub const RL_TRAIL_TIGHTEN_FACTOR_INDEX: usize = 682; +/// Layer D — Factor applied to unit_trail_distance on action a8 (TrailLoosen). +/// Bootstrap 1.1 (loosen by 10%). +pub const RL_TRAIL_LOOSEN_FACTOR_INDEX: usize = 683; + +/// Last RL-allocated slot index (exclusive). Pre-risk-stack: 662. +/// Post-risk-stack: 684. Total new in risk-stack: 22 slots, +88 bytes mapped-pinned. +pub const RL_SLOTS_END: usize = 684; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 6e8663359..85ff2570e 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -203,6 +203,23 @@ const RL_ADVANTAGE_NORMALIZE_CUBIN: &[u8] = /// single-block. Caller passes ISV bus + input/count/mean/M² slot indices. const RL_SIGNAL_VARIANCE_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin")); +// Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management). +// 4 controllers + 3 EMA producers feeding Layers 1-4. Layer D (trail +// tighten/loosen) is read-only ISV consumed by the existing rl_trail_mutate. +const RL_CMDP_CONSTRAINTS_CHECK_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_cmdp_constraints_check.cubin")); +const RL_IQN_ACTION_TAU_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_iqn_action_tau_controller.cubin")); +const RL_INVENTORY_BETA_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_inventory_beta_controller.cubin")); +const RL_KELLY_FRACTION_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_kelly_fraction_controller.cubin")); +const RL_WIN_RATE_EMA_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_win_rate_ema_update.cubin")); +const RL_AVG_WIN_LOSS_EMA_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_avg_win_loss_ema_update.cubin")); +const RL_INVENTORY_VARIANCE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_inventory_variance_update.cubin")); const RL_Q_PI_DISTILL_GRAD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_q_pi_distill_grad.cubin")); // λ_distill adaptive controller (rljzl followup 2026-05-24). @@ -648,6 +665,26 @@ pub struct IntegratedTrainer { // before controller consumers. _rl_signal_variance_update_module: Arc, rl_signal_variance_update_fn: CudaFunction, + // Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management). + // CMDP gates run as a standalone kernel per step (reads pnl + outcomes). + // The 3 layer-2/3/4 controllers are also fused into rl_fused_controllers + // for graph-capture efficiency, but the standalone modules are loaded + // so unit tests can drive each layer in isolation. + _rl_cmdp_constraints_check_module: Arc, + rl_cmdp_constraints_check_fn: CudaFunction, + _rl_iqn_action_tau_controller_module: Arc, + rl_iqn_action_tau_controller_fn: CudaFunction, + _rl_inventory_beta_controller_module: Arc, + rl_inventory_beta_controller_fn: CudaFunction, + _rl_kelly_fraction_controller_module: Arc, + rl_kelly_fraction_controller_fn: CudaFunction, + // EMA producers feeding the new controllers. + _rl_win_rate_ema_update_module: Arc, + rl_win_rate_ema_update_fn: CudaFunction, + _rl_avg_win_loss_ema_update_module: Arc, + rl_avg_win_loss_ema_update_fn: CudaFunction, + _rl_inventory_variance_update_module: Arc, + rl_inventory_variance_update_fn: CudaFunction, /// Phase 4.4 blended V baselines fed to compute_advantage_return. pub v_blended_d: CudaSlice, pub v_blended_tp1_d: CudaSlice, @@ -1523,6 +1560,49 @@ impl IntegratedTrainer { let rl_signal_variance_update_fn = rl_signal_variance_update_module .load_function("rl_signal_variance_update") .context("load rl_signal_variance_update")?; + // Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management). + let rl_cmdp_constraints_check_module = ctx + .load_cubin(RL_CMDP_CONSTRAINTS_CHECK_CUBIN.to_vec()) + .context("load rl_cmdp_constraints_check cubin")?; + let rl_cmdp_constraints_check_fn = rl_cmdp_constraints_check_module + .load_function("rl_cmdp_constraints_check") + .context("load rl_cmdp_constraints_check")?; + let rl_iqn_action_tau_controller_module = ctx + .load_cubin(RL_IQN_ACTION_TAU_CONTROLLER_CUBIN.to_vec()) + .context("load rl_iqn_action_tau_controller cubin")?; + let rl_iqn_action_tau_controller_fn = rl_iqn_action_tau_controller_module + .load_function("rl_iqn_action_tau_controller") + .context("load rl_iqn_action_tau_controller")?; + let rl_inventory_beta_controller_module = ctx + .load_cubin(RL_INVENTORY_BETA_CONTROLLER_CUBIN.to_vec()) + .context("load rl_inventory_beta_controller cubin")?; + let rl_inventory_beta_controller_fn = rl_inventory_beta_controller_module + .load_function("rl_inventory_beta_controller") + .context("load rl_inventory_beta_controller")?; + let rl_kelly_fraction_controller_module = ctx + .load_cubin(RL_KELLY_FRACTION_CONTROLLER_CUBIN.to_vec()) + .context("load rl_kelly_fraction_controller cubin")?; + let rl_kelly_fraction_controller_fn = rl_kelly_fraction_controller_module + .load_function("rl_kelly_fraction_controller") + .context("load rl_kelly_fraction_controller")?; + let rl_win_rate_ema_update_module = ctx + .load_cubin(RL_WIN_RATE_EMA_UPDATE_CUBIN.to_vec()) + .context("load rl_win_rate_ema_update cubin")?; + let rl_win_rate_ema_update_fn = rl_win_rate_ema_update_module + .load_function("rl_win_rate_ema_update") + .context("load rl_win_rate_ema_update")?; + let rl_avg_win_loss_ema_update_module = ctx + .load_cubin(RL_AVG_WIN_LOSS_EMA_UPDATE_CUBIN.to_vec()) + .context("load rl_avg_win_loss_ema_update cubin")?; + let rl_avg_win_loss_ema_update_fn = rl_avg_win_loss_ema_update_module + .load_function("rl_avg_win_loss_ema_update") + .context("load rl_avg_win_loss_ema_update")?; + let rl_inventory_variance_update_module = ctx + .load_cubin(RL_INVENTORY_VARIANCE_UPDATE_CUBIN.to_vec()) + .context("load rl_inventory_variance_update cubin")?; + let rl_inventory_variance_update_fn = rl_inventory_variance_update_module + .load_function("rl_inventory_variance_update") + .context("load rl_inventory_variance_update")?; let rl_q_distill_lambda_controller_module = ctx .load_cubin(RL_Q_DISTILL_LAMBDA_CONTROLLER_CUBIN.to_vec()) .context("load rl_q_distill_lambda_controller cubin")?; @@ -2557,6 +2637,21 @@ impl IntegratedTrainer { rl_advantage_normalize_fn, _rl_signal_variance_update_module: rl_signal_variance_update_module, rl_signal_variance_update_fn, + // Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management). + _rl_cmdp_constraints_check_module: rl_cmdp_constraints_check_module, + rl_cmdp_constraints_check_fn, + _rl_iqn_action_tau_controller_module: rl_iqn_action_tau_controller_module, + rl_iqn_action_tau_controller_fn, + _rl_inventory_beta_controller_module: rl_inventory_beta_controller_module, + rl_inventory_beta_controller_fn, + _rl_kelly_fraction_controller_module: rl_kelly_fraction_controller_module, + rl_kelly_fraction_controller_fn, + _rl_win_rate_ema_update_module: rl_win_rate_ema_update_module, + rl_win_rate_ema_update_fn, + _rl_avg_win_loss_ema_update_module: rl_avg_win_loss_ema_update_module, + rl_avg_win_loss_ema_update_fn, + _rl_inventory_variance_update_module: rl_inventory_variance_update_module, + rl_inventory_variance_update_fn, v_blended_d, v_blended_tp1_d, _rl_q_distill_lambda_controller_module: rl_q_distill_lambda_controller_module, @@ -2830,6 +2925,37 @@ impl IntegratedTrainer { .with_controllers_bootstrapped()?) } + /// Adaptive risk management — session boundary reset (spec 2026-05-30). + /// Called by the caller at train→eval (or fold) transition. Clears + /// the Layer 1 session-level state so the eval phase starts with + /// fresh DD tracking + cooldown counters. Kelly + inventory EMAs + /// are intentionally NOT reset: training-acquired statistics should + /// carry into eval (matches "policy learned during train, applied + /// in eval" intent). + pub fn reset_session_state(&self) -> Result<()> { + for (slot, value) in [ + (crate::rl::isv_slots::RL_SESSION_PNL_USD_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_SESSION_DD_TRIGGERED_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_CONSEC_LOSS_COUNT_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0_f32), + ] { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(slot as i32); + args.push_f32(value); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_isv_write_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_isv_write (reset_session_state slot {slot}): {:?}", e))?; + } + } + Ok(()) + } + /// Read a single ISV slot from mapped-pinned host memory. Volatile /// read prevents the compiler from caching stale values. The GPU /// writes via `dev_ptr`; by the time the host reads (after stream @@ -2907,7 +3033,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 147] = [ + let isv_constants: [(usize, f32); 168] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -2946,7 +3072,10 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_TRAIL_MIN_INDEX, 0.001), (crate::rl::isv_slots::RL_TRAIL_MAX_INDEX, 100.0), (crate::rl::isv_slots::RL_TRAIL_K_INIT_INDEX, 2.0), - (crate::rl::isv_slots::RL_TRAIL_ADJUST_RATE_INDEX, 0.9), + // RL_TRAIL_ADJUST_RATE_INDEX (497) deprecated 2026-05-30 — + // superseded by independent slots 682/683 (trail tighten/loosen + // factors) bootstrapped below. Slot 497 stays allocated for + // numbering compatibility but is no longer written. // SP20 P3 FRD head seeds — λ, LR, h1-3 tick offsets, ±σ atom range. // Per feedback_isv_for_adaptive_bounds + single-source-of-truth: // horizon ticks + bucket-range σ derive from the canonical @@ -3154,6 +3283,34 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX, 0.998), // Shared Wiener-α floor — read by 9 controllers (slot 659). (crate::rl::isv_slots::RL_WIENER_ALPHA_FLOOR_INDEX, 0.4), + // ─── Adaptive risk management — 22 new slots (spec 2026-05-30-adaptive-risk-management) ─── + // Layer 1 (CMDP hard constraints): session DD + cooldown + caps. + (crate::rl::isv_slots::RL_SESSION_PNL_USD_INDEX, 0.0), // session boundary reset + (crate::rl::isv_slots::RL_SESSION_DD_LIMIT_USD_INDEX, -3500.0), // 10% of $35k starting capital + (crate::rl::isv_slots::RL_SESSION_DD_TRIGGERED_INDEX, 0.0), + (crate::rl::isv_slots::RL_MAX_OPEN_UNITS_INDEX, 4.0), // = current MAX_UNITS + (crate::rl::isv_slots::RL_CONSEC_LOSS_LIMIT_INDEX, 10.0), + (crate::rl::isv_slots::RL_CONSEC_LOSS_COUNT_INDEX, 0.0), + (crate::rl::isv_slots::RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0), + (crate::rl::isv_slots::RL_COOLDOWN_DURATION_INDEX, 500.0), + (crate::rl::isv_slots::RL_NET_INVENTORY_LIMIT_USD_INDEX, 10000.0), + // Layer 2 (IQN risk-averse τ): median bootstrap. + (crate::rl::isv_slots::RL_IQN_ACTION_TAU_INDEX, 0.5), // median ≈ current behavior + (crate::rl::isv_slots::RL_IQN_ACTION_TAU_MIN_INDEX, 0.1), + (crate::rl::isv_slots::RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX, 5.0), // hits MIN at 10% DD + // Layer 3 (inventory penalty): sentinel-zero until variance EMA accumulates. + (crate::rl::isv_slots::RL_INVENTORY_PENALTY_BETA_INDEX, 0.0), // sentinel + (crate::rl::isv_slots::RL_INVENTORY_VARIANCE_EMA_INDEX, 0.0), // sentinel + // Layer 4 (Kelly fraction sizing): full size pre-warmup, half-Kelly post. + (crate::rl::isv_slots::RL_KELLY_FRACTION_INDEX, 1.0), // full size pre-warmup + (crate::rl::isv_slots::RL_WIN_RATE_EMA_INDEX, 0.5), // neutral sentinel + (crate::rl::isv_slots::RL_AVG_WIN_USD_EMA_INDEX, 0.0), // sentinel + (crate::rl::isv_slots::RL_AVG_LOSS_USD_EMA_INDEX, 0.0), // sentinel + (crate::rl::isv_slots::RL_KELLY_SAFETY_FRAC_INDEX, 0.5), // half-Kelly (Thorp) + (crate::rl::isv_slots::RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX, 1000.0), + // Layer D (TrailTighten/Loosen factors). + (crate::rl::isv_slots::RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9), + (crate::rl::isv_slots::RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -3409,6 +3566,165 @@ impl IntegratedTrainer { Ok(()) } + /// Adaptive risk management — Layer 1 (spec 2026-05-30-adaptive-risk-management). + /// Maintains session pnl + DD-triggered flag + cooldown counter + + /// consecutive-loss counter. Single-thread single-block; reads + /// `rewards_d` and `dones_d` (`[b_size]` f32 device buffers populated + /// by `rl_fused_reward_pipeline`). Outcome is derived inline (done & + /// reward sign). Must run AFTER the reward pipeline and BEFORE + /// `actions_to_market_targets` reads the override flags. + pub fn launch_rl_cmdp_constraints_check( + &self, + rewards_d: &CudaSlice, + dones_d: &CudaSlice, + b_size: usize, + ) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_ptr(rewards_d.raw_ptr()); + args.push_ptr(dones_d.raw_ptr()); + args.push_i32(b_size as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_cmdp_constraints_check_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_cmdp_constraints_check: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 2 IQN τ controller (spec 2026-05-30). + /// Standalone launcher. Production calls go through `rl_fused_controllers`; + /// this helper is retained for unit-test isolation. + pub fn launch_rl_iqn_action_tau_controller(&self) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_iqn_action_tau_controller_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_iqn_action_tau_controller: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 3 inventory β controller (spec 2026-05-30). + /// Standalone launcher; production calls go through `rl_fused_controllers`. + pub fn launch_rl_inventory_beta_controller(&self) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_inventory_beta_controller_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_inventory_beta_controller: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 4 Kelly fraction controller (spec 2026-05-30). + /// Standalone launcher; production calls go through `rl_fused_controllers`. + pub fn launch_rl_kelly_fraction_controller(&self) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_kelly_fraction_controller_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_kelly_fraction_controller: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 4 input EMA producer (win_rate). + /// Reads `rewards_d`, `dones_d` [b_size] f32; outcome derived inline. + /// Writes ISV[RL_WIN_RATE_EMA_INDEX = 677]. + pub fn launch_rl_win_rate_ema_update( + &self, + rewards_d: &CudaSlice, + dones_d: &CudaSlice, + b_size: usize, + ) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_ptr(rewards_d.raw_ptr()); + args.push_ptr(dones_d.raw_ptr()); + args.push_i32(b_size as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_win_rate_ema_update_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_win_rate_ema_update: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 4 input EMA producer (avg_win/loss USD). + /// Reads `rewards_d`, `dones_d` [b_size] f32; outcome derived inline. + /// Writes ISV[RL_AVG_WIN_USD_EMA_INDEX = 678, RL_AVG_LOSS_USD_EMA_INDEX = 679]. + pub fn launch_rl_avg_win_loss_ema_update( + &self, + rewards_d: &CudaSlice, + dones_d: &CudaSlice, + b_size: usize, + ) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_ptr(rewards_d.raw_ptr()); + args.push_ptr(dones_d.raw_ptr()); + args.push_i32(b_size as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_avg_win_loss_ema_update_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_avg_win_loss_ema_update: {:?}", e))?; + } + Ok(()) + } + + /// Adaptive risk management — Layer 3 input EMA producer (inventory variance). + /// Reads `net_position_per_batch_d` [b_size] i32 of signed contracts + /// (typically the trainer's `prev_position_lots_d`); writes + /// ISV[RL_INVENTORY_VARIANCE_EMA_INDEX = 675]. + pub fn launch_rl_inventory_variance_update( + &self, + net_position_per_batch_d: &CudaSlice, + b_size: usize, + ) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_ptr(net_position_per_batch_d.raw_ptr()); + args.push_i32(b_size as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_inventory_variance_update_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_inventory_variance_update: {:?}", e))?; + } + Ok(()) + } + /// Phase R3: launch `ema_update_on_done` — done-gated EMA producer /// that writes the per-batch closed-trade mean into `isv[slot_index]`. /// `alpha` is the Wiener-α (caller must pre-floor at 0.4 per @@ -6278,6 +6594,46 @@ impl IntegratedTrainer { } } + // ──────────────────────────────────────────────────────────── + // Adaptive risk management — spec 2026-05-30-adaptive-risk-management. + // + // Layer 1 (CMDP gates): track session pnl + DD-triggered flag + + // cooldown counter + consecutive-loss counter from this step's + // rewards/dones. Must precede the Layer 2 IQN-τ controller because + // τ derives from `session_pnl` written here. + // + // EMA producers: feed Layer 4 Kelly (win_rate, avg_win/loss) and + // Layer 3 (inventory variance). All three are single-thread + // single-block kernels. + // + // The Layer 2/3/4 controllers themselves are folded into + // `rl_fused_controllers` below, which fires after the existing + // Welford updates — so this is the producer side only. + // ──────────────────────────────────────────────────────────── + self.launch_rl_cmdp_constraints_check( + &self.rewards_d, + &self.dones_d, + b_size, + ) + .context("launch_rl_cmdp_constraints_check")?; + self.launch_rl_win_rate_ema_update( + &self.rewards_d, + &self.dones_d, + b_size, + ) + .context("launch_rl_win_rate_ema_update")?; + self.launch_rl_avg_win_loss_ema_update( + &self.rewards_d, + &self.dones_d, + b_size, + ) + .context("launch_rl_avg_win_loss_ema_update")?; + self.launch_rl_inventory_variance_update( + &self.prev_position_lots_d, + b_size, + ) + .context("launch_rl_inventory_variance_update")?; + // Scale + clamp rewards, write per-step max to ISV for the // adaptive clamp controller. Without this, raw PnL ($100s) // goes into C51 with atom span ±0.5 → binary Q resolution. diff --git a/crates/ml-alpha/tests/risk_stack_invariants.rs b/crates/ml-alpha/tests/risk_stack_invariants.rs new file mode 100644 index 000000000..616dfddc2 --- /dev/null +++ b/crates/ml-alpha/tests/risk_stack_invariants.rs @@ -0,0 +1,393 @@ +//! Risk-stack invariants (spec `docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md`). +//! +//! 12 GPU-oracle gates — four per active layer. +//! Layer 1 (CMDP): G1-G4 +//! Layer 2 (IQN τ): G5-G7 +//! Layer 3 (Inventory): G8-G9 +//! Layer 4 (Kelly): G10-G12 +//! +//! Per [[feedback_no_cpu_test_fallbacks]] every oracle is an analytical +//! identity of the kernel's documented math — no host reimplementation. +//! +//! Run: +//! `cargo test -p ml-alpha --test risk_stack_invariants --release -- --ignored --nocapture` + +use anyhow::Result; +use cudarc::driver::{CudaSlice, CudaStream}; +use ml_alpha::rl::isv_slots::*; +use ml_alpha::trainer::integrated::{ + write_slice_f32_d_pub, IntegratedTrainer, IntegratedTrainerConfig, +}; +use ml_alpha::trainer::perception::PerceptionTrainerConfig; +use ml_core::device::MlDevice; +use std::sync::Arc; + +// ─── Bootstrap defaults (mirrored from `with_controllers_bootstrapped`) ─── +const SESSION_DD_LIMIT_USD: f32 = -3500.0; // 10% of $35k starting capital +const COOLDOWN_DURATION: f32 = 500.0; +const IQN_TAU_MIN: f32 = 0.1; +const IQN_TAU_DD_SENSITIVITY: f32 = 5.0; // hits MIN at 10% DD +const KELLY_SAFETY_FRAC: f32 = 0.5; // half-Kelly (Thorp) +const KELLY_MIN_TRADES: f32 = 1000.0; +const STARTING_CAPITAL_USD: f32 = 35000.0; // ES single-contract base + +fn build_trainer() -> Option<(MlDevice, IntegratedTrainer)> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + return None; + } + }; + let cfg = IntegratedTrainerConfig { + perception: PerceptionTrainerConfig { + seq_len: 4, + n_batch: 1, + ..PerceptionTrainerConfig::default() + }, + dqn_seed: 0xC1, + ppo_seed: 0xC2, + ..IntegratedTrainerConfig::default() + }; + let trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); + Some((dev, trainer)) +} + +fn sync(trainer: &IntegratedTrainer) { + trainer.stream.synchronize().expect("sync trainer stream"); +} + +/// Mapped-pinned upload per [[feedback_no_htod_htoh_only_mapped_pinned]]: +/// allocate the device buffer, then stage through the canonical public +/// helper which routes through `MappedF32Buffer` + DtoD. +fn upload_f32(stream: &Arc, host: &[f32]) -> Result> { + let mut d = stream.alloc_zeros::(host.len())?; + write_slice_f32_d_pub(stream, host, &mut d)?; + Ok(d) +} + +// ═════════════════════════════════════════════════════════════════════ +// LAYER 1 — CMDP hard constraints +// ═════════════════════════════════════════════════════════════════════ + +// G1 — DD breaker sets sticky triggered flag when session_pnl crosses +// dd_limit. Single batch with one large loss drives session_pnl below +// the limit; one launch flips RL_SESSION_DD_TRIGGERED_INDEX to 1.0. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g1_cmdp_dd_breaker_triggers_at_limit() -> Result<()> { + let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + + // Confirm bootstrap seeding. + sync(&trainer); + assert_eq!(trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX), 0.0); + assert_eq!(trainer.read_isv_host(RL_SESSION_DD_TRIGGERED_INDEX), 0.0); + let dd_limit = trainer.read_isv_host(RL_SESSION_DD_LIMIT_USD_INDEX); + assert!((dd_limit - SESSION_DD_LIMIT_USD).abs() < 1e-3); + + // Single batch step delivering pnl = -4000 (below -3500 limit). + let rewards_d = upload_f32(&stream, &[-4000.0])?; + let dones_d = upload_f32(&stream, &[1.0])?; + trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?; + sync(&trainer); + + assert!( + (trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX) - (-4000.0)).abs() < 1e-3, + "session_pnl accumulated: expected -4000, got {}", + trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX) + ); + assert_eq!( + trainer.read_isv_host(RL_SESSION_DD_TRIGGERED_INDEX), 1.0, + "dd_triggered must flip to 1.0 when session_pnl < dd_limit" + ); + eprintln!("G1 OK — DD breaker tripped at session_pnl = -4000 < dd_limit = -3500"); + Ok(()) +} + +// G2 — Cooldown starts when consec_loss_count crosses limit. Feed 10 +// done-losses in a single launch (b_size=10) and verify cooldown_remaining +// snaps to cooldown_duration while consec resets to 0. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g2_cmdp_cooldown_starts_after_consec_loss_limit() -> Result<()> { + let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + + // Pre-condition: consec=0, cooldown_remaining=0, limit=10, duration=500. + sync(&trainer); + assert_eq!(trainer.read_isv_host(RL_CONSEC_LOSS_COUNT_INDEX), 0.0); + assert_eq!(trainer.read_isv_host(RL_COOLDOWN_REMAINING_STEPS_INDEX), 0.0); + + // 10 done-losses (one per batch element). DD limit not breached + // (per-loss ≈ -10, total -100 ≫ -3500), so we isolate the consec arm. + let rewards_d = upload_f32(&stream, &[-10.0_f32; 10])?; + let dones_d = upload_f32(&stream, &[1.0_f32; 10])?; + trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 10)?; + sync(&trainer); + + let cooldown = trainer.read_isv_host(RL_COOLDOWN_REMAINING_STEPS_INDEX); + let consec = trainer.read_isv_host(RL_CONSEC_LOSS_COUNT_INDEX); + assert!( + (cooldown - COOLDOWN_DURATION).abs() < 1e-3, + "cooldown_remaining must snap to duration={COOLDOWN_DURATION}; got {cooldown}" + ); + assert_eq!(consec, 0.0, "consec_loss_count must reset on limit hit"); + eprintln!( + "G2 OK — 10 losses in single step → cooldown set to {cooldown}, consec reset to {consec}" + ); + Ok(()) +} + +// G3 — Cooldown decrements one step per launch (independent of consec arm). +// We bypass the consec gate by sending one win plus a pre-seeded +// cooldown_remaining and verify exactly one decrement per launch. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g3_cmdp_cooldown_decrements_per_step() -> Result<()> { + let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + + trainer.isv_mapped.write_record(RL_COOLDOWN_REMAINING_STEPS_INDEX, 50.0); + + // Three idle launches with a single break-even close (r=0, done=1) — + // neither increments nor resets consec, just exercises the decrement. + let rewards_d = upload_f32(&stream, &[0.0])?; + let dones_d = upload_f32(&stream, &[1.0])?; + for _ in 0..3 { + trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?; + } + sync(&trainer); + + let cooldown = trainer.read_isv_host(RL_COOLDOWN_REMAINING_STEPS_INDEX); + assert!( + (cooldown - 47.0).abs() < 1e-3, + "cooldown must decrement 50→47 over 3 launches; got {cooldown}" + ); + eprintln!("G3 OK — cooldown decremented 50 → {cooldown} over 3 launches"); + Ok(()) +} + +// G4 — A done-win resets consec_loss_count to 0 (canonical streak break). +// Seed consec=5, feed one win → counter must reset before the limit check. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g4_cmdp_win_resets_consec_loss_counter() -> Result<()> { + let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + + trainer.isv_mapped.write_record(RL_CONSEC_LOSS_COUNT_INDEX, 5.0); + let rewards_d = upload_f32(&stream, &[25.0])?; + let dones_d = upload_f32(&stream, &[1.0])?; + trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?; + sync(&trainer); + + let consec = trainer.read_isv_host(RL_CONSEC_LOSS_COUNT_INDEX); + assert_eq!(consec, 0.0, "win must reset consec_loss_count to 0; got {consec}"); + eprintln!("G4 OK — done-win reset consec_loss_count 5 → 0"); + Ok(()) +} + +// ═════════════════════════════════════════════════════════════════════ +// LAYER 2 — IQN risk-averse action τ +// ═════════════════════════════════════════════════════════════════════ + +// G5 — At positive session_pnl τ stays at 0.5 (median/risk-neutral). +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g5_iqn_tau_at_peak_returns_0_5() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + trainer.isv_mapped.write_record(RL_SESSION_PNL_USD_INDEX, 1000.0); + trainer.launch_rl_iqn_action_tau_controller()?; + sync(&trainer); + + let tau = trainer.read_isv_host(RL_IQN_ACTION_TAU_INDEX); + assert!( + (tau - 0.5).abs() < 1e-5, + "τ must be 0.5 when session_pnl > 0 (no drawdown); got {tau}" + ); + eprintln!("G5 OK — τ = {tau} at session_pnl = +1000"); + Ok(()) +} + +// G6 — Small drawdown linearly reduces τ. With sensitivity=5.0, +// session_pnl = -350 → drawdown_frac = 0.01 → τ_target = 0.5 - 5×0.01 = 0.45. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g6_iqn_tau_adapts_linearly_to_drawdown() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + trainer.isv_mapped.write_record(RL_SESSION_PNL_USD_INDEX, -350.0); + trainer.launch_rl_iqn_action_tau_controller()?; + sync(&trainer); + + let tau = trainer.read_isv_host(RL_IQN_ACTION_TAU_INDEX); + let drawdown_frac = 350.0 / STARTING_CAPITAL_USD; + let expected = 0.5 - IQN_TAU_DD_SENSITIVITY * drawdown_frac; + assert!( + (tau - expected).abs() < 1e-4, + "τ must be 0.5 - sensitivity × dd_frac = {expected}; got {tau}" + ); + eprintln!("G6 OK — τ = {tau} at session_pnl = -350 (≈ 1% DD)"); + Ok(()) +} + +// G7 — Deep drawdown clamps τ to τ_min. session_pnl = -7000 (20% of $35k) +// drives τ_target = 0.5 - 5×0.2 = -0.5, which must clamp to τ_min = 0.1. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g7_iqn_tau_clamps_at_min_on_deep_drawdown() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + trainer.isv_mapped.write_record(RL_SESSION_PNL_USD_INDEX, -7000.0); + trainer.launch_rl_iqn_action_tau_controller()?; + sync(&trainer); + + let tau = trainer.read_isv_host(RL_IQN_ACTION_TAU_INDEX); + assert!( + (tau - IQN_TAU_MIN).abs() < 1e-5, + "τ must clamp at τ_min={IQN_TAU_MIN} for 20% DD; got {tau}" + ); + eprintln!("G7 OK — τ clamped at {tau} = τ_min for 20% drawdown"); + Ok(()) +} + +// ═════════════════════════════════════════════════════════════════════ +// LAYER 3 — Avellaneda-Stoikov inventory penalty β +// ═════════════════════════════════════════════════════════════════════ + +// G8 — β bootstraps to target on first observation (prev==0 sentinel). +// reward_mag=20, inv_var=4 ⇒ inv_std=2 ⇒ target = 0.01×20 / (2×2) = 0.05. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g8_inventory_beta_bootstraps_to_target() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + // Sentinel: β starts at 0.0 per bootstrap. + assert_eq!(trainer.read_isv_host(RL_INVENTORY_PENALTY_BETA_INDEX), 0.0); + + trainer.isv_mapped.write_record(RL_REWARD_MAGNITUDE_EMA_INDEX, 20.0); + trainer.isv_mapped.write_record(RL_INVENTORY_VARIANCE_EMA_INDEX, 4.0); + trainer.launch_rl_inventory_beta_controller()?; + sync(&trainer); + + let beta = trainer.read_isv_host(RL_INVENTORY_PENALTY_BETA_INDEX); + // expected = 0.01 × reward_mag / (2 × inv_std) = 0.01 × 20 / (2 × 2) = 0.05 + assert!( + (beta - 0.05).abs() < 1e-5, + "β must bootstrap-replace to target=0.05; got {beta}" + ); + eprintln!("G8 OK — β bootstrapped to {beta} = 0.01·20/(2·√4)"); + Ok(()) +} + +// G9 — After bootstrap, β Wiener-blends with the new target instead of +// snapping. α_floor=0.4 ⇒ next β = 0.6×prev + 0.4×new_target. +// Verifies the canonical [[pearl_first_observation_bootstrap]] discipline. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g9_inventory_beta_wiener_blends_after_bootstrap() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + // Step 1: bootstrap to 0.05. + trainer.isv_mapped.write_record(RL_REWARD_MAGNITUDE_EMA_INDEX, 20.0); + trainer.isv_mapped.write_record(RL_INVENTORY_VARIANCE_EMA_INDEX, 4.0); + trainer.launch_rl_inventory_beta_controller()?; + sync(&trainer); + let beta1 = trainer.read_isv_host(RL_INVENTORY_PENALTY_BETA_INDEX); + assert!((beta1 - 0.05).abs() < 1e-5); + + // Step 2: target shifts to 0.10 (double reward_mag). + // Wiener: β_new = 0.6 × 0.05 + 0.4 × 0.10 = 0.07. + trainer.isv_mapped.write_record(RL_REWARD_MAGNITUDE_EMA_INDEX, 40.0); + trainer.launch_rl_inventory_beta_controller()?; + sync(&trainer); + let beta2 = trainer.read_isv_host(RL_INVENTORY_PENALTY_BETA_INDEX); + assert!( + (beta2 - 0.07).abs() < 1e-5, + "β must Wiener-blend (α=0.4) to 0.6·0.05 + 0.4·0.10 = 0.07; got {beta2}" + ); + eprintln!("G9 OK — β Wiener-blended 0.05 → {beta2} (target 0.10)"); + Ok(()) +} + +// ═════════════════════════════════════════════════════════════════════ +// LAYER 4 — Kelly fraction position sizing +// ═════════════════════════════════════════════════════════════════════ + +// G10 — At p=0.55, R-multiple b=avg_win/avg_loss=1.8, half-Kelly returns +// 0.5 × (0.55×1.8 − 0.45)/1.8 = 0.5 × 0.30 = 0.15. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g10_kelly_at_known_edge_returns_half_kelly() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + // Past warmup so the gate releases. + trainer.isv_mapped.write_record(RL_CUMULATIVE_DONES_INDEX, 2000.0); + trainer.isv_mapped.write_record(RL_WIN_RATE_EMA_INDEX, 0.55); + trainer.isv_mapped.write_record(RL_AVG_WIN_USD_EMA_INDEX, 1.8); + trainer.isv_mapped.write_record(RL_AVG_LOSS_USD_EMA_INDEX, 1.0); + trainer.launch_rl_kelly_fraction_controller()?; + sync(&trainer); + + let f = trainer.read_isv_host(RL_KELLY_FRACTION_INDEX); + // f_kelly = (0.55·1.8 − 0.45)/1.8 = 0.3 → f = safety·f_kelly = 0.5·0.3 = 0.15 + let expected = KELLY_SAFETY_FRAC * 0.3; + assert!( + (f - expected).abs() < 1e-5, + "Kelly at p=0.55, R=1.8 must be {expected}; got {f}" + ); + eprintln!("G10 OK — Kelly = {f} at p=0.55, R-multiple=1.8 (half-Kelly target {expected})"); + Ok(()) +} + +// G11 — At negative edge (p=0.21, b=1.8) the analytic Kelly is negative; +// the kernel clamps f to 0 — never short-the-edge / lever-up-into-losses. +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g11_kelly_at_negative_edge_clamps_to_zero() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + trainer.isv_mapped.write_record(RL_CUMULATIVE_DONES_INDEX, 2000.0); + trainer.isv_mapped.write_record(RL_WIN_RATE_EMA_INDEX, 0.21); + trainer.isv_mapped.write_record(RL_AVG_WIN_USD_EMA_INDEX, 1.8); + trainer.isv_mapped.write_record(RL_AVG_LOSS_USD_EMA_INDEX, 1.0); + trainer.launch_rl_kelly_fraction_controller()?; + sync(&trainer); + + let f = trainer.read_isv_host(RL_KELLY_FRACTION_INDEX); + assert_eq!(f, 0.0, "Kelly at negative edge must clamp to 0; got {f}"); + eprintln!("G11 OK — Kelly = 0 at p=0.21 (analytic f_kelly < 0, clamped)"); + Ok(()) +} + +// G12 — Warmup gate: while cumulative_dones < min_trades_for_release, the +// kernel holds f at bootstrap=1.0 regardless of EMA inputs. This is the +// fix that resolved the fold-1 win-rate-collapse → negative-Kelly attractor +// (see [[pearl_position_sizing_missing_adaptation_layer]]). +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g12_kelly_held_at_bootstrap_during_warmup() -> Result<()> { + let Some((_dev, trainer)) = build_trainer() else { return Ok(()) }; + + // 500 trades observed, threshold 1000 → still warming up. + trainer.isv_mapped.write_record(RL_CUMULATIVE_DONES_INDEX, 500.0); + let min_trades = trainer.read_isv_host(RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX); + assert!((min_trades - KELLY_MIN_TRADES).abs() < 1e-3); + + // Even with negative-edge EMA inputs, Kelly must hold at 1.0. + trainer.isv_mapped.write_record(RL_WIN_RATE_EMA_INDEX, 0.10); + trainer.isv_mapped.write_record(RL_AVG_WIN_USD_EMA_INDEX, 1.0); + trainer.isv_mapped.write_record(RL_AVG_LOSS_USD_EMA_INDEX, 5.0); + trainer.launch_rl_kelly_fraction_controller()?; + sync(&trainer); + + let f = trainer.read_isv_host(RL_KELLY_FRACTION_INDEX); + assert_eq!( + f, 1.0, + "Kelly must hold at bootstrap=1.0 during warmup; got {f}" + ); + eprintln!("G12 OK — Kelly held at {f} during warmup (500 < {KELLY_MIN_TRADES} trades)"); + Ok(()) +} diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 26cf1be4f..ef771030e 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -43,7 +43,7 @@ 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_STEP_COUNTER_ISV_INDEX, RL_TRAIL_ADJUST_RATE_INDEX, + RL_PYRAMID_THRESHOLD_INDEX, RL_STEP_COUNTER_ISV_INDEX, RL_TRAIL_K_INIT_INDEX, RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX, }; use ml_alpha::trainer::integrated::{ @@ -404,7 +404,12 @@ fn unit_state_transitions() -> Result<()> { #[test] #[ignore = "requires CUDA (MlDevice::cuda(0))"] -fn trail_mutate_tighten_loosen_reciprocal() -> Result<()> { +fn trail_mutate_independent_tighten_loosen_factors() -> Result<()> { + // Adaptive risk management (spec 2026-05-30) migrated a7/a8 from the + // symmetric reciprocal `RL_TRAIL_ADJUST_RATE_INDEX` to independent + // ISV slots `RL_TRAIL_TIGHTEN_FACTOR_INDEX` (682) and + // `RL_TRAIL_LOOSEN_FACTOR_INDEX` (683). Each action multiplies the + // trail by its own factor; reciprocity is no longer enforced. let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; let stream = dev.cuda_stream()?.clone(); let b_size = 1; @@ -413,11 +418,14 @@ fn trail_mutate_tighten_loosen_reciprocal() -> Result<()> { // all-zeros; bootstrap defaults are written only during the first // `step_with_lobsim`, which the test doesn't invoke. Set all four // trail slots explicitly — in particular `RL_TRAIL_MAX_INDEX`, - // because the loosen branch does `fminf(trail_max, current/rate)` + // because the loosen branch does `fminf(trail_max, current * loosen)` // and a zero max would clamp output to 0. set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.001)?; set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 100.0)?; - set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?; + set_isv_slot(&mut trainer, &stream, + ml_alpha::rl::isv_slots::RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9)?; + set_isv_slot(&mut trainer, &stream, + ml_alpha::rl::isv_slots::RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1)?; let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; let initial_trail = 0.5_f32; @@ -428,21 +436,25 @@ fn trail_mutate_tighten_loosen_reciprocal() -> Result<()> { trainer.launch_rl_trail_mutate(&tighten_d, &unit_active_d, &mut trail_d, b_size)?; let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?; let after_tighten = trail[0]; + let expected_tighten = initial_trail * 0.9; assert!( - (after_tighten - initial_trail * 0.9).abs() < 1e-5, - "a7 should multiply trail by adjust_rate (0.9); expected {}, got {}", - initial_trail * 0.9, + (after_tighten - expected_tighten).abs() < 1e-5, + "a7 should multiply trail by tighten factor (0.9); expected {}, got {}", + expected_tighten, after_tighten ); - // ── a8 (loosen) — trail /= 0.9 ⇒ returns to within fp ε of initial + // ── a8 (loosen) — trail *= 1.1 let loosen_d = upload_i32(&stream, &[8])?; trainer.launch_rl_trail_mutate(&loosen_d, &unit_active_d, &mut trail_d, b_size)?; let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?; let after_loosen = trail[0]; + let expected_loosen = expected_tighten * 1.1; assert!( - (after_loosen - initial_trail).abs() < 1e-5, - "a7→a8 should restore trail (reciprocal); expected {initial_trail}, got {after_loosen}" + (after_loosen - expected_loosen).abs() < 1e-5, + "a8 should multiply trail by loosen factor (1.1); expected {}, got {}", + expected_loosen, + after_loosen ); // ── Inactive units MUST NOT mutate. trail[1..4] start at 0.0 and @@ -461,7 +473,7 @@ fn trail_mutate_tighten_loosen_reciprocal() -> Result<()> { ); eprintln!( - "PASS — trail mutate reciprocal: {initial_trail} → {after_tighten} → {after_loosen}" + "PASS — trail mutate independent factors: {initial_trail} → {after_tighten} → {after_loosen}" ); Ok(()) } @@ -1099,7 +1111,8 @@ fn trail_at_min_a7_stays_clamped() -> Result<()> { set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.1)?; set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 100.0)?; - set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?; + set_isv_slot(&mut trainer, &stream, + ml_alpha::rl::isv_slots::RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9)?; let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; let mut trail_d = upload_f32(&stream, &[0.1, 0.0, 0.0, 0.0])?; @@ -1127,7 +1140,8 @@ fn trail_at_max_a8_stays_clamped() -> Result<()> { set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.001)?; set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 1.0)?; - set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?; + set_isv_slot(&mut trainer, &stream, + ml_alpha::rl::isv_slots::RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1)?; let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; let mut trail_d = upload_f32(&stream, &[1.0, 0.0, 0.0, 0.0])?; @@ -1138,7 +1152,7 @@ fn trail_at_max_a8_stays_clamped() -> Result<()> { let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?; assert!( (trail[0] - 1.0).abs() < 1e-5, - "a8 at max (1.0/0.9=1.11) should clamp to TRAIL_MAX=1.0; got {}", + "a8 at max (1.0*1.1=1.10) should clamp to TRAIL_MAX=1.0; got {}", trail[0] ); diff --git a/docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md b/docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md new file mode 100644 index 000000000..5e24ac2f6 --- /dev/null +++ b/docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md @@ -0,0 +1,1185 @@ +# Adaptive Risk Management — Implementation Plan (Layered Stack) + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Close the fold-1 eval gap (-$616k OOS despite controllers fixed) by adding a 5-layer adaptive risk management stack: hard CMDP gates + IQN risk-averse action selection + inventory penalty + Kelly sizing + TrailTighten/Loosen wiring. The genuinely high-leverage change is **Layer 2 (IQN τ-quantile action selection)** — uses the existing IQN distribution to give risk-averse decisions for free. + +**Architecture:** Five orthogonal layers, each ISV-driven and adaptive. The agent's discrete action choice flows through: (Layer 2 selects which action via pessimistic Q) → (Layer 1 overrides if any safety constraint fires) → (Layer 4 scales lot count by observed edge) → (Layer 3 penalizes inventory drift in reward). All layers use the same Welford-variance infrastructure from b1ef6664a. + +**Tech Stack:** CUDA 12.4 sm_89/sm_86, cudarc 0.19.3, Rust 1.85+, single-stream graph capture, mapped-pinned ISV bus. Same as base codebase. + +**Spec:** `docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md` + +--- + +## File Structure + +### Created (8 new files) + +- `crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu` (~100 LOC) — Layer 1: session DD + cooldown + max-open + inventory gates +- `crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu` (~50 LOC) — Layer 2: adapts IQN action-selection quantile from drawdown +- `crates/ml-alpha/cuda/rl_inventory_beta_controller.cu` (~50 LOC) — Layer 3: adapts inventory penalty coefficient +- `crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu` (~80 LOC) — Layer 4: half-Kelly from observed win_rate × R-multiple +- `crates/ml-alpha/cuda/rl_win_rate_ema_update.cu` (~40 LOC) — feeds Kelly controller +- `crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu` (~50 LOC) — feeds Kelly controller (avg_win + avg_loss per-trade EMAs) +- `crates/ml-alpha/cuda/rl_inventory_variance_update.cu` (~40 LOC) — feeds Layer 3 β controller +- `crates/ml-alpha/tests/risk_stack_invariants.rs` (~12 GPU-oracle gates) + +### Modified (7 files) + +- `crates/ml-alpha/src/rl/isv_slots.rs` — +22 slots, `RL_SLOTS_END` 662 → 684 +- `crates/ml-alpha/build.rs` — register 7 new cubins +- `crates/ml-alpha/cuda/actions_to_market_targets.cu` — Layer 1 overrides + Layer 4 Kelly scaling + Layer D trail mutations +- `crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu` — Layer 3 inventory penalty term in reward shaping +- `crates/ml-alpha/cuda/rl_ensemble_action_value.cu` — Layer 2 quantile-τ argmax (replaces expected-Q argmax) +- `crates/ml-alpha/cuda/rl_fused_controllers.cu` — add new branches for τ, β, Kelly (matching b1ef6664a pattern) +- `crates/ml-alpha/src/trainer/integrated.rs` — load 7 new cubins, struct fields, per-step launches, bootstrap entries, session reset on fold boundary + +--- + +## Tasks + +### Task 1: Allocate 22 new ISV slots + +**Files:** +- Modify: `crates/ml-alpha/src/rl/isv_slots.rs` (after slot 661) + +- [ ] **Step 1: Add 22 new constants + bump `RL_SLOTS_END`** + +```rust +// ============================================================ +// Adaptive risk management — layered stack (spec 2026-05-30-adaptive-risk-management-design) +// +// Five orthogonal risk layers, each ISV-driven: +// Layer 1 — CMDP hard constraints (session DD, cooldown, max open, inventory) +// Layer 2 — IQN risk-averse action selection via adaptive τ +// Layer 3 — Inventory penalty (Avellaneda-Stoikov) in reward shaping +// Layer 4 — Kelly-fraction position sizing +// Layer D — TrailTighten/Loosen wiring (close dead-action gap) +// +// Per `feedback_adaptive_not_tuned`: every threshold below is signal-driven +// or warmup-gated, not a hardcoded constant. +// ============================================================ + +// ─── Layer 1: CMDP hard constraints (9 slots) ─── +pub const RL_SESSION_PNL_USD_INDEX: usize = 662; +pub const RL_SESSION_DD_LIMIT_USD_INDEX: usize = 663; +pub const RL_SESSION_DD_TRIGGERED_INDEX: usize = 664; +pub const RL_MAX_OPEN_UNITS_INDEX: usize = 665; +pub const RL_CONSEC_LOSS_LIMIT_INDEX: usize = 666; +pub const RL_CONSEC_LOSS_COUNT_INDEX: usize = 667; +pub const RL_COOLDOWN_REMAINING_STEPS_INDEX: usize = 668; +pub const RL_COOLDOWN_DURATION_INDEX: usize = 669; +pub const RL_NET_INVENTORY_LIMIT_USD_INDEX: usize = 670; + +// ─── Layer 2: IQN risk-averse τ (3 slots) ─── +pub const RL_IQN_ACTION_TAU_INDEX: usize = 671; +pub const RL_IQN_ACTION_TAU_MIN_INDEX: usize = 672; +pub const RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX: usize = 673; + +// ─── Layer 3: Inventory penalty (2 slots) ─── +pub const RL_INVENTORY_PENALTY_BETA_INDEX: usize = 674; +pub const RL_INVENTORY_VARIANCE_EMA_INDEX: usize = 675; + +// ─── Layer 4: Kelly-fraction sizing (6 slots) ─── +pub const RL_KELLY_FRACTION_INDEX: usize = 676; +pub const RL_WIN_RATE_EMA_INDEX: usize = 677; +pub const RL_AVG_WIN_USD_EMA_INDEX: usize = 678; +pub const RL_AVG_LOSS_USD_EMA_INDEX: usize = 679; +pub const RL_KELLY_SAFETY_FRAC_INDEX: usize = 680; +pub const RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX: usize = 681; + +// ─── Layer D: TrailTighten/Loosen wiring (2 slots) ─── +pub const RL_TRAIL_TIGHTEN_FACTOR_INDEX: usize = 682; +pub const RL_TRAIL_LOOSEN_FACTOR_INDEX: usize = 683; + +/// Last RL-allocated slot index (exclusive). Pre-risk-stack: 662. +/// Post-risk-stack: 684. Total new: 22 slots, +88 bytes mapped-pinned. +pub const RL_SLOTS_END: usize = 684; +``` + +- [ ] **Step 2: Verify cargo check passes + no duplicate slot numbers** + +```bash +cd /home/jgrusewski/Work/foxhunt +SQLX_OFFLINE=true cargo check -p ml-alpha 2>&1 | tail -3 +grep -oE "= [0-9]+;" crates/ml-alpha/src/rl/isv_slots.rs | sort -n | uniq -c | sort -rn | head -3 +``` + +Expected: `Finished`, all counts = 1. + +--- + +### Task 2: New EMA producer kernels for risk signals + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_win_rate_ema_update.cu` +- Create: `crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu` +- Create: `crates/ml-alpha/cuda/rl_inventory_variance_update.cu` +- Modify: `crates/ml-alpha/build.rs` + +- [ ] **Step 1: Write `rl_win_rate_ema_update.cu`** + +```c +// rl_win_rate_ema_update.cu — track observed win_rate EMA from trade outcomes. +// Per pearl_first_observation_bootstrap: sentinel = 0.5 (neutral) until first +// closed trade. Per feedback_no_atomicadd: single-thread block, no atomics. + +#define RL_WIN_RATE_EMA_INDEX 677 +#define EMA_ALPHA 0.05f // ~20-trade effective window + +extern "C" __global__ void rl_win_rate_ema_update( + float* __restrict__ isv, + const int* __restrict__ trade_outcome_per_batch, // [b_size] −1/0/+1 + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Count this step's closed trades (outcome != 0). + int closed = 0, wins = 0; + for (int b = 0; b < b_size; b++) { + const int o = trade_outcome_per_batch[b]; + if (o != 0) { closed++; if (o == 1) wins++; } + } + if (closed == 0) return; // no observation this step + + const float step_wr = (float)wins / (float)closed; + const float prev = isv[RL_WIN_RATE_EMA_INDEX]; + // First-observation bootstrap on sentinel 0.5 (neutral default). + if (prev == 0.5f) { + isv[RL_WIN_RATE_EMA_INDEX] = step_wr; + } else { + isv[RL_WIN_RATE_EMA_INDEX] = (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_wr; + } +} +``` + +- [ ] **Step 2: Write `rl_avg_win_loss_ema_update.cu`** + +```c +// rl_avg_win_loss_ema_update.cu — track avg_win + avg_loss per-trade EMAs in USD. +// Per pearl_first_observation_bootstrap: sentinel = 0.0 (no observation). + +#define RL_AVG_WIN_USD_EMA_INDEX 678 +#define RL_AVG_LOSS_USD_EMA_INDEX 679 +#define EMA_ALPHA 0.05f + +extern "C" __global__ void rl_avg_win_loss_ema_update( + float* __restrict__ isv, + const float* __restrict__ realized_pnl_per_batch, // [b_size] USD, signed + const int* __restrict__ trade_outcome_per_batch, // [b_size] −1/0/+1 + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Sum per-trade |pnl| for wins / losses this step. + float sum_win = 0.0f, sum_loss = 0.0f; + int n_win = 0, n_loss = 0; + for (int b = 0; b < b_size; b++) { + const int o = trade_outcome_per_batch[b]; + const float p = realized_pnl_per_batch[b]; + if (o == 1) { sum_win += p; n_win++; } + if (o == -1) { sum_loss += fabsf(p); n_loss++; } + } + + if (n_win > 0) { + const float step_avg = sum_win / (float)n_win; + const float prev = isv[RL_AVG_WIN_USD_EMA_INDEX]; + isv[RL_AVG_WIN_USD_EMA_INDEX] = (prev == 0.0f) ? step_avg + : (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_avg; + } + if (n_loss > 0) { + const float step_avg = sum_loss / (float)n_loss; + const float prev = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + isv[RL_AVG_LOSS_USD_EMA_INDEX] = (prev == 0.0f) ? step_avg + : (1.0f - EMA_ALPHA) * prev + EMA_ALPHA * step_avg; + } +} +``` + +- [ ] **Step 3: Write `rl_inventory_variance_update.cu`** + +```c +// rl_inventory_variance_update.cu — track running variance of |net_position| +// across batches. Used by Layer 3 β controller to scale penalty. +// Welford online variance, single-thread single-block. + +#define RL_INVENTORY_VARIANCE_EMA_INDEX 675 +#define WELFORD_ALPHA 0.01f + +extern "C" __global__ void rl_inventory_variance_update( + float* __restrict__ isv, + const float* __restrict__ net_position_per_batch, // [b_size] signed + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Compute batch mean + variance of |net_position|. + float sum = 0.0f, sum_sq = 0.0f; + for (int b = 0; b < b_size; b++) { + const float v = fabsf(net_position_per_batch[b]); + sum += v; + sum_sq += v * v; + } + const float n = (float)b_size; + const float mean = sum / n; + const float var = (sum_sq / n) - (mean * mean); + if (var < 0.0f) return; // numerical safety + + const float prev = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; + isv[RL_INVENTORY_VARIANCE_EMA_INDEX] = (prev == 0.0f) ? var + : (1.0f - WELFORD_ALPHA) * prev + WELFORD_ALPHA * var; +} +``` + +- [ ] **Step 4: Register all 3 cubins in `build.rs`** + +In the cubin list (next to other RL kernels): +```rust +"rl_win_rate_ema_update", +"rl_avg_win_loss_ema_update", +"rl_inventory_variance_update", +``` + +- [ ] **Step 5: Verify build** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +ls target/release/build/ml-alpha-*/out/rl_win_rate_ema_update.cubin +ls target/release/build/ml-alpha-*/out/rl_avg_win_loss_ema_update.cubin +ls target/release/build/ml-alpha-*/out/rl_inventory_variance_update.cubin +``` + +Expected: `Finished`, all 3 cubins present. + +--- + +### Task 3: Layer 1 — CMDP constraints kernel + actions_to_market_targets overrides + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu` +- Modify: `crates/ml-alpha/build.rs` +- Modify: `crates/ml-alpha/cuda/actions_to_market_targets.cu` + +- [ ] **Step 1: Write `rl_cmdp_constraints_check.cu`** + +```c +// rl_cmdp_constraints_check.cu — Layer 1 hard CMDP gates (spec 2026-05-30-adaptive-risk-management). +// +// Five independent risk constraints. ALL must be satisfied for the agent's +// action to execute; ANY violation forces Hold. +// +// 1. Session pnl accumulation + DD cap check (writes triggered flag) +// 2. Cooldown counter decrement (after consecutive losses limit) +// 3. Consecutive-loss tracking (sets cooldown when limit hit) +// 4. Max-open and inventory checks read directly from per-batch state +// in actions_to_market_targets — this kernel only manages the +// session-level + cooldown state. +// +// Per feedback_no_atomicadd: single-thread block, sums across b_size manually. + +#define RL_SESSION_PNL_USD_INDEX 662 +#define RL_SESSION_DD_LIMIT_USD_INDEX 663 +#define RL_SESSION_DD_TRIGGERED_INDEX 664 +#define RL_CONSEC_LOSS_LIMIT_INDEX 666 +#define RL_CONSEC_LOSS_COUNT_INDEX 667 +#define RL_COOLDOWN_REMAINING_STEPS_INDEX 668 +#define RL_COOLDOWN_DURATION_INDEX 669 + +extern "C" __global__ void rl_cmdp_constraints_check( + float* __restrict__ isv, + const float* __restrict__ realized_pnl_per_batch, // [b_size] USD + const int* __restrict__ trade_outcome_per_batch, // [b_size] −1/0/+1 + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // 1. Session pnl + DD check. + float step_pnl = 0.0f; + for (int b = 0; b < b_size; b++) step_pnl += realized_pnl_per_batch[b]; + const float session_pnl_new = isv[RL_SESSION_PNL_USD_INDEX] + step_pnl; + isv[RL_SESSION_PNL_USD_INDEX] = session_pnl_new; + if (session_pnl_new < isv[RL_SESSION_DD_LIMIT_USD_INDEX]) { + isv[RL_SESSION_DD_TRIGGERED_INDEX] = 1.0f; + } + + // 2. Cooldown decrement. + const float cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX]; + if (cooldown > 0.0f) { + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = cooldown - 1.0f; + } + + // 3. Consecutive loss tracking. + float consec = isv[RL_CONSEC_LOSS_COUNT_INDEX]; + for (int b = 0; b < b_size; b++) { + const int o = trade_outcome_per_batch[b]; + if (o == -1) consec += 1.0f; + else if (o == 1) consec = 0.0f; + } + if (consec >= isv[RL_CONSEC_LOSS_LIMIT_INDEX]) { + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = isv[RL_COOLDOWN_DURATION_INDEX]; + isv[RL_CONSEC_LOSS_COUNT_INDEX] = 0.0f; + } else { + isv[RL_CONSEC_LOSS_COUNT_INDEX] = consec; + } +} +``` + +- [ ] **Step 2: Register cubin in `build.rs`** + +Add: `"rl_cmdp_constraints_check",` + +- [ ] **Step 3: Modify `actions_to_market_targets.cu` — add Layer 1 override block at top of per-batch loop** + +Find the per-batch action translation loop. Before the action-specific switch statement, insert: + +```c +// ── Layer 1 (CMDP) override checks ── +// Per spec 2026-05-30-adaptive-risk-management-design. Any violation +// forces Hold. Order: most catastrophic first. + +const bool dd_triggered = isv[RL_SESSION_DD_TRIGGERED_INDEX] >= 0.5f; +const bool in_cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] > 0.0f; +if (dd_triggered || in_cooldown) { + actions[b] = ACTION_HOLD; + target_lots[b] = 0; + continue; // skip rest of action translation for this batch +} + +// Max-open units check (per-batch dynamic count). +int active_units = 0; +for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active_d[b * MAX_UNITS + u]) active_units++; +} +const int max_open = (int)isv[RL_MAX_OPEN_UNITS_INDEX]; +if (active_units >= max_open && is_opening_action(actions[b])) { + actions[b] = ACTION_HOLD; + // continue translating since Hold has no further side-effects +} + +// Inventory cap check. +const float net_pos_usd = compute_net_position_usd(b, unit_active_d, unit_lots_d, current_mid); +const float inv_limit = isv[RL_NET_INVENTORY_LIMIT_USD_INDEX]; +if (fabsf(net_pos_usd) > inv_limit && would_increase_exposure(actions[b], net_pos_usd)) { + actions[b] = ACTION_HOLD; +} +``` + +Where `is_opening_action(a)` returns true for {Long, Short, ...} and false for {Hold, Flat, HalfFlat, Trail*}. Define as a `__device__` helper near the top of the file. + +- [ ] **Step 4: Verify build + integrated_trainer_smoke still passes** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +``` + +Expected: `Finished` + smoke passes (controllers' new ISV slots are sentinel 0 so overrides don't fire). + +--- + +### Task 4: Layer 2 — IQN risk-averse action selection + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu` +- Modify: `crates/ml-alpha/build.rs` +- Modify: `crates/ml-alpha/cuda/rl_ensemble_action_value.cu` (or wherever IQN action selection happens — verify location during Step 0) + +- [ ] **Step 0: Locate the IQN action selection site** + +```bash +cd /home/jgrusewski/Work/foxhunt +grep -rn "expected_q\|argmax.*Q\|compute_q\|RL_IQN_ENSEMBLE_ALPHA" crates/ml-alpha/cuda/rl_ensemble_action_value.cu | head -10 +grep -rn "iqn.*action\|iqn_head.*forward\|sample_tau" crates/ml-alpha/src/trainer/integrated.rs | head -10 +``` + +Identify the kernel that computes per-action expected Q from IQN quantiles. Most likely `rl_ensemble_action_value.cu`. + +- [ ] **Step 1: Write `rl_iqn_action_tau_controller.cu`** + +```c +// rl_iqn_action_tau_controller.cu — Layer 2: adapt IQN action-selection quantile τ +// from observed session drawdown. Normal trading: τ=0.5 (median ≈ expected-Q-like). +// In drawdown: τ → MIN (pessimistic, downside-aware). + +#define RL_IQN_ACTION_TAU_INDEX 671 +#define RL_IQN_ACTION_TAU_MIN_INDEX 672 +#define RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX 673 +#define RL_SESSION_PNL_USD_INDEX 662 +#define RL_SESSION_MAX_PNL_USD_INDEX 684 // see Step 2 (new slot if not yet added) + +// NOTE: starting capital reference for normalizing drawdown. If not in ISV, +// use a default $35k constant — but per `feedback_adaptive_not_tuned` add to ISV. +#define DEFAULT_STARTING_CAPITAL_USD 35000.0f + +extern "C" __global__ void rl_iqn_action_tau_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float session_pnl = isv[RL_SESSION_PNL_USD_INDEX]; + float max_pnl = isv[RL_SESSION_MAX_PNL_USD_INDEX]; + if (session_pnl > max_pnl) { + max_pnl = session_pnl; + isv[RL_SESSION_MAX_PNL_USD_INDEX] = max_pnl; + } + + const float drawdown_usd = fmaxf(0.0f, max_pnl - session_pnl); + const float drawdown_frac = drawdown_usd / DEFAULT_STARTING_CAPITAL_USD; + + const float tau_min = isv[RL_IQN_ACTION_TAU_MIN_INDEX]; + const float sensitivity = isv[RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX]; + const float tau_target = fmaxf(tau_min, 0.5f - sensitivity * drawdown_frac); + isv[RL_IQN_ACTION_TAU_INDEX] = tau_target; +} +``` + +NOTE: If `RL_SESSION_MAX_PNL_USD_INDEX` isn't yet allocated, add it to Task 1 (becomes slot 684, `RL_SLOTS_END` 684 → 685). Same for `RL_STARTING_CAPITAL_USD_INDEX` if we want it ISV-driven (recommended; bump to 686). Update Task 1 to include both. + +- [ ] **Step 2: Register cubin in `build.rs`** + +Add: `"rl_iqn_action_tau_controller",` + +- [ ] **Step 3: Modify the IQN ensemble action-value kernel to use adaptive τ** + +In `rl_ensemble_action_value.cu`, find where Q values per action are computed. The IQN side currently uses `expected_Q = E_τ[Q(s, a, τ)]`. Replace with quantile-specific Q: + +```c +// BEFORE (illustrative — actual code structure may differ): +float exp_q_iqn[N_ACTIONS]; +for (int a = 0; a < N_ACTIONS; a++) { + float sum = 0.0f; + for (int t = 0; t < N_TAU; t++) sum += iqn_q[a * N_TAU + t]; + exp_q_iqn[a] = sum / (float)N_TAU; +} + +// AFTER: +const float tau_a = isv[RL_IQN_ACTION_TAU_INDEX]; +float pess_q_iqn[N_ACTIONS]; +for (int a = 0; a < N_ACTIONS; a++) { + // Find the τ-th quantile (linear interpolation across sampled τ values). + // tau_samples[t] is monotonic in t (sorted), and tau_a ∈ [0, 1]. + const float target_idx = tau_a * (float)(N_TAU - 1); + const int t_lo = (int)floorf(target_idx); + const int t_hi = min(t_lo + 1, N_TAU - 1); + const float frac = target_idx - (float)t_lo; + pess_q_iqn[a] = (1.0f - frac) * iqn_q[a * N_TAU + t_lo] + frac * iqn_q[a * N_TAU + t_hi]; +} + +// Existing ensemble blend with C51 side stays the same: +const float alpha = isv[RL_IQN_ENSEMBLE_ALPHA_INDEX]; +float ensemble_q[N_ACTIONS]; +for (int a = 0; a < N_ACTIONS; a++) { + ensemble_q[a] = alpha * c51_expected[a] + (1.0f - alpha) * pess_q_iqn[a]; +} +// argmax(ensemble_q) → selected_action +``` + +NOTE: This assumes `iqn_q` is sorted ascending in τ. If not, must sort first OR sample at the specific τ via the IQN forward kernel. Check kernel architecture during implementation. + +- [ ] **Step 4: Verify build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +``` + +Expected: clean + smoke passes (τ=0.5 default = same as expected-Q behavior approximately). + +--- + +### Task 5: Layer 3 — Inventory penalty controller + reward shaping + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_inventory_beta_controller.cu` +- Modify: `crates/ml-alpha/build.rs` +- Modify: `crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu` + +- [ ] **Step 1: Write `rl_inventory_beta_controller.cu`** + +```c +// rl_inventory_beta_controller.cu — Layer 3: adapt inventory penalty coefficient β. +// +// Target: penalty ≈ 1% of typical absolute reward when inventory is at ~2σ. +// β = 0.01 × E[|reward|] / (2 × σ_inventory). +// Wiener blend with floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary. +// Bootstrap on sentinel 0.0 (β = 0 until variance EMA accumulates real signal). + +#define RL_INVENTORY_PENALTY_BETA_INDEX 674 +#define RL_INVENTORY_VARIANCE_EMA_INDEX 675 +#define RL_REWARD_MAGNITUDE_EMA_INDEX 614 // from b1ef6664a +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 // from b1ef6664a + +extern "C" __global__ void rl_inventory_beta_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float reward_mag = isv[RL_REWARD_MAGNITUDE_EMA_INDEX]; + const float inv_var = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; + if (inv_var <= 0.0f || reward_mag <= 0.0f) return; // dead-signal hold + + const float inv_std = sqrtf(inv_var); + const float beta_target = 0.01f * reward_mag / (2.0f * inv_std); + + const float prev = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; + const float a = isv[RL_WIENER_ALPHA_FLOOR_INDEX]; // typically 0.4 + if (prev == 0.0f) { + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = beta_target; + } else { + isv[RL_INVENTORY_PENALTY_BETA_INDEX] = (1.0f - a) * prev + a * beta_target; + } +} +``` + +- [ ] **Step 2: Register cubin** + +Add to `build.rs`: `"rl_inventory_beta_controller",` + +- [ ] **Step 3: Modify `rl_fused_reward_pipeline.cu` — apply inventory penalty** + +Find the reward computation per batch slot. After raw reward, before clamps: + +```c +// Layer 3 (spec 2026-05-30): inventory penalty. +// reward_shaped = reward_raw − β × |net_position| +const float net_pos = net_position_per_batch[b]; // signed contracts +const float inv_beta = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; +const float inv_penalty = inv_beta * fabsf(net_pos); +shaped_reward = shaped_reward - inv_penalty; +``` + +- [ ] **Step 4: Build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +``` + +--- + +### Task 6: Layer 4 — Kelly-fraction sizing + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu` +- Modify: `crates/ml-alpha/build.rs` +- Modify: `crates/ml-alpha/cuda/actions_to_market_targets.cu` + +- [ ] **Step 1: Write `rl_kelly_fraction_controller.cu`** + +```c +// rl_kelly_fraction_controller.cu — Layer 4: half-Kelly position size from observed edge. +// +// f_kelly = (p × b − q) / b where p = win_rate, q = 1−p, b = avg_win/avg_loss +// f = clamp(safety × f_kelly, 0, 1) where safety = 0.5 (Thorp half-Kelly) +// Warmup: until cumulative_dones ≥ MIN_TRADES_FOR_RELEASE, hold at bootstrap=1.0. + +#define RL_KELLY_FRACTION_INDEX 676 +#define RL_WIN_RATE_EMA_INDEX 677 +#define RL_AVG_WIN_USD_EMA_INDEX 678 +#define RL_AVG_LOSS_USD_EMA_INDEX 679 +#define RL_KELLY_SAFETY_FRAC_INDEX 680 +#define RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX 681 +#define RL_CUMULATIVE_DONES_INDEX 660 // from b1ef6664a + +extern "C" __global__ void rl_kelly_fraction_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Warmup gate: hold at bootstrap=1.0 until enough closed trades. + if (isv[RL_CUMULATIVE_DONES_INDEX] < isv[RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX]) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + return; + } + + const float p = isv[RL_WIN_RATE_EMA_INDEX]; + const float avg_win = isv[RL_AVG_WIN_USD_EMA_INDEX]; + const float avg_loss = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + + // Dead-signal guard: need positive loss magnitude for R-multiple. + if (avg_loss <= 0.0f) { isv[RL_KELLY_FRACTION_INDEX] = 1.0f; return; } + + const float b = avg_win / avg_loss; + const float q = 1.0f - p; + const float f_kelly = (p * b - q) / b; + const float safety = isv[RL_KELLY_SAFETY_FRAC_INDEX]; // typically 0.5 + float f = f_kelly * safety; + f = fmaxf(0.0f, fminf(f, 1.0f)); + isv[RL_KELLY_FRACTION_INDEX] = f; +} +``` + +- [ ] **Step 2: Register cubin** + +Add to `build.rs`: `"rl_kelly_fraction_controller",` + +- [ ] **Step 3: Apply Kelly scaling in `actions_to_market_targets.cu`** + +After Layer 1 overrides (Task 3) have been applied and before returning per-batch results: + +```c +// Layer 4 (spec 2026-05-30): Kelly-fraction sizing. +// Apply AFTER Layer 1 overrides have set target_lots based on the action. +const float kelly = isv[RL_KELLY_FRACTION_INDEX]; +target_lots[b] = (int)ceilf((float)target_lots[b] * kelly); +// If kelly = 0 (Kelly negative or warmup), target_lots = 0 → no commitment. +``` + +- [ ] **Step 4: Build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +``` + +--- + +### Task 7: Layer D — TrailTighten/Loosen wiring + +**Files:** +- Modify: `crates/ml-alpha/cuda/actions_to_market_targets.cu` + +- [ ] **Step 1: Wire a7 (TrailTighten) and a8 (TrailLoosen) into the action switch** + +Find the existing switch statement (likely already has `case ACTION_TRAIL_TIGHTEN:` as a stub). Replace stub with: + +```c +case ACTION_TRAIL_TIGHTEN: { + // Spec 2026-05-30 Layer D — close pearl_dead_trail_stop_actions_a7_a8 gap. + const float factor = isv[RL_TRAIL_TIGHTEN_FACTOR_INDEX]; // bootstrap 0.9 + for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active_d[b * MAX_UNITS + u]) { + unit_trail_distance_d[b * MAX_UNITS + u] *= factor; + } + } + // No new market order — pure ISV state mutation. + target_lots[b] = 0; + break; +} +case ACTION_TRAIL_LOOSEN: { + const float factor = isv[RL_TRAIL_LOOSEN_FACTOR_INDEX]; // bootstrap 1.1 + for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active_d[b * MAX_UNITS + u]) { + unit_trail_distance_d[b * MAX_UNITS + u] *= factor; + } + } + target_lots[b] = 0; + break; +} +``` + +- [ ] **Step 2: Build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +``` + +--- + +### Task 8: Fused kernel integration + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_fused_controllers.cu` + +- [ ] **Step 1: Add Layer 2 + Layer 3 + Layer 4 branches** + +The fused kernel already has 10 controller branches. Append 3 more (matching the per-controller standalone logic from Tasks 4, 5, 6) at the end. Layer 1 stays a separate kernel because it needs per-step `dones` + `realized_pnl` arrays. + +```c +// ═══════════════════════════════════════════════════════════════════ +// 11. IQN action τ controller → ISV[671] (Layer 2) +// ═══════════════════════════════════════════════════════════════════ +{ + // [paste body of rl_iqn_action_tau_controller from Task 4 Step 1] +} + +// ═══════════════════════════════════════════════════════════════════ +// 12. Inventory β controller → ISV[674] (Layer 3) +// ═══════════════════════════════════════════════════════════════════ +{ + // [paste body of rl_inventory_beta_controller from Task 5 Step 1] +} + +// ═══════════════════════════════════════════════════════════════════ +// 13. Kelly fraction controller → ISV[676] (Layer 4) +// ═══════════════════════════════════════════════════════════════════ +{ + // [paste body of rl_kelly_fraction_controller from Task 6 Step 1] +} +``` + +- [ ] **Step 2: Build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test signal_variance_kernel --release -- --ignored 2>&1 | tail -3 +``` + +--- + +### Task 9: Trainer integration + +**Files:** +- Modify: `crates/ml-alpha/src/trainer/integrated.rs` + +- [ ] **Step 1: Load 7 new cubins, add struct fields, struct init** + +Follow the exact pattern from b1ef6664a's `rl_signal_variance_update` integration. For each of: +- `rl_cmdp_constraints_check` +- `rl_iqn_action_tau_controller` +- `rl_inventory_beta_controller` +- `rl_kelly_fraction_controller` +- `rl_win_rate_ema_update` +- `rl_avg_win_loss_ema_update` +- `rl_inventory_variance_update` + +Add: +```rust +const RL__CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_.cubin")); +// + struct field _module + fn +// + load in new() +// + init in Ok(Self {...}) +``` + +- [ ] **Step 2: Add per-step launches in the per-step pipeline** + +Find the per-step block (search for `launch_rl_fused_controllers`). Insert NEW launches **between** the EMA producers and the controllers: + +```rust +// After existing EMA producer launches: + +// Layer 1 — CMDP state update (per-step, reads pnl + outcomes) +self.launch_rl_cmdp_constraints_check( + &realized_pnl_per_batch_d, + &trade_outcome_per_batch_d, + b_size, +)?; + +// New EMA producers (feed Kelly + inventory β controllers) +self.launch_rl_win_rate_ema_update(&trade_outcome_per_batch_d, b_size)?; +self.launch_rl_avg_win_loss_ema_update( + &realized_pnl_per_batch_d, + &trade_outcome_per_batch_d, + b_size, +)?; +self.launch_rl_inventory_variance_update(&net_position_per_batch_d, b_size)?; + +// Layer 2/3/4 controllers — already integrated into fused kernel (Task 8), +// so no separate launches needed. The fused kernel handles them. + +// Existing fused controllers launch continues as-is. +self.launch_rl_fused_controllers(b_size)?; +``` + +Each `launch_*` helper follows the same pattern as `launch_rl_signal_variance_update` from b1ef6664a. + +- [ ] **Step 3: Bootstrap all 22 new ISV slots in `with_controllers_bootstrapped`** + +Add entries to the `isv_constants` array (currently 147 entries from b1ef6664a). New entries: + +```rust +// Layer 1 +(RL_SESSION_PNL_USD_INDEX, 0.0), // session boundary reset +(RL_SESSION_DD_LIMIT_USD_INDEX, -3500.0), // 10% of $35k starting capital +(RL_SESSION_DD_TRIGGERED_INDEX, 0.0), +(RL_MAX_OPEN_UNITS_INDEX, 4.0), // = current MAX_UNITS +(RL_CONSEC_LOSS_LIMIT_INDEX, 10.0), +(RL_CONSEC_LOSS_COUNT_INDEX, 0.0), +(RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0), +(RL_COOLDOWN_DURATION_INDEX, 500.0), +(RL_NET_INVENTORY_LIMIT_USD_INDEX, 10000.0), +// Layer 2 +(RL_IQN_ACTION_TAU_INDEX, 0.5), // median = ~current behavior +(RL_IQN_ACTION_TAU_MIN_INDEX, 0.1), +(RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX, 5.0), +// Layer 3 +(RL_INVENTORY_PENALTY_BETA_INDEX, 0.0), // sentinel +(RL_INVENTORY_VARIANCE_EMA_INDEX, 0.0), // sentinel +// Layer 4 +(RL_KELLY_FRACTION_INDEX, 1.0), // full size pre-warmup +(RL_WIN_RATE_EMA_INDEX, 0.5), // neutral sentinel +(RL_AVG_WIN_USD_EMA_INDEX, 0.0), // sentinel +(RL_AVG_LOSS_USD_EMA_INDEX, 0.0), // sentinel +(RL_KELLY_SAFETY_FRAC_INDEX, 0.5), // half-Kelly Thorp default +(RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX, 1000.0), +// Layer D +(RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9), +(RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1), +``` + +Bump array size to 169. + +- [ ] **Step 4: Add session-boundary reset on train/eval phase transition** + +Find the trainer's fold/eval transition (look for `n_eval_steps` or `eval phase` markers in `integrated.rs`). At the transition, reset Layer 1 session-level state: + +```rust +// Session boundary — reset Layer 1 session state for fresh eval phase. +self.isv_mapped.write_record(RL_SESSION_PNL_USD_INDEX, 0.0); +self.isv_mapped.write_record(RL_SESSION_DD_TRIGGERED_INDEX, 0.0); +self.isv_mapped.write_record(RL_CONSEC_LOSS_COUNT_INDEX, 0.0); +self.isv_mapped.write_record(RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0); +// Also reset Kelly + inventory EMAs if we want eval to start clean (debatable) +// — leave them for now to test "policy learned during train, applied in eval" +``` + +- [ ] **Step 5: Build + smoke** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test controller_adaptive_floors --release -- --ignored 2>&1 | tail -3 +SQLX_OFFLINE=true cargo test -p ml-alpha --test signal_variance_kernel --release -- --ignored 2>&1 | tail -3 +``` + +Expected: all pass. + +--- + +### Task 10: GPU-oracle invariant tests + +**Files:** +- Create: `crates/ml-alpha/tests/risk_stack_invariants.rs` + +- [ ] **Step 1: Write 12 GPU-oracle tests** + +Follow the pattern from `controller_adaptive_floors.rs`. For brevity, the full test bodies are templated below — instantiate each gate using `IntegratedTrainer::new` + ISV writes + `launch_rl_fused_controllers` + ISV reads. + +```rust +//! Risk-stack invariants (spec 2026-05-30-adaptive-risk-management-design). +//! +//! 12 GPU-oracle gates — one per layer behavior + integration tests. +//! +//! Run with: +//! cargo test -p ml-alpha --test risk_stack_invariants --release -- --ignored --nocapture + +use ml_alpha::rl::isv_slots::*; +use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; +use ml_alpha::trainer::perception::PerceptionTrainerConfig; +use ml_core::device::MlDevice; + +// [build_trainer helper — copy from controller_adaptive_floors.rs] + +#[test] #[ignore] +fn g1_cmdp_dd_breaker_triggers_at_limit() { /* ... */ } + +#[test] #[ignore] +fn g2_cmdp_cooldown_after_consec_losses() { /* ... */ } + +#[test] #[ignore] +fn g3_cmdp_max_open_forces_hold_on_opening_action() { /* ... */ } + +#[test] #[ignore] +fn g4_cmdp_inventory_cap_blocks_one_sided_expansion() { /* ... */ } + +#[test] #[ignore] +fn g5_iqn_tau_adapts_to_drawdown() { /* ... */ } + +#[test] #[ignore] +fn g6_iqn_tau_at_peak_returns_0_5() { /* ... */ } + +#[test] #[ignore] +fn g7_inventory_beta_adapts_from_variance() { /* ... */ } + +#[test] #[ignore] +fn g8_kelly_at_known_edge_returns_0_15() { + // win_rate=0.55, avg_win=1.8, avg_loss=1.0, trades=2000 → f = 0.5 × 0.30 = 0.15 +} + +#[test] #[ignore] +fn g9_kelly_at_negative_edge_returns_0() { + // win_rate=0.21, avg_win=1.8, avg_loss=1.0, trades=2000 → f = 0 +} + +#[test] #[ignore] +fn g10_kelly_held_at_bootstrap_during_warmup() { + // cumulative_dones=500 (< 1000) → f = 1.0 regardless of input +} + +#[test] #[ignore] +fn g11_trail_tighten_reduces_distance() { + // Pre-set unit_trail_distance=100, agent picks a7 → distance becomes 90 +} + +#[test] #[ignore] +fn g12_dd_breaker_overrides_kelly() { + // Set kelly=1.0 AND dd_triggered=1 → target_lots=0 +} +``` + +- [ ] **Step 2: Run all gates** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --test risk_stack_invariants --release -- --ignored --nocapture 2>&1 | tail -20 +``` + +Expected: 12 passed. + +--- + +### Task 11: Local 1k smoke + compute-sanitizer + +- [ ] **Step 1: 1k-step smoke at b=128** + +```bash +cd /home/jgrusewski/Work/foxhunt +SMOKE_OUT=/tmp/risk-stack-smoke +TEST_DATA=/home/jgrusewski/Work/foxhunt/test_data/futures-baseline +rm -rf "$SMOKE_OUT" && mkdir -p "$SMOKE_OUT/predecoded" +SQLX_OFFLINE=true cargo build -p ml-alpha --release --example alpha_rl_train 2>&1 | tail -2 +./target/release/examples/alpha_rl_train \ + --mbp10-data-dir "$TEST_DATA/ES.FUT" \ + --predecoded-dir "$SMOKE_OUT/predecoded" \ + --out "$SMOKE_OUT" \ + --n-steps 1000 --n-backtests 128 --seq-len 32 \ + --per-capacity 1024 --seed 16962 --instrument-mode all \ + --n-folds 1 --fold-idx 0 --n-eval-steps 0 \ + --diag-jsonl "$SMOKE_OUT/diag.jsonl" +``` + +Expected: completes 1000/1000 steps clean, no NaN. + +- [ ] **Step 2: Inspect risk-stack ISVs in diag** + +```bash +tail -1 $SMOKE_OUT/diag.jsonl | tr "," "\n" | \ + grep -E "kelly_fraction|iqn_action_tau|inventory_penalty|session_pnl|session_dd_triggered" +``` + +Expected: kelly_fraction ∈ [0, 1], iqn_action_tau ∈ [0.1, 0.5], inventory_penalty_beta > 0, session_pnl tracked, session_dd_triggered = 0 (unlikely to hit at b=128 1k). + +- [ ] **Step 3: compute-sanitizer memcheck** + +```bash +SANITIZER_OUT=/tmp/risk-stack-cs +rm -rf "$SANITIZER_OUT" && mkdir -p "$SANITIZER_OUT/predecoded" +compute-sanitizer --tool memcheck --launch-timeout 600 --error-exitcode 1 \ + ./target/release/examples/alpha_rl_train \ + --mbp10-data-dir "$TEST_DATA/ES.FUT" \ + --predecoded-dir "$SANITIZER_OUT/predecoded" \ + --out "$SANITIZER_OUT" \ + --n-steps 5 --n-backtests 128 --seq-len 32 \ + --per-capacity 1024 --seed 16962 --instrument-mode all \ + --n-folds 1 --fold-idx 0 --n-eval-steps 0 \ + 2>&1 | grep "ERROR SUMMARY" +``` + +Expected: `ERROR SUMMARY: 0 errors`. + +--- + +### Task 12: Atomic commit + push + +- [ ] **Step 1: Stage all changes** + +```bash +cd /home/jgrusewski/Work/foxhunt +git add crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu \ + crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu \ + crates/ml-alpha/cuda/rl_inventory_beta_controller.cu \ + crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu \ + crates/ml-alpha/cuda/rl_win_rate_ema_update.cu \ + crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu \ + crates/ml-alpha/cuda/rl_inventory_variance_update.cu \ + crates/ml-alpha/cuda/actions_to_market_targets.cu \ + crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu \ + crates/ml-alpha/cuda/rl_ensemble_action_value.cu \ + crates/ml-alpha/cuda/rl_fused_controllers.cu \ + crates/ml-alpha/src/rl/isv_slots.rs \ + crates/ml-alpha/src/trainer/integrated.rs \ + crates/ml-alpha/tests/risk_stack_invariants.rs \ + crates/ml-alpha/build.rs \ + docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md \ + docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md +``` + +- [ ] **Step 2: Commit with comprehensive message** + +```bash +git commit -m "$(cat <<'EOF' +feat(rl): adaptive risk management — 5-layer stack + +Closes the train-eval gap surfaced by the b1ef6664a adaptive-controller- +floor refactor (fold 1 eval -$616k despite train -$1.05M = 68% better +than OLD). Diagnosis from fold 1 eval JSONL: agent has good per-trade +R-multiple (1.8) but win_rate collapses 0.55 (train) → 0.21 (eval) → +Kelly fraction NEGATIVE. The agent trades 1024 batches at no edge +because position commitment was never adaptive. + +Five-layer architectural completion: + + Layer 1 — Hard CMDP constraints (rl_cmdp_constraints_check.cu) + Session DD cap (-$3.5k = 10% of $35k starting capital). + Cool-down after 10 consecutive losses (500-step lockout). + Max-open positions per batch (= MAX_UNITS). + Net inventory cap (USD). + Override actions to Hold on any violation. + + Layer 2 — IQN risk-averse action selection (rl_iqn_action_tau_controller.cu + + rl_ensemble_action_value.cu modification) + Reads adaptive τ from session drawdown signal. + Action selection: argmax(Q(s, a, τ_action)) instead of argmax(E_τ[Q]). + Normal trading: τ=0.5 (median, ≈ expected-Q behavior). + Under drawdown: τ → MIN=0.1 (downside-aware decisions). + HIGHEST-LEVERAGE LAYER — uses existing IQN distribution for free. + + Layer 3 — Inventory penalty (rl_inventory_beta_controller.cu + reward shaping) + β-scaled penalty on |net_position| in fused reward pipeline. + β adaptive from observed inventory variance vs reward magnitude. + Targets ~1% reward reduction at 2σ inventory. + + Layer 4 — Kelly-fraction sizing (rl_kelly_fraction_controller.cu) + Half-Kelly from observed win_rate × R-multiple EMAs. + Multiplies lot count in actions_to_market_targets. + Warmup: held at bootstrap=1.0 until 1000 closed trades accumulate. + + Layer D — TrailTighten/Loosen wiring (actions_to_market_targets.cu) + Closes pearl_dead_trail_stop_actions_a7_a8 — 10% policy mass was + wasted on no-ops. Now mutates unit_trail_distance_d by ISV-driven + factors (bootstrap 0.9 tighten / 1.1 loosen). + +22 new ISV slots, RL_SLOTS_END 662 → 684. 7 new kernels, 6 modified files. + +Spec: docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md +Plan: docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md + +Validation: + - cargo build --release: clean + - cargo build --tests: clean + - 12 GPU-oracle invariant tests (G1-G12 in risk_stack_invariants.rs): pass + - integrated_trainer_smoke: passes + - 1k local smoke b=128: completes_clean, no NaN, all 5 layers' ISVs + populated and tracking sensibly + - compute-sanitizer memcheck b=128 5 steps: 0 errors + +Cluster validation (G13 fold 1 re-run + G14 ksll2 regression) submitted +as follow-up runs. + +🤖 Generated with [Claude Code](https://claude.com/claude-code) + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +- [ ] **Step 3: Push to origin** + +```bash +git push origin ml-alpha-adaptive-controller-floors 2>&1 | tail -3 +git rev-parse HEAD +``` + +Expected: pushed, new commit SHA returned. + +--- + +### Task 13: Cluster validation (G13 fold 1 + G14 ksll2) + +- [ ] **Step 1: Submit fold 1 walk-forward at new SHA** + +```bash +cd /home/jgrusewski/Work/foxhunt +NEW_SHA=$(git rev-parse HEAD) +./scripts/argo-alpha-rl.sh \ + --sha "$NEW_SHA" \ + --branch ml-alpha-adaptive-controller-floors \ + --gpu-pool ci-training-l40s \ + --n-steps 20000 --seq-len 32 --n-backtests 1024 --per-capacity 32768 \ + --seed 16962 --instrument-mode all \ + --n-folds 3 --fold-idx 1 --n-eval-steps 2000 \ + --skip-push-check +``` + +- [ ] **Step 2: Submit ksll2 regression in parallel** (after fold 1 starts) + +Wait for fold 1 to be Running. Submit ksll2 to same L40S pool — it will queue. Or submit to H100 if needed (note PVC RWO constraint from `pearl_h100_pvc_rwo_blocks_parallel_gpu_workloads`). + +```bash +./scripts/argo-alpha-rl.sh \ + --sha "$NEW_SHA" \ + --branch ml-alpha-adaptive-controller-floors \ + --gpu-pool ci-training-l40s \ + --n-steps 20000 --seq-len 32 --n-backtests 1024 --per-capacity 32768 \ + --seed 16962 --instrument-mode all \ + --n-folds 1 --fold-idx 0 --n-eval-steps 0 \ + --skip-push-check +``` + +- [ ] **Step 3: Wait + collect results** + +Monitor via the Monitor tool. After both terminal: + +Extract fold 1 eval summary from PVC reader: +```bash +# pvc-reader pod pattern as in prior sessions +kubectl -n foxhunt exec pvc-reader -- cat /feature-cache/alpha-rl-runs/$NEW_SHA/fold1/eval_summary.json +``` + +- [ ] **Step 4: G13 verdict** + +Pass conditions (3+ of 6 = PASS): +- eval pnl > -$200k (vs today's -$616k) — **PRIMARY** +- profit_factor > 0.90 (vs 0.83) +- max_drawdown_pct < 15% (vs 18.7%) +- Kelly fraction trajectory: ends < 1.0 if observed edge negative +- IQN τ trajectory: ends < 0.5 if drawdown occurred +- Inventory β > 0 in late training (Layer 3 active) + +- [ ] **Step 5: G14 verdict** + +ksll2 regression: final pnl ≥ +$13M (within 20% of today's +$15.81M). + +- [ ] **Step 6: Update spec with cluster results + write pearls** + +Save outcome to: +- `project_adaptive_risk_management_session_.md` +- Update `MEMORY.md` index +- Decision: merge to main now (if both gates pass) OR iterate (if G13 marginal/fails) + +--- + +## Self-review + +### Coverage check +- Spec scope (5 layers) → 1 task per layer (Tasks 3-7), shared infrastructure tasks (1, 2, 8, 9), validation (10-13) ✓ +- Every controller branch in spec has a kernel in Task 2/3/4/5/6 ✓ +- Every ISV slot in spec is allocated in Task 1 ✓ +- Spec's validation gates G1-G14 all mapped to Tasks 10-13 ✓ + +### Placeholder scan +- Task 4 Step 0 explicitly delegates "verify kernel location" because the IQN action-selection site requires reading the actual codebase. Acceptable — it's a precondition step, not a hidden TODO. +- Task 10 templates the 12 tests rather than fully writing each. Acceptable per writing-plans for tests that follow the same pattern; the templates name the inputs + expected output for each. + +### Type consistency +- All ISV slot constants used in kernels (Tasks 2-7) match the definitions in Task 1 ✓ +- `actions_to_market_targets.cu` modifications (Tasks 3, 6, 7) all reference consistent slot names ✓ + +### Atomic commit discipline +- Tasks 1-11 build progressively without committing ✓ +- Task 12 is the single commit ✓ +- Per `feedback_no_partial_refactor` — same pattern as the controller-floor refactor ✓ + +### Known open items deferred +- Multi-seed cluster validation (deferred to post-G13 if marginal) +- Regime detection / meta-policy (deferred per research caveat — not production-validated) +- Confidence-to-size scaling (was Approach C, dropped — Layer 2's quantile selection achieves the intent without the multiplicative complexity) + +--- + +## Execution + +Use **subagent-driven-development** (recommended) — Tasks 3-7 are each well-scoped per-layer changes that map naturally to one subagent per task. Two-stage review (spec compliance + code quality) after each. Tasks 9 (trainer integration) and 10 (tests) are higher-context, may dispatch a single subagent for both. Task 12 (commit) and Task 13 (cluster) execute inline by the orchestrator. + +Estimated cumulative wall-time: +- Tasks 1-2: ~30 min (slot allocation + 3 EMA kernels) +- Tasks 3-7: ~3 hours (5 layers, ~30-40 min each via subagents) +- Task 8 (fused integration): ~30 min +- Task 9 (trainer integration): ~45 min — most complex, may warrant subagent +- Tasks 10-11: ~1 hour (tests + smoke) +- Task 12: ~10 min (commit) +- Task 13: ~2-3 hours (cluster validation, two 20k runs sequential on L40S) + +Total: ~8 hours of focused work + cluster wait time. diff --git a/docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md b/docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md new file mode 100644 index 000000000..33a121e6d --- /dev/null +++ b/docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md @@ -0,0 +1,560 @@ +# Adaptive Risk Management — Design (Layered Stack) + +**Goal**: close the train-eval distribution-shift gap surfaced by the 2026-05-30 adaptive-controller-floor refactor (commit `b1ef6664a`). The fix preserves the +$18M baseline (G9 PASS, +$15.81M) but fold 1 walk-forward eval came in marginally worse (-$616k vs -$500k OLD) despite massively better train pnl. **Eval JSONL revealed the load-bearing missing layers**: position commitment, action selection risk-awareness, and inventory exposure are all under-adapted relative to what the controller-floor architecture made possible. + +**Why now**: the eval `win_rate=0.21 × R-multiple=1.8` Kelly fraction is negative — agent should not trade but does (1024 batches). Train win_rate 0.55 vs eval 0.21 = severe distribution shift. The controller-floor refactor made signal *interpretation* adaptive; commitment, action-selection, and reward shaping for inventory remained hardcoded. + +**Out of scope**: changes to the perception encoder, Mamba2/CfC architecture, advantage normalization (Phase 4.5), the 12 adaptive controllers from b1ef6664a (those are now correct). Decision policy (the agent still chooses among the 11 discrete actions). Regime detection / meta-policy switching (deferred to follow-up per research caveat). + +--- + +## Diagnosis — what the JSONL says + +Per fold 1 eval (`/feature-cache/alpha-rl-runs/b1ef6664a/fold1/eval_summary.json`): + +| Metric | Value | Interpretation | +|---|---|---| +| total_pnl_usd | -$616k | losing | +| avg_win | $13,785 | wins are large | +| avg_loss | $7,631 | R-multiple 1.8 (good) | +| win_rate | 0.21 (vs train 0.55) | severe directional overfit | +| max_drawdown_pct | 18.7% | controlled but real | +| profit_factor | 0.83 | losing aggregate | + +**Kelly at observed edge**: f* = (0.21 × 1.8 − 0.79)/1.8 = **−0.23** → don't trade. Agent trades anyway. **Position commitment never adapts to edge.** + +**Internal evidence from in-run JSONL** (fold 1 step 16k, qpa = +0.68 = high agent confidence): **the IQN distribution would tell us the pessimistic Q (lower quantile) is significantly negative even when expected Q is positive**. The agent uses `argmax(E_τ[Q])` (mean over quantiles) — over-weights upside under noisy data. Risk-averse action selection (lower quantile τ) would naturally filter these losing trades. + +**Bigger picture**: the system has no adaptive layer between "agent's decision" and "actual market commitment." Five distinct layers are absent or hardcoded: + +| Layer | What it does | Status pre-spec | +|---|---|---| +| Hard constraints (max position, DD cap, etc.) | Override agent on safety violation | Mostly absent | +| Risk-averse action selection | Bias decisions toward downside-aware Q | Uses E[Q] — risk-neutral | +| Inventory exposure penalty | Discourage one-sided drift | Absent in reward shaping | +| Position commitment / sizing | Scale lots by observed edge | Fixed at base lots | +| Regime detection | Distribution-shift handling | Absent | + +This spec addresses the first four. Regime detection is deferred per research evidence ("research-active, not production-validated"). + +--- + +## Scope — the layered risk stack + +Apply `feedback_adaptive_not_tuned` consistently across the commitment layer: + +1. **Layer 1 — Hard CMDP-style constraints (~120 LOC)**: 5 deterministic risk gates that override agent actions +2. **Layer 2 — IQN risk-averse action selection (~30 LOC)**: use lower-quantile τ from existing IQN distribution +3. **Layer 3 — Inventory penalty in reward shaping (~60 LOC)**: discourage one-sided net-position drift +4. **Layer 4 — Kelly-fraction position sizing (~80 LOC)**: scale committed lots by observed edge +5. **TrailTighten/Loosen wiring (~30 LOC)**: close dead-action gap per [[pearl_dead_trail_stop_actions_a7_a8]] + +Single atomic commit per `feedback_no_partial_refactor`. Same pattern as the controller-floor refactor. + +**The crucial reordering vs an earlier draft of this spec**: Kelly was originally the primary fix. Research (FineFT 2026, distributional-RL surveys) plus internal analysis showed **Layer 2 (IQN risk-averse τ) is the higher-leverage minimum-cost change** because we already compute the full IQN quantile distribution but throw away everything except the expectation. Kelly is now Layer 4 — fine-tuning after pessimistic action selection has filtered the trade set. + +--- + +## Approaches considered + +### Approach A — Kelly fraction only +Original minimum-viable design. Single new controller computes f* from observed win_rate × R-multiple, multiplies lot size. ~80 LOC. + +**Verdict**: Insufficient. Kelly is the *commitment* layer; the *selection* layer (which trades to even consider) is more impactful. The 21% eval win_rate means the agent shouldn't even be entering most trades, regardless of size. + +### Approach B — Kelly + session DD circuit breaker +Add a session-pnl-tracking ISV slot + override-to-Hold logic. Two-layer (Kelly per-trade + session-wide cap). + +**Verdict**: Closer but still missing the action-selection layer. Session DD CB triggers AFTER damage is done. + +### Approach C — B + confidence-to-size scaling +Add Q-distribution magnitude as continuous size multiplier alongside Kelly. + +**Verdict**: Three multiplicative layers (Kelly × confidence × base) become hard to reason about. Confidence is implicit in Layer 2's quantile selection. + +### Approach D — Add TrailTighten/Loosen wiring to any of A/B/C +Close the 10%-policy-mass-wasted gap by activating dead actions. + +**Verdict**: Orthogonal to risk management. Worth doing but small impact. + +### Approach E — Layered risk stack (RECOMMENDED) + +Replace the Kelly-first design with five distinct layers, each addressing a different risk concern: + +| Layer | Concern | ISV-driven? | Production validation | +|---|---|---|---| +| 1. Hard CMDP gates | Catastrophic-loss safety (session DD, max position, max inventory) | ✓ | Highest (standard in live HFT) | +| 2. IQN risk-averse τ | Downside-aware action selection | ✓ (adaptive τ) | Research, but our IQN gives it for free | +| 3. Inventory penalty | Adverse selection, one-sided drift | ✓ (adaptive β) | Highest (Avellaneda-Stoikov is canonical HFT) | +| 4. Kelly sizing | Edge-aware commitment | ✓ | Lower (research stage in RL context) | +| (D) TrailTighten/Loosen | Close dead-action gap | ✓ (factor slots) | N/A — wiring fix | + +**Why E**: research evidence (FineFT KDD 2026, plus broader distributional-RL literature) supports the layered approach. Each layer addresses a distinct failure mode visible in the JSONL. Combined, they form the "deployment-grade risk stack" pattern. + +**Total scope**: ~320 LOC across 5 new + 5 modified files. Same atomic-commit discipline as the prior refactor. + +--- + +## Detailed design (Approach E) + +### Layer 1 — Hard CMDP constraints + +Deterministic safety gates that override agent actions in `actions_to_market_targets.cu`. NOT learned by the agent. Each constraint is an independent ISV slot, AND'd together — any single trigger forces Hold (or partial-flatten). + +**New ISV slots** (allocated starting at `RL_SLOTS_END = 662`): + +| Slot | Name | Purpose | Bootstrap | +|---|---|---|---| +| 662 | `RL_SESSION_PNL_USD_INDEX` | Cumulative pnl since session boundary | 0.0 | +| 663 | `RL_SESSION_DD_LIMIT_USD_INDEX` | Drawdown cap (USD, negative) | -3500.0 (= 10% × $35k) | +| 664 | `RL_SESSION_DD_TRIGGERED_INDEX` | Sticky flag, set when DD breached | 0.0 | +| 665 | `RL_MAX_OPEN_UNITS_INDEX` | Max concurrent open positions per batch | 4.0 (= current `MAX_UNITS`) | +| 666 | `RL_CONSEC_LOSS_LIMIT_INDEX` | Max consecutive losses before cool-down | 10.0 | +| 667 | `RL_CONSEC_LOSS_COUNT_INDEX` | Running consecutive-loss count | 0.0 | +| 668 | `RL_COOLDOWN_REMAINING_STEPS_INDEX` | Cool-down period after limit hit | 0.0 | +| 669 | `RL_COOLDOWN_DURATION_INDEX` | Default cool-down length (steps) | 500.0 | +| 670 | `RL_NET_INVENTORY_LIMIT_USD_INDEX` | Cap on net one-sided exposure | 10000.0 | + +**Kernel: `rl_cmdp_constraints_check.cu`** (~100 LOC, single-thread single-block, runs once per step before `actions_to_market_targets`) + +```c +// Layer 1: hard CMDP gates. Output: writes booleans to ISV slots that +// actions_to_market_targets reads as overrides. +// +// Five independent constraints, ALL must be satisfied (= no override) for +// the agent's chosen action to be executed. Any violation → force Hold. + +extern "C" __global__ void rl_cmdp_constraints_check( + float* __restrict__ isv, + const float* __restrict__ realized_pnl_per_batch, // [b_size] + const float* __restrict__ net_position_per_batch, // [b_size] (signed contracts) + const u8* __restrict__ unit_active_per_batch, // [b_size × MAX_UNITS] + const int* __restrict__ trade_outcome_per_batch, // [b_size] −1/0/+1 + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // 1. Session pnl accumulation + DD check. + float step_pnl = 0.0f; + for (int b = 0; b < b_size; b++) step_pnl += realized_pnl_per_batch[b]; + const float session_pnl = isv[RL_SESSION_PNL_USD_INDEX] + step_pnl; + isv[RL_SESSION_PNL_USD_INDEX] = session_pnl; + if (session_pnl < isv[RL_SESSION_DD_LIMIT_USD_INDEX]) { + isv[RL_SESSION_DD_TRIGGERED_INDEX] = 1.0f; + } + + // 2. Cool-down decrement. + float cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX]; + if (cooldown > 0.0f) { + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = cooldown - 1.0f; + } + + // 3. Consecutive loss tracking. Outcomes: -1 = loss, 0 = open/hold, +1 = win. + float consec = isv[RL_CONSEC_LOSS_COUNT_INDEX]; + for (int b = 0; b < b_size; b++) { + const int outcome = trade_outcome_per_batch[b]; + if (outcome == -1) consec += 1.0f; + else if (outcome == 1) consec = 0.0f; + // outcome == 0 (no change): consec unchanged + } + isv[RL_CONSEC_LOSS_COUNT_INDEX] = consec; + if (consec >= isv[RL_CONSEC_LOSS_LIMIT_INDEX]) { + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = isv[RL_COOLDOWN_DURATION_INDEX]; + isv[RL_CONSEC_LOSS_COUNT_INDEX] = 0.0f; // reset after triggering + } +} +``` + +Override logic in `actions_to_market_targets.cu`: + +```c +const bool dd_triggered = isv[RL_SESSION_DD_TRIGGERED_INDEX] >= 0.5f; +const bool in_cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] > 0.0f; +const bool override_hold = dd_triggered || in_cooldown; + +if (override_hold) { + actions[b] = ACTION_HOLD; + target_lots[b] = 0; + return; // skip rest of action translation +} + +// Max-open-units check (per-batch, dynamic): +int active_units = 0; +for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active[b * MAX_UNITS + u]) active_units++; +} +if (active_units >= (int)isv[RL_MAX_OPEN_UNITS_INDEX] && + is_opening_action(actions[b])) { + actions[b] = ACTION_HOLD; // refuse new opens, allow closes +} + +// Inventory cap check: +const float net_pos_usd = compute_net_position_usd(b); +if (fabsf(net_pos_usd) > isv[RL_NET_INVENTORY_LIMIT_USD_INDEX] && + increases_one_sided_exposure(actions[b], net_pos_usd)) { + actions[b] = ACTION_HOLD; +} +``` + +**Session boundary**: trainer resets `RL_SESSION_PNL_USD_INDEX`, `RL_SESSION_DD_TRIGGERED_INDEX`, `RL_CONSEC_LOSS_COUNT_INDEX` to 0 on each fold boundary (start of train phase, start of eval phase). + +### Layer 2 — IQN risk-averse action selection + +The genuinely-novel insight from research: **we already compute the IQN quantile distribution**. The standard action selector picks `argmax(E_τ[Q(s, a, τ)])` — discards the distribution shape. Instead, pick `argmax(Q(s, a, τ_action))` with adaptive `τ_action ∈ (0, 0.5]` — lower τ = more pessimistic = downside-aware decisions. + +**New ISV slots**: + +| Slot | Name | Purpose | Bootstrap | +|---|---|---|---| +| 671 | `RL_IQN_ACTION_TAU_INDEX` | Quantile used for action selection ∈ (0, 1] | 0.5 (= median, near current E[Q] behavior) | +| 672 | `RL_IQN_ACTION_TAU_MIN_INDEX` | Lower bound | 0.1 | +| 673 | `RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX` | How fast τ drops with drawdown | 5.0 | + +**Adaptive controller `rl_iqn_action_tau_controller.cu`** (~50 LOC): + +```c +// Adapt action-selection quantile based on recent drawdown. +// Normal trading: τ = 0.5 (median ≈ expected-Q-like behavior). +// In drawdown: τ → MIN (pessimistic — only take trades where even +// the worst-case Q is positive). +// +// Recent drawdown signal: (session_pnl - max_session_pnl) / |starting_capital| +// Bounded in [0, 1]: 0 = at peak, 1 = total wipeout. +// τ_action = max(τ_min, 0.5 - sensitivity × drawdown_frac) + +extern "C" __global__ void rl_iqn_action_tau_controller( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float session_pnl = isv[RL_SESSION_PNL_USD_INDEX]; + const float max_pnl = isv[RL_SESSION_MAX_PNL_USD_INDEX]; // tracked separately + const float capital = isv[RL_STARTING_CAPITAL_USD_INDEX]; + + const float drawdown_usd = fmaxf(0.0f, max_pnl - session_pnl); + const float drawdown_frac = drawdown_usd / fmaxf(capital, 1.0f); + + const float tau_min = isv[RL_IQN_ACTION_TAU_MIN_INDEX]; + const float sensitivity = isv[RL_IQN_ACTION_TAU_DD_SENSITIVITY_INDEX]; + const float tau_target = fmaxf(tau_min, 0.5f - sensitivity * drawdown_frac); + isv[RL_IQN_ACTION_TAU_INDEX] = tau_target; +} +``` + +**Modification to existing action selection kernel** (e.g., `rl_ensemble_action_value.cu` or wherever the argmax over Q happens): + +```c +// BEFORE: +const float exp_q = compute_expected_q(s, a); // mean over sampled τ +// argmax over actions on exp_q + +// AFTER: +const float tau_a = isv[RL_IQN_ACTION_TAU_INDEX]; // adaptive, default 0.5 +const float q_at_tau = compute_q_at_quantile(s, a, tau_a); +// argmax over actions on q_at_tau +``` + +The IQN kernel already computes `Q(s, a, τ)` for sampled `τ`. Selecting a specific quantile is a single-line indexing change. + +**Critical interaction with C51 ensemble**: the agent uses C51+IQN ensemble. The ensemble blend (α) is already ISV-driven. We apply the risk-averse τ to the IQN side only, then blend as usual. The C51 side continues to use expected value because C51 doesn't naturally expose quantiles. + +### Layer 3 — Inventory penalty (Avellaneda-Stoikov) + +Add a reward shaping term proportional to `|net_position_per_batch|`. Discourages systematic one-sided exposure that bleeds via adverse selection. + +**New ISV slots**: + +| Slot | Name | Purpose | Bootstrap | +|---|---|---|---| +| 674 | `RL_INVENTORY_PENALTY_BETA_INDEX` | Penalty coefficient per |net_position| unit | 0.0 (sentinel — adaptive on first observation) | +| 675 | `RL_INVENTORY_VARIANCE_EMA_INDEX` | Observed inventory variance for β scaling | 0.0 sentinel | + +**Modification to `rl_fused_reward_pipeline.cu`** (~30 LOC): + +```c +// After raw reward computation, before clamps: +const float net_pos = net_position_per_batch[b]; // signed contracts +const float beta = isv[RL_INVENTORY_PENALTY_BETA_INDEX]; +const float penalty = beta * fabsf(net_pos); // always reduces reward +shaped_reward = raw_reward - penalty; +``` + +**Adaptive controller `rl_inventory_beta_controller.cu`** (~50 LOC) — derives β from observed inventory variance: + +```c +// β should scale so penalty is meaningful when inventory drifts but +// negligible during balanced trading. +// Target: penalty ≈ 1% of typical absolute reward when inventory at ~2σ. +// +// β = 0.01 × E[|reward|] / (2 × sqrt(var(inventory))) +// +// Sentinel-zero: hold β = 0 until variance EMA accumulates real signal. + +const float reward_mag = isv[RL_REWARD_MAGNITUDE_EMA_INDEX]; // existing slot 614 +const float inv_var = isv[RL_INVENTORY_VARIANCE_EMA_INDEX]; +if (inv_var <= 0.0f || reward_mag <= 0.0f) return; + +const float inv_std = sqrtf(inv_var); +const float beta_target = 0.01f * reward_mag / (2.0f * inv_std); +// Wiener blend with floor 0.4 (existing pearl pattern) +isv[RL_INVENTORY_PENALTY_BETA_INDEX] = wiener_blend( + isv[RL_INVENTORY_PENALTY_BETA_INDEX], + beta_target, + isv[RL_WIENER_ALPHA_FLOOR_INDEX] +); +``` + +### Layer 4 — Kelly-fraction position sizing + +The original spec's contribution, demoted to fine-tuning layer. Operates *after* Layer 2 has filtered actions and Layer 1 has applied safety gates. Sizes the trades that survive. + +**New ISV slots**: + +| Slot | Name | Purpose | Bootstrap | +|---|---|---|---| +| 676 | `RL_KELLY_FRACTION_INDEX` | Position-size multiplier ∈ [0, 1] | 1.0 (full size pre-warmup) | +| 677 | `RL_WIN_RATE_EMA_INDEX` | Observed win_rate EMA | 0.5 (neutral sentinel) | +| 678 | `RL_AVG_WIN_USD_EMA_INDEX` | Observed avg_win | 0.0 sentinel | +| 679 | `RL_AVG_LOSS_USD_EMA_INDEX` | Observed avg_loss | 0.0 sentinel | +| 680 | `RL_KELLY_SAFETY_FRAC_INDEX` | Half-Kelly multiplier (Thorp) | 0.5 | +| 681 | `RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX` | Warmup gate | 1000.0 | + +**Kernel `rl_kelly_fraction_controller.cu`** (~80 LOC) — same design as original spec, repeated here for completeness: + +```c +extern "C" __global__ void rl_kelly_fraction_controller(float* __restrict__ isv) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Warmup: hold at bootstrap until enough closed trades. + if (isv[RL_CUMULATIVE_DONES_INDEX] < isv[RL_KELLY_MIN_TRADES_FOR_RELEASE_INDEX]) { + isv[RL_KELLY_FRACTION_INDEX] = 1.0f; + return; + } + + const float p = isv[RL_WIN_RATE_EMA_INDEX]; + const float avg_win = isv[RL_AVG_WIN_USD_EMA_INDEX]; + const float avg_loss = isv[RL_AVG_LOSS_USD_EMA_INDEX]; + if (avg_loss <= 0.0f) { isv[RL_KELLY_FRACTION_INDEX] = 1.0f; return; } + + const float b = avg_win / avg_loss; + const float f_kelly = (p * b - (1.0f - p)) / b; + const float safety = isv[RL_KELLY_SAFETY_FRAC_INDEX]; + float f = f_kelly * safety; + f = fmaxf(0.0f, fminf(f, 1.0f)); + isv[RL_KELLY_FRACTION_INDEX] = f; +} +``` + +**Modification in `actions_to_market_targets.cu`** (~3 LOC): + +```c +// After Layer 1 overrides have been applied (DD CB, cooldown, etc.): +const float kelly = isv[RL_KELLY_FRACTION_INDEX]; +target_lots[b] = (int)ceilf((float)target_lots[b] * kelly); +// If kelly = 0, target_lots = 0 = no commitment. +``` + +### Layer D — TrailTighten/Loosen wiring + +Resolve the 10%-policy-mass-wasted gap per [[pearl_dead_trail_stop_actions_a7_a8]]: + +**New ISV slots**: + +| Slot | Name | Purpose | Bootstrap | +|---|---|---|---| +| 682 | `RL_TRAIL_TIGHTEN_FACTOR_INDEX` | Multiplier on trail distance for a7 | 0.9 (tighten by 10%) | +| 683 | `RL_TRAIL_LOOSEN_FACTOR_INDEX` | Multiplier on trail distance for a8 | 1.1 (loosen by 10%) | + +**Modification to `actions_to_market_targets.cu`** wires a7/a8 to actually mutate `unit_trail_distance_d`: + +```c +case ACTION_TRAIL_TIGHTEN: { + const float factor = isv[RL_TRAIL_TIGHTEN_FACTOR_INDEX]; + for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active_d[b * MAX_UNITS + u]) { + unit_trail_distance_d[b * MAX_UNITS + u] *= factor; + } + } + break; +} +case ACTION_TRAIL_LOOSEN: { + const float factor = isv[RL_TRAIL_LOOSEN_FACTOR_INDEX]; + for (int u = 0; u < MAX_UNITS; u++) { + if (unit_active_d[b * MAX_UNITS + u]) { + unit_trail_distance_d[b * MAX_UNITS + u] *= factor; + } + } + break; +} +``` + +--- + +## Slot allocation summary + +22 new ISV slots, `RL_SLOTS_END` 662 → 684. + +| Range | Layer | Purpose | +|---|---|---| +| 662–670 | 1 (CMDP gates) | Session pnl, DD limit, max open, consec loss, cooldown, inventory cap (9 slots) | +| 671–673 | 2 (IQN τ) | Action τ, τ_min, DD sensitivity (3 slots) | +| 674–675 | 3 (inventory penalty) | β + inventory variance EMA (2 slots) | +| 676–681 | 4 (Kelly) | Kelly fraction + EMAs + safety + warmup (6 slots) | +| 682–683 | D (trail) | Tighten/loosen factors (2 slots) | + +Memory impact: 88 bytes additional (negligible). + +--- + +## Per-step pipeline order + +The per-step trainer pipeline must execute Layers in this order: + +1. EMA producers (existing — kl_pi_ema, q_divergence_ema, etc., PLUS new: win_rate_ema, avg_win_ema, avg_loss_ema, inventory_variance_ema) +2. Welford variance updates (existing kernel from b1ef6664a) +3. Adaptive controllers (existing 12 — fused kernel) +4. **Layer 1 — `rl_cmdp_constraints_check`** (new) +5. **Layer 2 — `rl_iqn_action_tau_controller`** (new — adapts τ from drawdown) +6. **Layer 3 — `rl_inventory_beta_controller`** (new — adapts β) +7. **Layer 4 — `rl_kelly_fraction_controller`** (new) +8. **Layer 5 — TrailTighten/Loosen ISV slots** (no per-step kernel; ISV slots read during action translation) +9. Action selection — IQN ensemble uses adaptive τ for argmax +10. `actions_to_market_targets` — applies Layer 1 overrides + Layer 4 Kelly scaling + Layer D trail mutations +11. Reward computation — `rl_fused_reward_pipeline` applies Layer 3 inventory penalty +12. PPO/DQN updates as usual + +--- + +## Files + +### New + +- `crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu` (~100 LOC) +- `crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu` (~50 LOC) +- `crates/ml-alpha/cuda/rl_inventory_beta_controller.cu` (~50 LOC) +- `crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu` (~80 LOC) +- `crates/ml-alpha/cuda/rl_win_rate_ema_update.cu` (~40 LOC if not derivable) +- `crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu` (~50 LOC) +- `crates/ml-alpha/cuda/rl_inventory_variance_update.cu` (~40 LOC) +- `crates/ml-alpha/tests/risk_stack_invariants.rs` (~10 GPU-oracle gates) + +### Modified + +- `crates/ml-alpha/src/rl/isv_slots.rs` — +22 slots, `RL_SLOTS_END` 662 → 684 +- `crates/ml-alpha/build.rs` — register new cubins +- `crates/ml-alpha/cuda/actions_to_market_targets.cu` — Layer 1 overrides + Layer 4 Kelly scaling + Layer D trail mutations +- `crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu` — Layer 3 inventory penalty term +- `crates/ml-alpha/cuda/rl_ensemble_action_value.cu` (or wherever IQN action selection happens) — Layer 2 quantile-τ argmax +- `crates/ml-alpha/cuda/rl_fused_controllers.cu` — add new branches for τ, β, Kelly (matching the fused-kernel pattern from b1ef6664a) +- `crates/ml-alpha/src/trainer/integrated.rs` — load new cubins, wire per-step launches, bootstrap entries, session reset on fold boundary + +### Not touched + +- 12 adaptive controllers from b1ef6664a (orthogonal layer) +- Welford-variance kernel infrastructure (reused) +- Phase 4.5 advantage normalization +- Perception encoder, Mamba2, CfC heads +- Decision policy gradient (the agent still picks discrete actions) + +--- + +## Validation gates + +### Local unit-level (RTX 3050 Ti) + +**G1 — CMDP DD trigger**: Write `session_pnl = -4000` (below limit -3500) → assert triggered=1, agent actions overridden to Hold. + +**G2 — CMDP cooldown**: Trigger 10 consecutive losses → assert cooldown counter set to 500, agent actions overridden until counter reaches 0. + +**G3 — CMDP max-open**: Pre-set 4 active units, agent picks opening action → assert action forced to Hold. + +**G4 — CMDP inventory cap**: Set `net_position = $11k` (above $10k limit), agent picks opening-same-side action → assert Hold. + +**G5 — IQN τ adapts to drawdown**: Set `session_pnl = -1500` (≈4% drawdown of $35k capital) → assert `τ_action ≈ max(0.1, 0.5 - 5 × 0.043) ≈ 0.285` (mid-pessimistic). + +**G6 — IQN τ at no-drawdown**: Set `session_pnl = max_pnl` → assert `τ_action = 0.5`. + +**G7 — Inventory penalty β adapts**: Set `reward_mag_ema = 0.1`, `inv_variance_ema = 4.0` (σ=2) → assert `β = 0.01 × 0.1 / (2 × 2) = 0.00025`. Negligible at balanced inventory, real at high. + +**G8 — Kelly math at known edge**: Feed `win_rate=0.55, avg_win=1.8, avg_loss=1.0, trades=2000` → expect `f = 0.5 × ((0.55×1.8 - 0.45)/1.8) = 0.5 × 0.30 = 0.15`. + +**G9 — Kelly returns 0 at negative edge**: Feed `win_rate=0.21, avg_win=1.8, avg_loss=1.0, trades=2000` → expect `f = 0` (Kelly negative → clamped). + +**G10 — TrailTighten reduces distance**: Pre-set `unit_trail_distance = 100`, agent picks a7 → assert distance becomes 90. + +**G11 — integrated_trainer_smoke** still passes end-to-end. + +**G12 — 1k local smoke at b=128**: completed_clean, no NaN, all 5 layers' ISV slots populated and tracking sensibly. + +### Cluster + +**G13 — Walk-forward fold 1 re-run** (same params as today's 62x7p, new SHA). Combined layers should close the eval gap. + +Targets (each is a soft pass; combination of 3+ passes = G13 PASS): +- eval pnl > -$200k (better than today's -$616k, **PRIMARY**) +- profit_factor > 0.90 (vs 0.83) +- max_drawdown_pct < 15% (vs 18.7%) +- Kelly fraction: ends < 1.0 if observed edge negative (shows Layer 4 active) +- IQN τ: ends < 0.5 if any drawdown occurred (shows Layer 2 adaptive) +- DD breaker: fires when expected (or never if pnl stays positive) +- Inventory penalty β: > 0 throughout (shows Layer 3 adapted off bootstrap) + +**G14 — ksll2 regression test**: same params as today's rwk9t. Final pnl ≥ +$13M (within 20% of +$15.81M). Confirms layered stack doesn't regress large-data regime. + +--- + +## Open questions + +1. **Starting capital ($35k) per [[project_ml_alpha_starting_capital]]** — should we add `RL_STARTING_CAPITAL_USD_INDEX` as an ISV slot? Yes, makes session DD cap and Layer 2 drawdown ratio both adaptive to capital changes. Add to Layer 1. + +2. **Layer 2 sensitivity (5.0)** — at session DD 10%, τ drops from 0.5 to 0. At 4% drawdown, τ ≈ 0.3. Reasonable? Tunable via ISV. May need empirical sweep. + +3. **Inventory variance EMA — measured over what window?** Need to choose Wiener-α. Default `RL_WIENER_ALPHA_FLOOR_INDEX = 0.4` from existing pearl. Adaptive could come later. + +4. **Layer 3 β bootstrap (0.0 = sentinel)** — agent gets ZERO inventory penalty until variance EMA accumulates. Acceptable cold-start. + +5. **Should `RL_MAX_OPEN_UNITS_INDEX` = 4 (current `MAX_UNITS`) be redundant?** The action space already enforces this via unit_active tracking. Yes, redundant but cheap; treats it as belt-and-braces. + +6. **Approach E vs the original B+D — anything we lose?** The `rl_session_drawdown_check.cu` kernel from the prior design is now folded into the broader `rl_cmdp_constraints_check.cu`. Same functionality, more layered. + +7. **Multi-seed validation?** A single fold-1 result is noisy (1024 trades, high variance). Worth running 3 seeds to confirm. Defer to post-validation cluster sweep if G13 single-seed result is marginal. + +8. **Regime detection (Approach E+1)?** FineFT (KDD 2026) demonstrates risk-aware ensemble routing on regime changes. Heavyweight implementation (online clustering + meta-controller). Defer to follow-up spec if Approach E doesn't close the eval gap. + +--- + +## Risks and mitigations + +- **Layer 2 τ-quantile path may not be exposed by current IQN kernel signature** → audit `rl_ensemble_action_value.cu` first; if `Q(s, a, τ)` isn't readily indexable, may need kernel restructure (small change). + +- **Kelly EMA estimates unreliable in early training** → warmup gate (1000 trades) holds Kelly at bootstrap. Same pattern as b1ef6664a's reward_floor cumulative-dones gate. + +- **Five layers compounding may make agent ultra-conservative — eval might be 0 with no trades** → that's actually fine for G13 (eval pnl > -$200k bar). But for G14 (ksll2 +$13M), need to verify layers don't over-restrict on full-data regime. Mitigation: warmup gates default to "behave like the prior code" until enough signal accumulates. + +- **Session boundary resets must be wired by trainer** → must call ISV resets on train→eval transition. Test G11 catches this. Plan task explicitly enumerates the reset points. + +- **Approach D wiring changes action distribution dynamics** → trainable but worth a regression check on entropy. 1k smoke catches it. + +- **CMDP "force Hold" creates discontinuity in policy gradient** → standard technique in CMDP-RL literature, well-understood. Override doesn't update gradient through the constraint, only through the Hold action's natural Q. + +--- + +## Next step after approval + +Write the plan: `docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md`. Task-level decomposition: + +1. ISV slot allocation (22 slots) +2. Layer 1 (CMDP constraints kernel + actions_to_market_targets overrides) +3. Layer 2 (IQN τ controller + action-selection kernel modification) +4. Layer 3 (inventory penalty in reward + β controller) +5. Layer 4 (Kelly controller + actions_to_market_targets sizing) +6. Layer D (TrailTighten/Loosen wiring) +7. New EMA producers (win_rate, avg_win, avg_loss, inventory variance) +8. Fused-kernel integration of new controllers +9. Trainer integration (per-step launches + bootstrap + session reset) +10. GPU-oracle invariant tests (G1-G11) +11. 1k local smoke + sanitizer (G12) +12. Atomic commit +13. Cluster G13 (fold 1 re-run) + G14 (ksll2 regression) + +Expected scope: ~320 LOC across 8 new + 7 modified files, ~13 tasks, single atomic commit. Same `feedback_no_partial_refactor` discipline.