diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index a56a7281a..9739d147e 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -112,6 +112,140 @@ __device__ __forceinline__ int argmax_n(const __nv_bfloat16* arr, int n) { /* exposure_idx_to_fraction DELETED — replaced by compute_target_position() * from trade_physics.cuh which uses dynamic step = 2/(b0_size-1). */ +/* ================================================================== */ +/* Kernel 0: saboteur_generate_params (#33 GPU-native) */ +/* ================================================================== */ + +/** + * Generate per-episode adversarial market parameters using GPU-resident LCG RNG. + * Each episode gets independent spread/fill/slippage multipliers centered + * around the current best params with random perturbation. + * + * Grid: ceil(N / 256), Block: 256. One thread per episode. + * + * @param saboteur_params [N, 3] output: (spread_mult, fill_prob, slippage_mult) + * @param rng_states [N] LCG RNG state (read-write) + * @param base_params [3] current best params (center of perturbation) + * @param perturbation_scale controls exploration width (decays over epochs) + * @param N number of episodes + */ +extern "C" __global__ void saboteur_generate_params( + float* __restrict__ saboteur_params, /* [N, 3] output */ + unsigned int* __restrict__ rng_states, /* [N] read-write */ + const float* __restrict__ base_params, /* [3] center */ + float perturbation_scale, + int N +) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; + + unsigned int rng = rng_states[i]; + + float base_spread = base_params[0]; + float base_fill = base_params[1]; + float base_slip = base_params[2]; + float ps = perturbation_scale; + + /* Box-Muller-lite for Gaussian-like perturbation */ + rng = rng * 1664525u + 1013904223u; + float u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; + rng = rng * 1664525u + 1013904223u; + float u2 = (float)(rng & 0xFFFF) / 65536.0f; + float g1 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); + + rng = rng * 1664525u + 1013904223u; + u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; + rng = rng * 1664525u + 1013904223u; + u2 = (float)(rng & 0xFFFF) / 65536.0f; + float g2 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); + + rng = rng * 1664525u + 1013904223u; + u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; + rng = rng * 1664525u + 1013904223u; + u2 = (float)(rng & 0xFFFF) / 65536.0f; + float g3 = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); + + /* Perturb from base with Gaussian noise, clamp to valid ranges */ + float spread = fminf(fmaxf(base_spread + g1 * ps * 2.0f, 0.5f), 5.0f); + float fill = fminf(fmaxf(base_fill + g2 * ps * 0.3f, 0.3f), 0.95f); + float slip = fminf(fmaxf(base_slip + g3 * ps * 1.5f, 0.5f), 3.0f); + + saboteur_params[i * 3 + 0] = spread; + saboteur_params[i * 3 + 1] = fill; + saboteur_params[i * 3 + 2] = slip; + + rng_states[i] = rng; +} + +/** + * Select the saboteur params that caused the worst trader performance. + * Reads per-episode raw_returns (from experience collection) and finds + * the episode with the lowest cumulative return. Its params become the + * next epoch's baseline. + * + * Grid: (1, 1, 1), Block: (256, 1, 1). Single-block reduction. + * + * @param saboteur_params [N, 3] per-episode params (from generate) + * @param raw_returns [N, L] per-bar returns (from experience collection) + * @param best_output [3] output: winning (worst for trader) params + * @param best_return_out [1] output: the worst cumulative return + * @param N number of episodes + * @param L timesteps per episode + */ +extern "C" __global__ void saboteur_select_best( + const float* __restrict__ saboteur_params, /* [N, 3] */ + const __nv_bfloat16* __restrict__ raw_returns, /* [N, L] */ + float* __restrict__ best_output, /* [3] */ + float* __restrict__ best_return_out, /* [1] */ + int N, + int L +) { + __shared__ float s_worst_return[256]; + __shared__ int s_worst_idx[256]; + + int tid = threadIdx.x; + + /* Each thread finds worst episode in its tile */ + float my_worst = 1e9f; + int my_idx = 0; + + for (int i = tid; i < N; i += 256) { + /* Sum per-bar returns for this episode */ + float cumul = 0.0f; + const __nv_bfloat16* ep = raw_returns + (long long)i * L; + for (int t = 0; t < L; t++) { + cumul += (float)ep[t]; + } + if (cumul < my_worst) { + my_worst = cumul; + my_idx = i; + } + } + + s_worst_return[tid] = my_worst; + s_worst_idx[tid] = my_idx; + __syncthreads(); + + /* Block reduction: find global worst */ + for (int s = 128; s > 0; s >>= 1) { + if (tid < s) { + if (s_worst_return[tid + s] < s_worst_return[tid]) { + s_worst_return[tid] = s_worst_return[tid + s]; + s_worst_idx[tid] = s_worst_idx[tid + s]; + } + } + __syncthreads(); + } + + if (tid == 0) { + int winner = s_worst_idx[0]; + best_output[0] = saboteur_params[winner * 3 + 0]; + best_output[1] = saboteur_params[winner * 3 + 1]; + best_output[2] = saboteur_params[winner * 3 + 2]; + best_return_out[0] = s_worst_return[0]; + } +} + /* ================================================================== */ /* Kernel 1: experience_state_gather */ /* ================================================================== */ @@ -619,11 +753,28 @@ extern "C" __global__ void experience_env_step( * at episode end (done=1) and added to final reward. */ float* __restrict__ position_histogram, float position_entropy_weight, /* #19: reward += weight * H(histogram). 0=disabled. */ - float regret_blend /* #17: blend factor for counterfactual regret (0=pure PnL, 1=pure regret) */ + float regret_blend, /* #17: blend factor for counterfactual regret (0=pure PnL, 1=pure regret) */ + /* #33 Per-episode saboteur params [N, 3]: (spread_mult, fill_prob, slippage_mult). + * NULL = disabled (use global scalars). When non-NULL, overrides + * spread_cost, fill_ioc_fill_prob, tx_cost_multiplier per episode. */ + const float* __restrict__ saboteur_params ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; + /* #33 Per-episode saboteur overrides */ + if (saboteur_params != NULL) { + float sab_spread = saboteur_params[i * 3 + 0]; + float sab_fill = saboteur_params[i * 3 + 1]; + float sab_slip = saboteur_params[i * 3 + 2]; + spread_cost *= sab_spread; + tx_cost_multiplier *= sab_slip; + /* fill_ioc_fill_prob is not a kernel param — it's baked into the + * portfolio sim via trade_physics.cuh. We can't override it per-episode + * without adding it as a kernel parameter. For now, the saboteur affects + * spread and slippage (the two most impactful microstructure levers). */ + } + int t = current_timesteps[i]; int bar_idx = episode_starts[i] + t; diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 5205ebdd5..001660450 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -542,6 +542,23 @@ pub struct GpuExperienceCollector { exp_bn_hidden: Option>, /// #31 Bottleneck concat buffer [alloc_episodes, bn_dim + portfolio_dim] bf16. exp_bn_concat: Option>, + /// #33 GPU-native saboteur: per-episode params [alloc_episodes, 3] f32. + saboteur_params_buf: CudaSlice, + /// #33 Saboteur base params [3] f32 (center of perturbation). + saboteur_base_buf: CudaSlice, + /// #33 Saboteur best output [3] f32 (result of selection). + saboteur_best_buf: CudaSlice, + /// #33 Saboteur best return [1] f32. + saboteur_best_return_buf: CudaSlice, + /// #33 Saboteur generate kernel. + saboteur_generate_kernel: CudaFunction, + /// #33 Saboteur select kernel. + saboteur_select_kernel: CudaFunction, + /// #33 Whether saboteur is active this epoch. + saboteur_active: bool, + /// #33 Perturbation scale (decays over epochs). + saboteur_perturbation_scale: f32, + /// #31 Bottleneck dimension (0 = disabled). bottleneck_dim: usize, /// #31 Market feature dimension for bottleneck separation. @@ -915,6 +932,34 @@ impl GpuExperienceCollector { let position_histogram = stream.alloc_zeros::(alloc_episodes * 9) .map_err(|e| MLError::ModelError(format!("alloc position_histogram: {e}")))?; + // #33 Saboteur GPU buffers + let saboteur_params_buf = stream.alloc_zeros::(alloc_episodes * 3) + .map_err(|e| MLError::ModelError(format!("alloc saboteur_params: {e}")))?; + let mut saboteur_base_buf = stream.alloc_zeros::(3) + .map_err(|e| MLError::ModelError(format!("alloc saboteur_base: {e}")))?; + // Init base params to neutral (no adversarial effect) + stream.memcpy_htod(&[1.0_f32, 0.85, 1.0], &mut saboteur_base_buf) + .map_err(|e| MLError::ModelError(format!("init saboteur_base: {e}")))?; + let saboteur_best_buf = stream.alloc_zeros::(3) + .map_err(|e| MLError::ModelError(format!("alloc saboteur_best: {e}")))?; + let saboteur_best_return_buf = stream.alloc_zeros::(1) + .map_err(|e| MLError::ModelError(format!("alloc saboteur_best_return: {e}")))?; + // Load saboteur kernels from experience cubin + let saboteur_generate_kernel = { + let context = stream.context(); + let module = context.load_cubin(EXPERIENCE_KERNELS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("saboteur cubin: {e}")))?; + module.load_function("saboteur_generate_params") + .map_err(|e| MLError::ModelError(format!("saboteur_generate load: {e}")))? + }; + let saboteur_select_kernel = { + let context = stream.context(); + let module = context.load_cubin(EXPERIENCE_KERNELS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("saboteur cubin2: {e}")))?; + module.load_function("saboteur_select_best") + .map_err(|e| MLError::ModelError(format!("saboteur_select load: {e}")))? + }; + // #31 Load bottleneck tanh+concat kernel from utility cubin (if active) let bn_tanh_concat_fn = if bn_dim_from_params > 0 { use super::gpu_dqn_trainer::DQN_UTILITY_CUBIN; @@ -1005,6 +1050,14 @@ impl GpuExperienceCollector { curiosity_trainer, feature_mask_buf: None, position_histogram, + saboteur_params_buf, + saboteur_base_buf, + saboteur_best_buf, + saboteur_best_return_buf, + saboteur_generate_kernel, + saboteur_select_kernel, + saboteur_active: false, + saboteur_perturbation_scale: 0.3, exp_bn_hidden, exp_bn_concat, bottleneck_dim: bn_dim_from_params, @@ -1028,6 +1081,18 @@ impl GpuExperienceCollector { self.cvar_scales_ptr = device_ptr; } + /// #33 Enable/disable the GPU-native saboteur for this epoch. + /// When active, per-episode adversarial parameters are generated by CUDA kernel + /// and the best (worst for trader) params are selected after collection. + pub fn set_saboteur_active(&mut self, active: bool) { + self.saboteur_active = active; + } + + /// #33 Set perturbation scale for the saboteur's evolutionary search. + pub fn set_saboteur_perturbation_scale(&mut self, scale: f32) { + self.saboteur_perturbation_scale = scale; + } + /// Upload pre-computed expert demonstration actions to GPU. /// /// `expert_actions[bar_index]` = expert exposure action index (-1 = no opinion). @@ -1303,6 +1368,24 @@ impl GpuExperienceCollector { let mirror_i32: i32 = if config.mirror_active { 1 } else { 0 }; + // #33 Saboteur: generate per-episode adversarial params on GPU + if self.saboteur_active { + let ps = self.saboteur_perturbation_scale; + unsafe { + self.stream + .launch_builder(&self.saboteur_generate_kernel) + .arg(&mut self.saboteur_params_buf) + .arg(&mut self.rng_states) + .arg(&self.saboteur_base_buf) + .arg(&ps) + .arg(&n_i32) + .launch(launch_cfg) + .map_err(|e| MLError::ModelError(format!( + "saboteur_generate_params: {e}" + )))?; + } + } + // #19 Position entropy: zero histogram at epoch start if config.position_entropy_weight > 0.0 { self.stream.memset_zeros(&mut self.position_histogram) @@ -1566,6 +1649,14 @@ impl GpuExperienceCollector { .arg(&mut self.position_histogram) // #19 position entropy histogram .arg(&config.position_entropy_weight) // #19 position entropy weight .arg(&config.regret_blend) // #17 counterfactual regret + // #33 Per-episode saboteur params (0 = NULL = disabled) + .arg(&{ + if self.saboteur_active { + self.saboteur_params_buf.device_ptr(&self.stream).0 + } else { + 0u64 + } + }) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" @@ -1573,6 +1664,42 @@ impl GpuExperienceCollector { } } + // #33 Saboteur: select best (worst for trader) params from this epoch + if self.saboteur_active { + let l_i32 = timesteps as i32; + unsafe { + self.stream + .launch_builder(&self.saboteur_select_kernel) + .arg(&self.saboteur_params_buf) + .arg(&self.raw_returns_out) + .arg(&mut self.saboteur_best_buf) + .arg(&mut self.saboteur_best_return_buf) + .arg(&n_i32) + .arg(&l_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "saboteur_select_best: {e}" + )))?; + } + // Copy best params → base for next epoch's perturbation center (GPU DtoD) + let best_ptr = self.saboteur_best_buf.raw_ptr(); + let base_ptr = self.saboteur_base_buf.raw_ptr(); + unsafe { + cudarc::driver::sys::cuMemcpyDtoDAsync_v2( + base_ptr, best_ptr, + 3 * std::mem::size_of::(), + self.stream.cu_stream(), + ); + } + // Decay perturbation scale + self.saboteur_perturbation_scale *= 0.995; + self.saboteur_perturbation_scale = self.saboteur_perturbation_scale.max(0.05); + } + // Auto-clear reset flags after launch self.reset_flags = 0; diff --git a/crates/ml/src/trainers/dqn/adversarial_self_play.rs b/crates/ml/src/trainers/dqn/adversarial_self_play.rs index 8adcecfbf..3d9e16549 100644 --- a/crates/ml/src/trainers/dqn/adversarial_self_play.rs +++ b/crates/ml/src/trainers/dqn/adversarial_self_play.rs @@ -1,53 +1,34 @@ //! #33 Adversarial Self-Play with Past Self (THE KING) //! -//! Trains a saboteur network that controls market microstructure (spread, -//! fill probability, slippage) to MAXIMIZE the trader's losses. The saboteur -//! is initialized from a past checkpoint of the trader's weights — it knows -//! exactly what the trader USED to do and specifically attacks those patterns. +//! GPU-native evolutionary saboteur that controls market microstructure +//! (spread, fill probability, slippage) to MAXIMIZE the trader's losses. //! -//! The trader must evolve AWAY from its own past strategies to survive. -//! The only stable equilibrium is a policy that works against ALL possible -//! adversaries — including itself. +//! All random number generation and parameter search happens on GPU: +//! - `saboteur_generate_params` kernel: per-episode params from GPU LCG RNG +//! - `saboteur_select_best` kernel: reduction to find worst-for-trader params +//! - Per-episode overrides in `experience_env_step` kernel via `saboteur_params` ptr //! -//! Architecture: -//! Saboteur: [mean_reward_last_epoch, mean_sharpe, epoch_fraction] → [32] ReLU → [3] sigmoid -//! Outputs: (spread_mult ∈ [0.5, 5.0], fill_prob ∈ [0.3, 0.95], slippage_mult ∈ [0.5, 3.0]) +//! This Rust struct only tracks epoch-level state (warmup, cycle phase). +//! Zero CPU-side RNG. Zero host↔device transfers for the evolutionary search. //! //! Self-play cycle: //! Phase 0 (epochs 0..warmup): Normal training, no saboteur -//! Phase 1 (odd epochs after warmup): Saboteur trains (maximize trader loss) -//! Phase 2 (even epochs after warmup): Trader trains against frozen saboteur +//! Phase 1 (odd epochs after warmup): Saboteur explores (GPU kernel generates params) +//! Phase 2 (even epochs after warmup): Trader trains against best saboteur params use tracing::info; -/// Saboteur state: controls market microstructure adversarially. +/// Epoch-level saboteur state. All computation is GPU-native. pub(crate) struct AdversarialSaboteur { /// Whether self-play is enabled pub(crate) enabled: bool, - /// Warmup epochs before self-play begins (let trader learn basics first) + /// Warmup epochs before self-play begins warmup_epochs: usize, - /// How often to save a new trader checkpoint as the saboteur's reference + /// How often to reset the saboteur's search (epochs) checkpoint_interval: usize, - - /// Saboteur's current output: spread multiplier [0.5, 5.0] - pub(crate) spread_mult: f32, - /// Saboteur's current output: fill probability [0.3, 0.95] - pub(crate) fill_prob: f32, - /// Saboteur's current output: slippage multiplier [0.5, 3.0] - pub(crate) slippage_mult: f32, - - /// Simple exponential moving average of trader's epoch reward. - /// The saboteur tries to MINIMIZE this (maximize losses). - trader_reward_ema: f32, - /// Saboteur's learning rate for gradient-free optimization (CMA-ES lite) - saboteur_lr: f32, - /// Best saboteur params found so far (spread, fill, slippage) - best_params: [f32; 3], - /// Best (worst for trader) reward achieved - best_reward: f32, - /// Random perturbation scale for exploration - perturbation_scale: f32, - /// Step counter for perturbation scheduling + /// Current perturbation scale (decays over epochs) + pub(crate) perturbation_scale: f32, + /// Epoch counter for decay scheduling step_count: usize, } @@ -57,13 +38,6 @@ impl AdversarialSaboteur { enabled, warmup_epochs, checkpoint_interval, - spread_mult: 1.0, - fill_prob: 0.85, - slippage_mult: 1.0, - trader_reward_ema: 0.0, - saboteur_lr: 0.1, - best_params: [1.0, 0.85, 1.0], - best_reward: f32::MAX, perturbation_scale: 0.3, step_count: 0, } @@ -74,81 +48,43 @@ impl AdversarialSaboteur { self.enabled && epoch >= self.warmup_epochs } - /// Check if the saboteur is training this epoch (odd epochs after warmup). - pub(crate) fn is_saboteur_training(&self, epoch: usize) -> bool { + /// Check if the saboteur is exploring this epoch (odd epochs after warmup). + pub(crate) fn is_saboteur_epoch(&self, epoch: usize) -> bool { self.is_active(epoch) && (epoch - self.warmup_epochs) % 2 == 1 } - /// Update the saboteur based on the trader's epoch performance. - /// - /// Uses gradient-free optimization (evolutionary strategy): - /// - Perturb current params → evaluate trader → keep if trader did worse - /// - This is CMA-ES without the covariance matrix (1D perturbations) - /// - /// Called at the end of each epoch with the trader's Sharpe ratio. - pub(crate) fn update(&mut self, epoch: usize, trader_sharpe: f32) { + /// Called at epoch end. Updates perturbation scale decay. + /// The actual evolutionary search (param generation + selection) is done + /// entirely in CUDA kernels inside the experience collector. + pub(crate) fn update(&mut self, epoch: usize, _trader_sharpe: f32) { if !self.is_active(epoch) { return; } - // EMA of trader reward - let alpha = 0.3_f32; - self.trader_reward_ema = alpha * trader_sharpe + (1.0 - alpha) * self.trader_reward_ema; - self.step_count += 1; - if self.is_saboteur_training(epoch) { - // Saboteur's goal: MINIMIZE trader_sharpe (make it as negative as possible) - // If trader did worse than best (lower Sharpe), keep these params - if trader_sharpe < self.best_reward { - self.best_reward = trader_sharpe; - self.best_params = [self.spread_mult, self.fill_prob, self.slippage_mult]; - info!( - epoch, trader_sharpe, - spread = self.spread_mult, fill = self.fill_prob, slippage = self.slippage_mult, - "SABOTEUR found better attack params" - ); - } + // Decay perturbation scale (exploration → exploitation) + self.perturbation_scale *= 0.995; + self.perturbation_scale = self.perturbation_scale.max(0.05); - // Perturb for next evaluation - use rand::Rng; - let mut rng = rand::thread_rng(); - let decay = 1.0 / (1.0 + self.step_count as f32 * 0.01); // Slow decay - let ps = self.perturbation_scale * decay; - - self.spread_mult = (self.best_params[0] + rng.gen_range(-ps..ps) * 2.0) - .clamp(0.5, 5.0); - self.fill_prob = (self.best_params[1] + rng.gen_range(-ps..ps) * 0.3) - .clamp(0.3, 0.95); - self.slippage_mult = (self.best_params[2] + rng.gen_range(-ps..ps) * 1.5) - .clamp(0.5, 3.0); + // Reset search every checkpoint_interval epochs + if self.checkpoint_interval > 0 + && self.step_count % self.checkpoint_interval == 0 + { + self.perturbation_scale = 0.3; // Reset exploration + info!( + epoch, step_count = self.step_count, + "SABOTEUR search reset (checkpoint interval)" + ); + } + if epoch % 10 == 0 { info!( epoch, - next_spread = self.spread_mult, - next_fill = self.fill_prob, - next_slippage = self.slippage_mult, - "SABOTEUR next perturbation" + perturbation_scale = %format!("{:.3}", self.perturbation_scale), + is_exploring = self.is_saboteur_epoch(epoch), + "Saboteur epoch state (GPU-native)" ); - } else { - // Trader epoch: apply best saboteur params (frozen adversary) - self.spread_mult = self.best_params[0]; - self.fill_prob = self.best_params[1]; - self.slippage_mult = self.best_params[2]; } } - - /// Apply saboteur outputs to the experience collector config. - /// Called during ExperienceCollectorConfig construction. - pub(crate) fn apply_to_spread(&self, base_spread: f32) -> f32 { - base_spread * self.spread_mult - } - - pub(crate) fn apply_to_fill_prob(&self, base_fill: f32) -> f32 { - (base_fill * self.fill_prob / 0.85).clamp(0.1, 0.99) - } - - pub(crate) fn apply_to_tx_cost(&self, base_cost: f32) -> f32 { - base_cost * self.slippage_mult - } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index fbeaf9233..a6aeaa460 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -947,6 +947,17 @@ impl DQNTrainer { } } + // #33 GPU-native saboteur: set active state BEFORE borrowing collector + { + let active = self.saboteur.is_active(self.current_epoch); + if let Some(ref mut c) = self.gpu_experience_collector { + c.set_saboteur_active(active); + if active { + c.set_saboteur_perturbation_scale(self.saboteur.perturbation_scale); + } + } + } + let ( Some(ref mut collector), Some(ref features_buf), @@ -1098,7 +1109,7 @@ impl DQNTrainer { info!("Adversarial regime ACTIVE this epoch: 3x spread, 2x tx_cost, 0.5x fill"); } - let mut config = ExperienceCollectorConfig { + let config = ExperienceCollectorConfig { n_episodes, timesteps_per_episode: timesteps, total_bars, @@ -1202,21 +1213,7 @@ impl DQNTrainer { ..Default::default() }; - // #33 Adversarial self-play: apply saboteur multipliers on top of everything - if self.saboteur.is_active(self.current_epoch) { - config.spread_cost = self.saboteur.apply_to_spread(config.spread_cost); - config.fill_ioc_fill_prob = self.saboteur.apply_to_fill_prob(config.fill_ioc_fill_prob); - config.tx_cost_multiplier = self.saboteur.apply_to_tx_cost(config.tx_cost_multiplier); - if self.current_epoch % 10 == 0 { - tracing::info!( - epoch = self.current_epoch + 1, - spread_mult = self.saboteur.spread_mult, - fill_prob = self.saboteur.fill_prob, - slippage_mult = self.saboteur.slippage_mult, - "Saboteur active: modified experience collection params" - ); - } - } + // #33 Saboteur state already set above (before collector borrow) // Zero-roundtrip GPU path — GPU PER is always active in CUDA builds let gpu_batch = collector.collect_experiences_gpu(