feat(rl): R2 — fee model + loader pair API

Ports two scope-complete fixes from the ml-alpha-phase-f-g-flawed
reference branch:

A7 fee model:
- order_match.cu::submit_market_immediate gains 2 new kernel args
  (cost_per_lot_per_side, total_fees_per_b) mirroring the per-fill
  fee deduction in resting_orders.cu::apply_fill_to_pos:213-219.
  Fee deducts from pos.realized_pnl on EVERY fill (open, scale-in,
  counter); total_fees_per_b accumulates for telemetry.
  Single-writer-per-block plain += is safe — line 79 has
  `threadIdx.x != 0 → return` (feedback_no_atomicadd).
- LobSimCuda::submit_market launch updated to pass the 2 new args.
- LobSimCuda::upload_cost_per_lot_per_side host API lets callers
  configure ES-realistic fees (≈$1.25/contract/side). Default
  alloc_zeros = $0; production decision-policy path is unchanged
  (uploads its own cost via step_decision_with_latency).

A8 loader pair API:
- MultiHorizonLoader::next_sequence_pair returns
  (LabeledSequence, LabeledSequence) at adjacent anchors in the
  same source file. anchor_t sampled from [min_anchor, max_anchor−1)
  so anchor+1 also fits the upper-bound. Counts as ONE yielded
  sequence against n_max_sequences.
- next_sequence_random and next_sequence_pair share a new private
  helper build_sequence_at(lf, anchor) -> LabeledSequence that
  contains the multi-resolution windowing logic. Single source of
  truth for the build (feedback_single_source_of_truth_no_duplicates).
- next_sequence's caller-facing contract (random anchor, one
  sequence per call) is unchanged — alpha_train.rs supervised
  pipeline keeps working as-is.

No new local tests this phase per the rebuild plan (R2): the fee
deduction with default cost=0 is a no-op for existing callers, and
G7 in R6 covers the with-fees path via a rebuilt reward_calibration
test driving LobSimCuda directly (no LobEnv adapter).

cargo check -p ml-alpha -p ml-backtesting + cargo build --tests on
ml-backtesting both green; baseline test suite unaffected.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

crates/ml-alpha/src/data/loader.rs

@@ -679,6 +679,55 @@ impl MultiHorizonLoader {
         self.next_sequence_random()
     }
 
