fix(rl): wire TrailTighten/TrailLoosen + SP20 P1+P5 foundation

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 == <idx>` and `action == ACTION_<UPPER_SNAKE>` 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 <noreply@anthropic.com>

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

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 <stdint.h>
+
+#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;
+}

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 <stdint.h>
+
+#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;
+    }
+}

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 <stdint.h>
+
+#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;
+}

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;

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<f32>;
+
+    /// SP20 P5: read-only accessor for the shared best-ask price
+    /// array. Pair to `bid_px_d`.
+    fn ask_px_d(&self) -> &CudaSlice<f32>;
 }

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<CudaModule>,
     rl_q_distill_lambda_controller_fn: CudaFunction,
+    // SP20 P1+P5 (audit-wiring fix for dead a7/a8).
+    _rl_unit_state_update_module: Arc<CudaModule>,
+    rl_unit_state_update_fn: CudaFunction,
+    _rl_trail_mutate_module: Arc<CudaModule>,
+    rl_trail_mutate_fn: CudaFunction,
+    _rl_trail_stop_check_module: Arc<CudaModule>,
+    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<f32>,        // [B * 4]
+    pub unit_entry_step_d: CudaSlice<i32>,         // [B * 4]
+    pub unit_lots_d: CudaSlice<i32>,               // [B * 4]
+    pub unit_initial_r_d: CudaSlice<f32>,          // [B * 4]
+    pub unit_trail_distance_d: CudaSlice<f32>,     // [B * 4]
+    pub unit_active_d: CudaSlice<u8>,              // [B * 4]
+    pub pyramid_units_count_d: CudaSlice<i32>,     // [B]
+    pub unit_prev_pos_lots_d: CudaSlice<i32>,      // [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<CudaModule>,
     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::<i32>(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::<f32>(b_size * 4)
+            .context("alloc unit_entry_price_d")?;
+        let unit_entry_step_d = stream
+            .alloc_zeros::<i32>(b_size * 4)
+            .context("alloc unit_entry_step_d")?;
+        let unit_lots_d = stream
+            .alloc_zeros::<i32>(b_size * 4)
+            .context("alloc unit_lots_d")?;
+        let unit_initial_r_d = stream
+            .alloc_zeros::<f32>(b_size * 4)
+            .context("alloc unit_initial_r_d")?;
+        let unit_trail_distance_d = stream
+            .alloc_zeros::<f32>(b_size * 4)
+            .context("alloc unit_trail_distance_d")?;
+        let unit_active_d = stream
+            .alloc_zeros::<u8>(b_size * 4)
+            .context("alloc unit_active_d")?;
+        let pyramid_units_count_d = stream
+            .alloc_zeros::<i32>(b_size)
+            .context("alloc pyramid_units_count_d")?;
+        let unit_prev_pos_lots_d = stream
+            .alloc_zeros::<i32>(b_size)
+            .context("alloc unit_prev_pos_lots_d")?;
         let rewards_d = stream
             .alloc_zeros::<f32>(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<u8> = 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(&current_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`

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<f32> {
+        &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<f32> {
+        &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<f32> {
+        self.bid_px_d()
+    }
+
+    fn ask_px_d(&self) -> &CudaSlice<f32> {
+        self.ask_px_d()
+    }
 }
 
 // Re-open the impl block for `LobSimCuda` so subsequent methods (if any

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_<NAME> = <int>`) 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))

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

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

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

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 == <idx>`
-  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 == <idx>` or `actions[..] == <idx>`
+  #   * Named-constant compare: `action == ACTION_<UPPER_SNAKE>` 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")