feat(rl): adaptive risk-management stack — 5 layers, all ISV-driven
Layered risk stack per spec docs/superpowers/specs/2026-05-30-adaptive-risk-management-design.md:
Layer 1 (CMDP) hard session-DD, cooldown, max-open, inventory limits
Layer 2 (IQN τ) risk-averse action selection adapts to session drawdown
Layer 3 (Inventory) Avellaneda-Stoikov penalty β scales with reward magnitude
and inventory variance
Layer 4 (Kelly) half-Kelly fraction sizing from observed win-rate +
R-multiple, warmup-gated until cumulative_dones >= 1000
Layer D (Trail) wire dead a7/a8 (TrailTighten/Loosen) via independent
ISV factors, replacing the symmetric reciprocal
Architecture: every threshold ISV-driven (22 new slots, RL_SLOTS_END 662→684).
Every adaptive bound follows the canonical Wiener-α blend with floor 0.4,
sentinel-zero bootstrap, and asymmetric Schulman where applicable.
Kernels:
rl_cmdp_constraints_check session pnl + cooldown + consec-loss tracking
rl_iqn_action_tau_controller τ = clamp(0.5 - 5·dd_frac, τ_min, 1.0)
rl_inventory_beta_controller β_target = 0.01·E|reward| / (2·σ_inventory)
rl_kelly_fraction_controller f = clamp(safety · (p·b - q)/b, 0, 1)
rl_win_rate_ema_update closed-trade win-rate EMA from rewards + dones
rl_avg_win_loss_ema_update separate avg-win and avg-loss EMAs
rl_inventory_variance_update Welford variance of net-position-per-batch
Integration:
- All 7 new cubins loaded in IntegratedTrainer + launched per step in spec order
(CMDP after reward pipeline, before actions_to_market_targets reads override
flags; Layer 2/3/4 controllers in rl_fused_controllers.cu).
- actions_to_market_targets.cu: Layer 1 hard overrides (DD-triggered → 0 lots;
cooldown → 0 lots; max-open → block opening actions; inventory cap → block
one-sided expansion) and Layer 4 Kelly fraction scaling on target lots.
- rl_fused_reward_pipeline.cu: Layer 3 inventory penalty term in reward shaping.
- rl_trail_mutate.cu: a7 multiplies trail by RL_TRAIL_TIGHTEN_FACTOR_INDEX,
a8 by RL_TRAIL_LOOSEN_FACTOR_INDEX (was symmetric reciprocal — pearls
pearl_dead_trail_stop_actions_a7_a8).
Validation:
- 12 GPU-oracle invariants pass (tests/risk_stack_invariants.rs):
G1-G4 CMDP, G5-G7 IQN τ, G8-G9 inventory β, G10-G12 Kelly.
- 20/20 trade_management_kernels.rs tests pass (a7/a8 migrated).
- 5/5 controller_adaptive_floors.rs tests still pass.
- integrated_trainer_smoke passes (end-to-end pipeline launch).
- Local 1k smoke b=128: completes 1000/1000 steps, no NaN, controllers steady.
- compute-sanitizer memcheck (5 steps b=128): ERROR SUMMARY: 0 errors.
Plan: docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
77
crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu
Normal file
77
crates/ml-alpha/cuda/rl_avg_win_loss_ema_update.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
}
|
||||
92
crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu
Normal file
92
crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
|
||||
|
||||
57
crates/ml-alpha/cuda/rl_inventory_beta_controller.cu
Normal file
57
crates/ml-alpha/cuda/rl_inventory_beta_controller.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
59
crates/ml-alpha/cuda/rl_inventory_variance_update.cu
Normal file
59
crates/ml-alpha/cuda/rl_inventory_variance_update.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
53
crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu
Normal file
53
crates/ml-alpha/cuda/rl_iqn_action_tau_controller.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
70
crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu
Normal file
70
crates/ml-alpha/cuda/rl_kelly_fraction_controller.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
@@ -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 <stdint.h>
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
63
crates/ml-alpha/cuda/rl_win_rate_ema_update.cu
Normal file
63
crates/ml-alpha/cuda/rl_win_rate_ema_update.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
@@ -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")?;
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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<CudaModule>,
|
||||
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<CudaModule>,
|
||||
rl_cmdp_constraints_check_fn: CudaFunction,
|
||||
_rl_iqn_action_tau_controller_module: Arc<CudaModule>,
|
||||
rl_iqn_action_tau_controller_fn: CudaFunction,
|
||||
_rl_inventory_beta_controller_module: Arc<CudaModule>,
|
||||
rl_inventory_beta_controller_fn: CudaFunction,
|
||||
_rl_kelly_fraction_controller_module: Arc<CudaModule>,
|
||||
rl_kelly_fraction_controller_fn: CudaFunction,
|
||||
// EMA producers feeding the new controllers.
|
||||
_rl_win_rate_ema_update_module: Arc<CudaModule>,
|
||||
rl_win_rate_ema_update_fn: CudaFunction,
|
||||
_rl_avg_win_loss_ema_update_module: Arc<CudaModule>,
|
||||
rl_avg_win_loss_ema_update_fn: CudaFunction,
|
||||
_rl_inventory_variance_update_module: Arc<CudaModule>,
|
||||
rl_inventory_variance_update_fn: CudaFunction,
|
||||
/// Phase 4.4 blended V baselines fed to compute_advantage_return.
|
||||
pub v_blended_d: CudaSlice<f32>,
|
||||
pub v_blended_tp1_d: CudaSlice<f32>,
|
||||
@@ -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<f32>,
|
||||
dones_d: &CudaSlice<f32>,
|
||||
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<f32>,
|
||||
dones_d: &CudaSlice<f32>,
|
||||
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<f32>,
|
||||
dones_d: &CudaSlice<f32>,
|
||||
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<i32>,
|
||||
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.
|
||||
|
||||
393
crates/ml-alpha/tests/risk_stack_invariants.rs
Normal file
393
crates/ml-alpha/tests/risk_stack_invariants.rs
Normal file
@@ -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<CudaStream>, host: &[f32]) -> Result<CudaSlice<f32>> {
|
||||
let mut d = stream.alloc_zeros::<f32>(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(())
|
||||
}
|
||||
@@ -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]
|
||||
);
|
||||
|
||||
|
||||
1185
docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md
Normal file
1185
docs/superpowers/plans/2026-05-30-adaptive-risk-management-plan.md
Normal file
File diff suppressed because it is too large
Load Diff
@@ -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.
|
||||
Reference in New Issue
Block a user