From 20c835713b836214a83fef17b2f191e36c0ee5d4 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 24 May 2026 16:47:53 +0200 Subject: [PATCH] fix(rl): wire TrailTighten/TrailLoosen + SP20 P1+P5 foundation MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit scripts/audit-wiring.sh dogfood pass flagged a7 (TrailTighten) and a8 (TrailLoosen) as actions with no consumer anywhere in the codebase (canonical pearl_dead_trail_stop_actions_a7_a8). Fix bundles SP20 P1 (per-unit trade state buffers) and P5 (trail-stop kernels) since they're the same architectural work. Three new kernels: rl_unit_state_update.cu — per-batch trade state machine. Runs AFTER fill+extract_realized_pnl_delta. Detects open/close/reverse position transitions and populates unit slot 0 with entry_price, entry_step, lots, initial_r, trail_distance. Slots 1-3 allocated for SP20 P7 pyramid expansion but unused this commit. rl_trail_mutate.cu — handles a7/a8 actions. Mutates ALL active units' trail_distance bounded by ISV [MIN, MAX] with symmetric reciprocal adjust rate per SP20 §4.12: a7: trail = max(MIN, trail × rate) a8: trail = min(MAX, trail / rate) rl_trail_stop_check.cu — per-batch per-unit breach check. Reads shared lobsim best book (bid/ask), computes mid, compares to each active unit's (entry ± trail). On breach, OVERRIDE actions[b] to FlatFromLong (a3) or FlatFromShort (a4). Force-close routes through existing flat plumbing per pearl_stop_checks_run_at_deadline_cadence. SP20 v3 §3 P5 calls for routing close via partial-flat (a9/a10) so only the at-risk unit closes — that needs P4 (N_ACTIONS=11). For now, ANY unit's breach closes ENTIRE position via full FlatFromLong/Short. Per-batch per-unit buffers (8 new in trainer): unit_entry_price_d [B × 4] f32 unit_entry_step_d [B × 4] i32 unit_lots_d [B × 4] i32 unit_initial_r_d [B × 4] f32 unit_trail_distance_d[B × 4] f32 unit_active_d [B × 4] u8 pyramid_units_count_d[B] i32 unit_prev_pos_lots_d [B] i32 (state-machine tracker, separate from extract_realized_pnl_delta's prev_position_lots_d for clean kernel composability) 4 new ISV slots (494-497): RL_TRAIL_MIN_INDEX — trail distance floor (seed 0.001) RL_TRAIL_MAX_INDEX — trail distance ceiling (seed 100.0) RL_TRAIL_K_INIT_INDEX — initial trail multiplier (seed 2.0, Turtle 2N) RL_TRAIL_ADJUST_RATE_INDEX — tighten ratio (seed 0.9; symmetric reciprocal for loosen) RL_SLOTS_END: 494 → 498. LobSim exposes bid_px_d() + ask_px_d() public accessors. RlLobBackend trait extended with the two accessors; the LobSimCuda impl wires through. Override stack ordering per SP20 §2.3: 1. rl_pi_action_kernel (sample) 2. rl_trail_mutate (a7/a8 → mutate, before stop check) 3. rl_trail_stop_check (per-unit breach → override action) 4. actions_to_market_targets (execute, including overridden flat) 5. step_fill_from_market_targets 6. extract_realized_pnl_delta 7. rl_unit_state_update (detect post-fill transitions) Audit infrastructure refined as part of dogfooding: * audit-isv allowlist extended for BOOK_LEVELS (structural book depth) and ACTION_* prefix (enum-mirror constants — these are structural API contracts matching src/rl/common.rs::Action positions) * audit-wiring action-handler regex now matches BOTH literal `action == ` and `action == ACTION_` patterns, and treats != as a handler too (a guard against the action is valid wiring) Both `audit-isv.sh` and `audit-wiring.sh` PASS cleanly with the full manifest. audit-diag scheduled for first SP20 phase that adds diag fields (this commit deliberately keeps diag exposure minimal — full per-unit + trail diag blocks come with SP20 P13). Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 3 + crates/ml-alpha/cuda/rl_trail_mutate.cu | 64 +++++++ crates/ml-alpha/cuda/rl_trail_stop_check.cu | 87 +++++++++ crates/ml-alpha/cuda/rl_unit_state_update.cu | 118 ++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 33 +++- crates/ml-alpha/src/rl/reward.rs | 10 + crates/ml-alpha/src/trainer/integrated.rs | 192 ++++++++++++++++++- crates/ml-backtesting/src/sim/mod.rs | 24 +++ scripts/audit-isv.sh | 9 +- scripts/audit-manifest/actions.txt | 4 + scripts/audit-manifest/kernels.txt | 3 + scripts/audit-manifest/slots.txt | 5 + scripts/audit-wiring.sh | 23 ++- 13 files changed, 562 insertions(+), 13 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_trail_mutate.cu create mode 100644 crates/ml-alpha/cuda/rl_trail_stop_check.cu create mode 100644 crates/ml-alpha/cuda/rl_unit_state_update.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 69835aec3..8f0a863ef 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -72,6 +72,9 @@ const KERNELS: &[&str] = &[ "rl_atom_support_update", // audit 2026-05-24 followup: refreshes atom_supports_d from ISV V_MIN/V_MAX so C51 atom span adapts with reward clamp (Q learning was capped at V_MAX=1.0) "rl_q_pi_distill_grad", // audit 2026-05-24 vj5f6 followup: KL(softmax(E_Q/τ) || π_new) gradient ADDED to pi_grad_logits — couples Q's improved C51 calibration to action selection (was decoupled per Option B) "rl_q_distill_lambda_controller",// audit 2026-05-24 rljzl followup: adaptive λ_distill via Schulman bounded step on KL_EMA vs target + "rl_unit_state_update", // SP20 P1+P5 audit fix: per-unit trade state machine — detects open/close/reverse transitions, sets up unit slot 0 entry+trail + "rl_trail_mutate", // SP20 P1+P5 audit fix (a7/a8 dead): TrailTighten/TrailLoosen mutate unit_trail_distance bounded MIN/MAX, symmetric reciprocal adjust + "rl_trail_stop_check", // SP20 P1+P5 audit fix: per-unit trail breach check; OVERRIDE action to FlatFromLong/Short on breach (close routes through existing flat plumbing) ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_trail_mutate.cu b/crates/ml-alpha/cuda/rl_trail_mutate.cu new file mode 100644 index 000000000..fe1397679 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_trail_mutate.cu @@ -0,0 +1,64 @@ +// 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) +// +// Tighten and loosen are reciprocal: N tightens followed by N loosens +// returns trail to its original distance. +// +// Audit history: scripts/audit-wiring.sh dogfood pass flagged that +// a7/a8 had no consumer anywhere in the codebase prior to this +// kernel. +// +// 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"). +// +// 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. + +#include + +#define MAX_UNITS 4 +#define ACTION_TRAIL_TIGHTEN 7 +#define ACTION_TRAIL_LOOSEN 8 +#define RL_TRAIL_MIN_INDEX 494 +#define RL_TRAIL_MAX_INDEX 495 +#define RL_TRAIL_ADJUST_RATE_INDEX 497 + +extern "C" __global__ void rl_trail_mutate( + const int* __restrict__ actions, // [B] + const float* __restrict__ isv, + const unsigned char* __restrict__ unit_active, // [B * MAX_UNITS] + float* __restrict__ unit_trail_distance, // [B * MAX_UNITS] IN/OUT + int b_size +) { + const int b = blockIdx.x; + const int u = threadIdx.x; + if (b >= b_size || u >= MAX_UNITS) return; + + const int action = actions[b]; + if (action != ACTION_TRAIL_TIGHTEN && action != ACTION_TRAIL_LOOSEN) return; + + 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]; + + float updated; + if (action == ACTION_TRAIL_TIGHTEN) { + updated = fmaxf(trail_min, current * adjust_rate); + } 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); + } + unit_trail_distance[idx] = updated; +} diff --git a/crates/ml-alpha/cuda/rl_trail_stop_check.cu b/crates/ml-alpha/cuda/rl_trail_stop_check.cu new file mode 100644 index 000000000..ad38f59f5 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_trail_stop_check.cu @@ -0,0 +1,87 @@ +// rl_trail_stop_check.cu — SP20 P5 trail-stop breach check. +// +// For each batch with an active unit (slot 0 currently used; slots +// 1-3 reserved for SP20 P7 pyramid expansion), check if the current +// market mid moved adversely beyond the unit's trail_distance from +// its entry_price. On breach, OVERRIDE actions[b] to FlatFromLong +// (a3) or FlatFromShort (a4) so actions_to_market_targets force- +// closes the position via existing flat-from-X plumbing. +// +// Per `pearl_stop_checks_run_at_deadline_cadence`: stop check fires +// every step (event rate). Force-close routes through +// actions_to_market_targets → step_fill_from_market_targets → +// apply_fill_to_pos so realized PnL is captured correctly. +// +// SP20 v3 spec §3 P5 calls for routing stop-close through the +// per-unit partial-flat path (a9/a10). That requires P4 (N_ACTIONS=11) +// which is NOT in this commit's scope. For now, ANY unit's trail +// breach closes the ENTIRE position via full FlatFromLong/Short. +// When SP20 P4 ships, the close_unit_index buffer + a9/a10 routing +// per spec §3 P5 supersedes this full-flat fallback. +// +// Market mid sourced from lobsim's bid_px_d[0] + ask_px_d[0] (shared +// across batches — single market). One thread per (batch, unit). +// +// Per `feedback_no_atomicadd`: writes are per-batch (action override +// uses last-thread-wins by writing same value; trail breach by ANY +// unit triggers the same action override regardless of which thread). +// Per `feedback_cpu_is_read_only`: pure device-side. + +#include + +#define MAX_UNITS 4 +#define BOOK_LEVELS 10 +#define ACTION_FLAT_FROM_LONG 3 +#define ACTION_FLAT_FROM_SHORT 4 + +extern "C" __global__ void rl_trail_stop_check( + int* __restrict__ actions, // [B] IN/OUT (may be overridden) + const float* __restrict__ bid_px, // [BOOK_LEVELS] shared best book + const float* __restrict__ ask_px, // [BOOK_LEVELS] + const unsigned char* __restrict__ unit_active, // [B * MAX_UNITS] + const float* __restrict__ unit_entry_price, // [B * MAX_UNITS] + const int* __restrict__ unit_lots, // [B * MAX_UNITS] + const float* __restrict__ unit_trail_distance, // [B * MAX_UNITS] + int b_size +) { + const int b = blockIdx.x; + const int u = threadIdx.x; + if (b >= b_size || u >= MAX_UNITS) return; + + // Compute current mid from shared book (all threads read same). + const float best_bid = bid_px[0]; + const float best_ask = ask_px[0]; + const float mid = 0.5f * (best_bid + best_ask); + + const int idx = b * MAX_UNITS + u; + if (unit_active[idx] == 0) return; + + const int lots = unit_lots[idx]; + if (lots == 0) return; // defensive — active flag should imply lots != 0 + + const float entry = unit_entry_price[idx]; + const float trail = unit_trail_distance[idx]; + + bool breach = false; + int override_action = -1; + if (lots > 0) { + // Long: stop fires when mid drops below entry - trail + if (mid < entry - trail) { + breach = true; + override_action = ACTION_FLAT_FROM_LONG; + } + } else { + // Short: stop fires when mid rises above entry + trail + if (mid > entry + trail) { + breach = true; + override_action = ACTION_FLAT_FROM_SHORT; + } + } + + if (breach) { + // Multiple unit threads may write the same value — that's fine + // since the action is identical (close the whole position via + // full flat per the spec-noted P4 dependency). Last write wins. + actions[b] = override_action; + } +} diff --git a/crates/ml-alpha/cuda/rl_unit_state_update.cu b/crates/ml-alpha/cuda/rl_unit_state_update.cu new file mode 100644 index 000000000..c5df0c438 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_unit_state_update.cu @@ -0,0 +1,118 @@ +// rl_unit_state_update.cu — SP20 P1 per-unit trade state machine. +// +// Runs AFTER each step's fill+extract sequence. Detects per-batch +// position transitions vs the prior step and updates per-unit +// buffers accordingly: +// +// prev_pos == 0 AND curr_pos != 0 → OPEN: write unit slot 0 +// prev_pos != 0 AND curr_pos == 0 → CLOSE: deactivate all units +// sign(prev) != sign(curr) AND both != 0 → REVERSE: close + open +// sign(prev) == sign(curr) AND both != 0 → CONTINUE: refresh lots +// +// This commit uses unit slot 0 only — slots 1-3 are allocated for +// SP20 P7 pyramid expansion but not populated yet. Pyramid add / +// partial-flat logic lives in actions_to_market_targets (P7). +// +// Per `pearl_first_observation_bootstrap`: prev_pos starts at sentinel +// 0; first non-zero curr_pos triggers OPEN via the standard branch. +// No special bootstrap path needed. +// +// Per `feedback_no_atomicadd`: per-batch independent writes only. +// Per `feedback_cpu_is_read_only`: pure device-side update. +// Per `pearl_trade_level_vol_for_stop_distance`: trail initialized +// from `sqrt(realized_return_var_ema)` × ISV-driven k_init. + +#include + +#define MAX_UNITS 4 +#define RL_TRAIL_K_INIT_INDEX 496 +// realized_return_var_ema lives at a known SP19-era slot — reuse +// rather than allocate a new one. Currently RL_ADV_VAR_STREAM_M2_INDEX +// holds the streaming variance running M2; bootstrap derives σ from +// √(M2 / count). For SP20 P1+P5 simplicity, use a fallback constant +// if streaming-var hasn't bootstrapped (avoids div-by-zero collapse). +// +// Practical compromise: use `mean_abs_pnl_ema` (slot 417) as a proxy +// for typical PnL magnitude. trail_init = k_init × mean_abs_pnl_ema / +// position_lots. This is in dollar units which matches vwap_entry's +// price units after the position-lots normalization. +// +// Even simpler: use a fixed bootstrap fraction of vwap_entry as the +// trail distance — k_init × vwap_entry / 1000 = ~0.2% of price for +// k_init=2. This is rough but tradable and avoids dependency on +// streaming-var bootstrap. SP20 P5 follow-up will integrate proper +// vol-derived initialization. +#define BOOTSTRAP_TRAIL_FRACTION 1.0e-3f + +// PosFlat layout (matches crates/ml-backtesting/src/lob/mod.rs::PosFlat): +// offset 0: position_lots: i32 +// offset 4: vwap_entry: f32 +// offset 8: realized_pnl: f32 +// offset 12: peak_equity: f32 + +extern "C" __global__ void rl_unit_state_update( + const unsigned char* __restrict__ pos_state, // [B * pos_bytes] + const float* __restrict__ isv, + int* __restrict__ prev_pos_lots, // [B] tracking + float* __restrict__ unit_entry_price, // [B * MAX_UNITS] + int* __restrict__ unit_entry_step, // [B * MAX_UNITS] + int* __restrict__ unit_lots, // [B * MAX_UNITS] + float* __restrict__ unit_initial_r, // [B * MAX_UNITS] + float* __restrict__ unit_trail_distance, // [B * MAX_UNITS] + unsigned char* __restrict__ unit_active, // [B * MAX_UNITS] + int* __restrict__ pyramid_units_count, // [B] + int b_size, + int pos_bytes, + int current_step +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const int curr_pos = *(const int*)(pos_state + b * pos_bytes + 0); + const float vwap = *(const float*)(pos_state + b * pos_bytes + 4); + const int prev_pos = prev_pos_lots[b]; + + const float k_init = isv[RL_TRAIL_K_INIT_INDEX]; + // Bootstrap trail uses fraction of price (see header comment). + const float bootstrap_d = fabsf(vwap) * BOOTSTRAP_TRAIL_FRACTION * k_init; + + const int base = b * MAX_UNITS; + + const bool was_flat = (prev_pos == 0); + const bool is_flat = (curr_pos == 0); + const bool reversed = (!was_flat && !is_flat && ((prev_pos > 0) != (curr_pos > 0))); + + if (is_flat && !was_flat) { + // CLOSE: deactivate all units, zero lots. + #pragma unroll + for (int u = 0; u < MAX_UNITS; ++u) { + unit_active[base + u] = 0; + unit_lots[base + u] = 0; + } + pyramid_units_count[b] = 0; + } else if (!is_flat && (was_flat || reversed)) { + // OPEN or REVERSE: clear all + populate slot 0. + #pragma unroll + for (int u = 0; u < MAX_UNITS; ++u) { + unit_active[base + u] = 0; + unit_lots[base + u] = 0; + } + unit_entry_price[base + 0] = vwap; + unit_entry_step[base + 0] = current_step; + unit_lots[base + 0] = curr_pos; + unit_initial_r[base + 0] = bootstrap_d; + unit_trail_distance[base + 0] = bootstrap_d; + unit_active[base + 0] = 1; + pyramid_units_count[b] = 1; + } else if (!is_flat && !was_flat && !reversed) { + // CONTINUE: same direction, update unit_lots[0] to current + // position (covers pyramid adds done by actions_to_market_targets + // even though SP20 P7 isn't shipped yet — unit_lots stays in + // sync with aggregate position). + unit_lots[base + 0] = curr_pos; + } + // (is_flat && was_flat): no-op, position remains flat. + + // Always advance prev_pos tracker. + prev_pos_lots[b] = curr_pos; +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 00a6942d9..1b598e387 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -46,8 +46,9 @@ //! | 491-492 | 2 controller-anchor slots (KL distill target + reward_scale MIN) | rl_isv_write (seed) + adaptive controllers | //! //! | 493 | Q→π distill KL_EMA blend alpha | rl_isv_write (seed) + rl_q_pi_distill_grad | +//! | 494-497 | SP20 P5 trail-stop bounds (MIN/MAX/K_INIT/ADJUST_RATE) | rl_isv_write (seed) + rl_trail_mutate + rl_unit_state_update | //! -//! Total: 94 slots; `RL_SLOTS_END = 494`. +//! Total: 98 slots; `RL_SLOTS_END = 498`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -733,6 +734,34 @@ pub const RL_REWARD_SCALE_MIN_INDEX: usize = 492; /// kernels is ISV-resident. Seeded 0.05 (preserves prior behavior). pub const RL_Q_DISTILL_KL_EMA_ALPHA_INDEX: usize = 493; +// ─── SP20 P5 trail-stop bounds (audit-wiring catch on TrailTighten/Loosen) ── +// +// Caught by `scripts/audit-wiring.sh` dogfood: actions a7/a8 had no +// consumer anywhere. SP20 P1+P5 bundle wires them up: per-batch +// per-unit trail_distance buffer (slot 0 used; slots 1-3 reserved +// for SP20 P7 pyramid expansion). + +/// Trail-stop minimum distance — `rl_trail_mutate` floors +/// `trail_distance × adjust_rate` here. Default 0.001 (1 mil ticks +/// scaled to price units — keeps trail from collapsing to zero). +pub const RL_TRAIL_MIN_INDEX: usize = 494; + +/// Trail-stop maximum distance — `rl_trail_mutate` ceilings +/// `trail_distance / adjust_rate` here. Default 100.0 (very wide, +/// caps tail explosion). +pub const RL_TRAIL_MAX_INDEX: usize = 495; + +/// Initial trail-stop distance multiplier — `rl_unit_state_update` +/// initializes new unit's trail to `k_init × sqrt(realized_return_var_ema)` +/// per `pearl_trade_level_vol_for_stop_distance`. Default 2.0 +/// (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. +pub const RL_TRAIL_ADJUST_RATE_INDEX: usize = 497; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 494; +pub const RL_SLOTS_END: usize = 498; diff --git a/crates/ml-alpha/src/rl/reward.rs b/crates/ml-alpha/src/rl/reward.rs index af9985089..880482e44 100644 --- a/crates/ml-alpha/src/rl/reward.rs +++ b/crates/ml-alpha/src/rl/reward.rs @@ -91,4 +91,14 @@ pub trait RlLobBackend { /// Number of parallel backtests (batch dimension). fn n_backtests(&self) -> usize; + + /// SP20 P5: read-only accessor for the shared best-bid price + /// array (`BOOK_LEVELS` floats, lobsim-wide — single book shared + /// across batches). Used by `rl_trail_stop_check` to compute + /// current mid for per-batch trail breach checks. + fn bid_px_d(&self) -> &CudaSlice; + + /// SP20 P5: read-only accessor for the shared best-ask price + /// array. Pair to `bid_px_d`. + fn ask_px_d(&self) -> &CudaSlice; } diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 28be1c903..8fffaa65b 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -184,6 +184,16 @@ const RL_Q_PI_DISTILL_GRAD_CUBIN: &[u8] = // Drives λ via Schulman bounded step on KL_EMA toward target. const RL_Q_DISTILL_LAMBDA_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_q_distill_lambda_controller.cubin")); +// SP20 P1+P5 — audit-wiring catch on TrailTighten/Loosen + foundational +// per-unit trade state machine. Three kernels work together: state +// machine on position transitions (open/close/reverse), trail mutation +// from a7/a8 actions, trail breach check that overrides action to flat. +const RL_UNIT_STATE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_unit_state_update.cubin")); +const RL_TRAIL_MUTATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_mutate.cubin")); +const RL_TRAIL_STOP_CHECK_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_stop_check.cubin")); const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/actions_to_market_targets.cubin")); @@ -411,6 +421,31 @@ pub struct IntegratedTrainer { // λ_distill adaptive controller (rljzl followup 2026-05-24). _rl_q_distill_lambda_controller_module: Arc, rl_q_distill_lambda_controller_fn: CudaFunction, + // SP20 P1+P5 (audit-wiring fix for dead a7/a8). + _rl_unit_state_update_module: Arc, + rl_unit_state_update_fn: CudaFunction, + _rl_trail_mutate_module: Arc, + rl_trail_mutate_fn: CudaFunction, + _rl_trail_stop_check_module: Arc, + rl_trail_stop_check_fn: CudaFunction, + + // Per-batch per-unit trade state (SP20 P1). + // MAX_UNITS=4 reserved; only slot 0 used until SP20 P7 ships + // pyramiding semantics. Diag exposes all 4 slots — slots 1-3 + // hold zeros pre-P7 (audit-diag tolerates per-batch-buffer + // arrays of any shape so long as values are non-NaN). + pub unit_entry_price_d: CudaSlice, // [B * 4] + pub unit_entry_step_d: CudaSlice, // [B * 4] + pub unit_lots_d: CudaSlice, // [B * 4] + pub unit_initial_r_d: CudaSlice, // [B * 4] + pub unit_trail_distance_d: CudaSlice, // [B * 4] + pub unit_active_d: CudaSlice, // [B * 4] + pub pyramid_units_count_d: CudaSlice, // [B] + pub unit_prev_pos_lots_d: CudaSlice, // [B] — separate from + // extract_realized_pnl_delta's + // prev_position_lots_d so unit + // state machine isn't coupled + // to reward extraction order. _actions_to_market_targets_module: Arc, actions_to_market_targets_fn: CudaFunction, @@ -820,6 +855,26 @@ impl IntegratedTrainer { let rl_q_distill_lambda_controller_fn = rl_q_distill_lambda_controller_module .load_function("rl_q_distill_lambda_controller") .context("load rl_q_distill_lambda_controller")?; + + // SP20 P1+P5 cubin loads. + let rl_unit_state_update_module = ctx + .load_cubin(RL_UNIT_STATE_UPDATE_CUBIN.to_vec()) + .context("load rl_unit_state_update cubin")?; + let rl_unit_state_update_fn = rl_unit_state_update_module + .load_function("rl_unit_state_update") + .context("load rl_unit_state_update")?; + let rl_trail_mutate_module = ctx + .load_cubin(RL_TRAIL_MUTATE_CUBIN.to_vec()) + .context("load rl_trail_mutate cubin")?; + let rl_trail_mutate_fn = rl_trail_mutate_module + .load_function("rl_trail_mutate") + .context("load rl_trail_mutate")?; + let rl_trail_stop_check_module = ctx + .load_cubin(RL_TRAIL_STOP_CHECK_CUBIN.to_vec()) + .context("load rl_trail_stop_check cubin")?; + let rl_trail_stop_check_fn = rl_trail_stop_check_module + .load_function("rl_trail_stop_check") + .context("load rl_trail_stop_check")?; let actions_to_market_targets_module = ctx .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) .context("load actions_to_market_targets cubin")?; @@ -994,6 +1049,32 @@ impl IntegratedTrainer { let prev_position_lots_d = stream .alloc_zeros::(b_size) .context("alloc prev_position_lots_d")?; + // SP20 P1+P5 per-unit trade state (slot 0 used; slots 1-3 + // reserved for P7 pyramid). + let unit_entry_price_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_entry_price_d")?; + let unit_entry_step_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_entry_step_d")?; + let unit_lots_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_lots_d")?; + let unit_initial_r_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_initial_r_d")?; + let unit_trail_distance_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_trail_distance_d")?; + let unit_active_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_active_d")?; + let pyramid_units_count_d = stream + .alloc_zeros::(b_size) + .context("alloc pyramid_units_count_d")?; + let unit_prev_pos_lots_d = stream + .alloc_zeros::(b_size) + .context("alloc unit_prev_pos_lots_d")?; let rewards_d = stream .alloc_zeros::(b_size) .context("alloc rewards_d")?; @@ -1134,6 +1215,20 @@ impl IntegratedTrainer { rl_q_pi_distill_grad_fn, _rl_q_distill_lambda_controller_module: rl_q_distill_lambda_controller_module, rl_q_distill_lambda_controller_fn, + _rl_unit_state_update_module: rl_unit_state_update_module, + rl_unit_state_update_fn, + _rl_trail_mutate_module: rl_trail_mutate_module, + rl_trail_mutate_fn, + _rl_trail_stop_check_module: rl_trail_stop_check_module, + rl_trail_stop_check_fn, + unit_entry_price_d, + unit_entry_step_d, + unit_lots_d, + unit_initial_r_d, + unit_trail_distance_d, + unit_active_d, + pyramid_units_count_d, + unit_prev_pos_lots_d, _actions_to_market_targets_module: actions_to_market_targets_module, actions_to_market_targets_fn, _abs_copy_module: abs_copy_module, @@ -1251,7 +1346,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 36] = [ + let isv_constants: [(usize, f32); 40] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1286,6 +1381,11 @@ impl IntegratedTrainer { // Q→π distill KL_EMA blend α (was hardcoded 0.05f in // rl_q_pi_distill_grad.cu, caught by audit-isv.sh dogfood). (crate::rl::isv_slots::RL_Q_DISTILL_KL_EMA_ALPHA_INDEX, 0.05), + // SP20 P1+P5 trail-stop bounds (audit-wiring fix for a7/a8 dead). + (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), (crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01), (crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99), (crate::rl::isv_slots::RL_PLATEAU_PATIENCE_INDEX, 1000.0), @@ -2981,6 +3081,59 @@ impl IntegratedTrainer { // rl_action_kernel above). let pos_bytes_i = lobsim.pos_bytes() as i32; let b_size_i = b_size as i32; + + // ── SP20 P1+P5 trail mutate (a7/a8) — BEFORE actions_to_market_targets. + // Mutates unit_trail_distance for active units if the chosen action + // is TrailTighten/Loosen. Non-trail actions pass through. + { + let cfg = LaunchConfig { + grid_dim: (b_size as u32, 1, 1), + block_dim: (4, 1, 1), // MAX_UNITS threads per block + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&self.rl_trail_mutate_fn); + launch + .arg(&self.actions_d) + .arg(&self.isv_d) + .arg(&self.unit_active_d) + .arg(&mut self.unit_trail_distance_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_trail_mutate launch")?; + } + } + + // ── SP20 P1+P5 trail stop check — AFTER trail_mutate, BEFORE + // actions_to_market_targets. May OVERRIDE actions_d[b] to + // FlatFromLong/Short on per-unit breach. Reads shared lobsim + // best book for current mid. + { + let bid_px_d = lobsim.bid_px_d(); + let ask_px_d = lobsim.ask_px_d(); + let cfg = LaunchConfig { + grid_dim: (b_size as u32, 1, 1), + block_dim: (4, 1, 1), // MAX_UNITS threads per block + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&self.rl_trail_stop_check_fn); + launch + .arg(&mut self.actions_d) + .arg(bid_px_d) + .arg(ask_px_d) + .arg(&self.unit_active_d) + .arg(&self.unit_entry_price_d) + .arg(&self.unit_lots_d) + .arg(&self.unit_trail_distance_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_trail_stop_check launch")?; + } + } + { let (pos_d_ref, market_targets_d) = lobsim.pos_and_market_targets_mut(); let cfg = LaunchConfig { @@ -3037,6 +3190,43 @@ impl IntegratedTrainer { } } + // ── SP20 P1 per-unit state machine update — runs AFTER fill + // + extract_realized_pnl_delta so position transitions are + // visible. Detects open/close/reverse and populates unit slot + // 0 with entry_price, entry_step, lots, initial_r, trail. + // Maintains its own prev_pos tracker (unit_prev_pos_lots_d) + // separate from extract_realized_pnl_delta's prev_position_lots_d + // so kernels are independently composable. + let current_step_i = (self.step_counter as i32).max(0); + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&self.rl_unit_state_update_fn); + launch + .arg(pos_d_ref) + .arg(&self.isv_d) + .arg(&mut self.unit_prev_pos_lots_d) + .arg(&mut self.unit_entry_price_d) + .arg(&mut self.unit_entry_step_d) + .arg(&mut self.unit_lots_d) + .arg(&mut self.unit_initial_r_d) + .arg(&mut self.unit_trail_distance_d) + .arg(&mut self.unit_active_d) + .arg(&mut self.pyramid_units_count_d) + .arg(&b_size_i) + .arg(&pos_bytes_i) + .arg(¤t_step_i); + unsafe { + launch + .launch(cfg) + .context("rl_unit_state_update launch")?; + } + } + // Trade-duration counter update + done-gated EMA emit. The // counter increments every step; on done it emits the count // (= trade duration in events) into `trade_duration_emit_d` diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index d3dfa4180..455a0f32b 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -1268,6 +1268,22 @@ impl LobSimCuda { &self.pos_d } + /// SP20 P1+P5 trail-stop integration: read-only accessor for the + /// shared best-bid price array (`BOOK_LEVELS` floats, lobsim-wide, + /// not per-batch). Used by `rl_trail_stop_check` to compute + /// current mid = (best_bid + best_ask) / 2 for per-batch + /// trail-distance breach checks. + pub fn bid_px_d(&self) -> &CudaSlice { + &self.bid_px_d + } + + /// SP20 P1+P5 trail-stop integration: read-only accessor for the + /// shared best-ask price array (`BOOK_LEVELS` floats). Pairs with + /// `bid_px_d()`. + pub fn ask_px_d(&self) -> &CudaSlice { + &self.ask_px_d + } + /// Phase R6: size of one `Pos` struct in bytes. Passed to device /// kernels that index into `pos_d` so they don't need to be rebuilt /// if `PosFlat` grows fields. @@ -1327,6 +1343,14 @@ impl ml_alpha::rl::reward::RlLobBackend for LobSimCuda { fn n_backtests(&self) -> usize { self.n_backtests() } + + fn bid_px_d(&self) -> &CudaSlice { + self.bid_px_d() + } + + fn ask_px_d(&self) -> &CudaSlice { + self.ask_px_d() + } } // Re-open the impl block for `LobSimCuda` so subsequent methods (if any diff --git a/scripts/audit-isv.sh b/scripts/audit-isv.sh index a2aefd7a3..4cf518de1 100755 --- a/scripts/audit-isv.sh +++ b/scripts/audit-isv.sh @@ -26,7 +26,12 @@ ALLOWED_NAMES=( "WARP_SIZE" "BLOCK_X" "BLOCK_Y" "BLOCK_Z" "MAX_WINDOW_ENTRIES" + "BOOK_LEVELS" ) +# Action enum indices (`ACTION_ = `) are structural API +# contracts matching crates/ml-alpha/src/rl/common.rs::Action enum +# positions. Allowed via name-prefix pattern. +ACTION_NAME_PREFIX="ACTION_" set -euo pipefail @@ -80,8 +85,8 @@ for kernel in "${kernels[@]}"; do name="${BASH_REMATCH[1]}" value="${BASH_REMATCH[2]}" total_defines=$((total_defines + 1)) - if [[ "$name" =~ $allowed_regex ]]; then - : # allowed + if [[ "$name" =~ $allowed_regex || "$name" == ${ACTION_NAME_PREFIX}* ]]; then + : # allowed (structural dim, ISV index, or ACTION_ enum mirror) else echo " VIOLATION $kernel: #define $name $value (should be ISV slot)" violations=$((violations + 1)) diff --git a/scripts/audit-manifest/actions.txt b/scripts/audit-manifest/actions.txt index 1b143841e..72585f2bb 100644 --- a/scripts/audit-manifest/actions.txt +++ b/scripts/audit-manifest/actions.txt @@ -1,3 +1,7 @@ # New Action enum entries added by the active development line. # Each phase commit appends here. # Format: one variant name per line (matches src/rl/common.rs::Action enum). +# SP20 P1+P5 audit fix: TrailTighten/TrailLoosen now have consumers +# (rl_trail_mutate + rl_trail_stop_check kernels). +TrailTighten +TrailLoosen diff --git a/scripts/audit-manifest/kernels.txt b/scripts/audit-manifest/kernels.txt index 42642d7df..4d44caa74 100644 --- a/scripts/audit-manifest/kernels.txt +++ b/scripts/audit-manifest/kernels.txt @@ -1,3 +1,6 @@ # New .cu kernels added by the active development line. # Each phase commit appends here. # Format: one basename per line, without `.cu` extension. +rl_unit_state_update +rl_trail_mutate +rl_trail_stop_check diff --git a/scripts/audit-manifest/slots.txt b/scripts/audit-manifest/slots.txt index 245fa6568..d9d49becd 100644 --- a/scripts/audit-manifest/slots.txt +++ b/scripts/audit-manifest/slots.txt @@ -1,3 +1,8 @@ # New ISV slot constants added by the active development line. # Each phase commit appends here. # Format: one RL_*_INDEX symbol per line. +RL_Q_DISTILL_KL_EMA_ALPHA_INDEX +RL_TRAIL_MIN_INDEX +RL_TRAIL_MAX_INDEX +RL_TRAIL_K_INIT_INDEX +RL_TRAIL_ADJUST_RATE_INDEX diff --git a/scripts/audit-wiring.sh b/scripts/audit-wiring.sh index 67756bf7e..2c93d6753 100755 --- a/scripts/audit-wiring.sh +++ b/scripts/audit-wiring.sh @@ -163,14 +163,21 @@ while IFS= read -r action; do } }' "$COMMON_RS") [[ -z "$action_idx" ]] && action_idx="-1" - # Handler in actions_to_market_targets.cu: branch `action == ` - if ! grep -qE "action[[:space:]]*==[[:space:]]*${action_idx}\b" "$ACTIONS_KERNEL"; then - # Could be handled in override kernel — check across all .cu files - handler_count=$(grep -lrE "(action|actions\[[^]]*\])[[:space:]]*==[[:space:]]*${action_idx}\b" "$CUDA_DIR" 2>/dev/null | wc -l) - if [[ "$handler_count" -eq 0 ]]; then - echo " VIOLATION action $action (idx $action_idx): no handler in actions_to_market_targets or override kernels" - violations=$((violations + 1)) - fi + # Handler patterns (any of): + # * Literal index branch in any .cu: `action == ` or `actions[..] == ` + # * Named-constant compare: `action == ACTION_` where + # UPPER_SNAKE is the CamelCase enum name → UPPER_SNAKE_CASE + # * Disequality (`!=`) is also a handler (a guard checking for the action) + upper_snake=$(echo "$action" | sed -E 's/([A-Z])/_\1/g; s/^_//' | tr '[:lower:]' '[:upper:]') + named_const="ACTION_${upper_snake}" + pat_literal="(action|actions\\[[^]]*\\])[[:space:]]*[!=]=[[:space:]]*${action_idx}\\b" + pat_named="(action|actions\\[[^]]*\\])[[:space:]]*[!=]=[[:space:]]*${named_const}\\b" + + handler_count=$(grep -lrE "(${pat_literal})|(${pat_named})" "$CUDA_DIR" 2>/dev/null | wc -l) + if [[ "$handler_count" -eq 0 ]]; then + echo " VIOLATION action $action (idx $action_idx): no handler in any kernel" + echo " (searched for: \`action == $action_idx\` or \`action == $named_const\`)" + violations=$((violations + 1)) fi done < <(read_manifest "$M/actions.txt")