+    /// Phase R2 (was G fix): return a `(s_t, s_{t+1})` pair of
+    /// consecutive sequences at adjacent anchors in the same file.
+    /// The first sequence ends one snapshot earlier than the second —
+    /// the canonical RL Bellman-target requirement that
+    /// `IntegratedTrainer::step_with_lobsim` consumes as
+    /// `(snapshots, next_snapshots)`.
+    ///
+    /// Both yielded sequences come from the SAME random file so the
+    /// underlying market regime / time-of-day context is consistent.
+    /// The anchor for `s_t` is sampled uniformly from `[min_anchor,
+    /// max_anchor − 1)` so `anchor + 1` also satisfies the
+    /// upper-bound constraint. Counts as ONE yielded sequence against
+    /// `n_max_sequences` (the loader views the pair as one training
+    /// step's worth of input).
+    pub fn next_sequence_pair(
+        &mut self,
+    ) -> Result<Option<(LabeledSequence, LabeledSequence)>> {
+        if self.yielded >= self.cfg.n_max_sequences {
+            return Ok(None);
+        }
+        let n_files = self.files_loaded.len();
+        let file_idx = self.rng.gen_range(0..n_files);
+        let lf = &self.files_loaded[file_idx];
+
+        let max_horizon = *self.cfg.horizons.iter().max().expect("non-empty horizons");
+        let total = self.cfg.multi_resolution.total_positions();
+        let lookback = self.cfg.multi_resolution.required_lookback_ticks();
+        let min_anchor = lookback.saturating_sub(total);
+        // Reserve one extra event past the standard max_anchor so
+        // `anchor + 1`'s window also fits within the file.
+        let max_anchor = lf
+            .snapshots
+            .len()
+            .saturating_sub(total + max_horizon + 1);
+        anyhow::ensure!(
+            max_anchor > min_anchor,
+            "file too short for paired-window sampling (total={total}, lookback={lookback}, max_horizon={max_horizon}): \
+             snapshots.len()={} cannot span [{min_anchor}, {max_anchor})",
+            lf.snapshots.len()
+        );
+        let anchor_t: usize = self.rng.gen_range(min_anchor..max_anchor);
+
+        let s_t = self.build_sequence_at(lf, anchor_t)?;
+        let s_tp1 = self.build_sequence_at(lf, anchor_t + 1)?;
+
+        self.yielded += 1;
+        Ok(Some((s_t, s_tp1)))
+    }
+
     fn next_sequence_random(&mut self) -> Result<Option<LabeledSequence>> {
         if self.yielded >= self.cfg.n_max_sequences {
             return Ok(None);
@@ -712,6 +761,24 @@ impl MultiHorizonLoader {
         );
         let anchor: usize = self.rng.gen_range(min_anchor..max_anchor);
 
+        let seq = self.build_sequence_at(lf, anchor)?;
+        self.yielded += 1;
+        Ok(Some(seq))
+    }
+
+    /// Construct a `LabeledSequence` at a specific `(file, anchor)`
+    /// pair. Pure function of the loader config + file contents; does
+    /// not advance the loader's `yielded` counter (callers update it).
+    ///
+    /// Phase R2: this helper exists so `next_sequence_random` (random
+    /// anchor) and `next_sequence_pair` (paired adjacent anchors)
+    /// share the multi-resolution windowing logic verbatim per
+    /// `feedback_single_source_of_truth_no_duplicates`.
+    fn build_sequence_at(
+        &self,
+        lf: &LoadedFile,
+        anchor: usize,
+    ) -> Result<LabeledSequence> {
         let mut labels: [Vec<f32>; N_HORIZONS] = Default::default();
         let mut outcome_prof_long: [Vec<f32>; N_HORIZONS] = Default::default();
         let mut outcome_prof_short: [Vec<f32>; N_HORIZONS] = Default::default();
@@ -777,8 +844,7 @@ impl MultiHorizonLoader {
         let sequence: Vec<Mbp10RawInput> = sequence_rev.into_iter().rev().collect();
         debug_assert_eq!(sequence.len(), total);
 
-        self.yielded += 1;
-        Ok(Some(LabeledSequence {
+        Ok(LabeledSequence {
             snapshots: sequence,
             labels,
             outcome_prof_long,
@@ -787,7 +853,7 @@ impl MultiHorizonLoader {
             outcome_size_short,
             sigma_k,
             pos_fraction: lf.pos_fraction.clone(),
-        }))
+        })
     }
 }
 

crates/ml-backtesting/cuda/order_match.cu

@@ -67,12 +67,21 @@ __device__ static float walk_bid_for_sell(
 //
 // S1.22: min_reasonable_px / max_reasonable_px added so walk_* can reject
 // sentinel-priced depth levels (zero-priced or huge-priced empty MBP-10 slots).
+//
+// Phase R2 (was F.3a): cost_per_lot_per_side + total_fees_per_b mirror
+// the per-fill fee deduction in resting_orders.cu::apply_fill_to_pos
+// (lines 213-219). Fee is deducted from pos.realized_pnl on every fill
+// (open, scale-in, counter) and accumulated to total_fees_per_b for
+// telemetry. Default cost is zero (LobSimCuda::new), so this is a
+// no-op for callers that don't call upload_cost_per_lot_per_side.
 extern "C" __global__ void submit_market_immediate(
     const Book*         __restrict__ books,
     const int*          __restrict__ targets,         // [n_backtests, 2]
     Pos*                __restrict__ positions,        // [n_backtests]
     const float*        __restrict__ min_reasonable_px, // [n_backtests]
     const float*        __restrict__ max_reasonable_px, // [n_backtests]
+    const float*        __restrict__ cost_per_lot_per_side, // [n_backtests]
+    float*              __restrict__ total_fees_per_b,      // [n_backtests]
     int                              n_backtests
 ) {
     int b = blockIdx.x;
@@ -127,4 +136,13 @@ extern "C" __global__ void submit_market_immediate(
         // If position is now flat, vwap_entry becomes irrelevant.
     }
     pos.position_lots = (int)new_pos_f;
+
+    // Phase R2 (was F.3a): per-fill cost deduction
+    // (single-writer-per-block, plain += safe per feedback_no_atomicadd).
+    // Mirrors apply_fill_to_pos at resting_orders.cu:213-219 so the
+    // simple submit_market path matches the production resting-order
+    // path's economic model.
+    const float fill_cost = filled * cost_per_lot_per_side[b];
+    pos.realized_pnl -= fill_cost;
+    total_fees_per_b[b] += fill_cost;
 }

crates/ml-backtesting/src/sim/mod.rs

@@ -1189,6 +1189,8 @@ impl LobSimCuda {
                 .arg(&mut self.pos_d)
                 .arg(&self.min_reasonable_px_d)
                 .arg(&self.max_reasonable_px_d)
+                .arg(&self.cost_per_lot_per_side_d)
+                .arg(&mut self.total_fees_per_b_d)
                 .arg(&n)
                 .launch(cfg)?;
         }
@@ -1199,6 +1201,22 @@ impl LobSimCuda {
         Ok(())
     }
 
+    /// Upload the per-batch fee constant (USD per contract per side).
+    /// Default after `new()` is zero — callers should set this before
+    /// the first `submit_market` call if they want fees to apply. ES
+    /// futures default per AMP/CME pricing is ≈$1.25/contract/side
+    /// (~$2.50 round-trip). The production `step_decision_with_latency`
+    /// path uploads its own cost via `sim_cfg.cost_per_lot_per_side`.
+    pub fn upload_cost_per_lot_per_side(&mut self, cost: &[f32]) -> Result<()> {
+        anyhow::ensure!(
+            cost.len() == self.n_backtests,
+            "upload_cost_per_lot_per_side: len {} != n_backtests {}",
+            cost.len(), self.n_backtests
+        );
+        self.stream.memcpy_htod(cost, &mut self.cost_per_lot_per_side_d)?;
+        Ok(())
+    }
+
     /// Run pnl_track_step kernel: detect segment_complete, emit TradeRecord.
     /// Called automatically by submit_market; exposed for callers that
     /// orchestrate matching themselves.