feat(generalization): #33 GPU-native saboteur — zero CPU-side RNG

Replace CPU rand::thread_rng() saboteur with fully GPU-native implementation:

Two new CUDA kernels in experience_kernels.cu:
- saboteur_generate_params: per-episode adversarial parameters [N, 3]
  via LCG GPU RNG with Box-Muller Gaussian perturbation. Each of N
  episodes gets independent (spread_mult, fill_prob, slippage_mult).
- saboteur_select_best: single-block reduction finds the episode
  whose params caused the WORST trader performance (lowest cumulative
  return). Winner's params become next epoch's perturbation center.

experience_env_step modified: reads per-episode saboteur_params[i, 3]
pointer. When non-NULL, overrides spread_cost and tx_cost_multiplier
per episode. NULL = disabled (standard global scalars).

GPU data flow (zero CPU involvement):
  generate_params (GPU LCG) → env_step reads per-episode →
  select_best (GPU reduction) → DtoD copy to base_params →
  next epoch generate_params centered on winner

Rust AdversarialSaboteur simplified to epoch-level state tracker.
All randomness, evaluation, and selection on GPU.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-30 10:43:22 +02:00
parent 152f4c16d9
commit 8bf1cdca13
4 changed files with 331 additions and 120 deletions

View File

@@ -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;

View File

@@ -542,6 +542,23 @@ pub struct GpuExperienceCollector {
exp_bn_hidden: Option<CudaSlice<half::bf16>>,
/// #31 Bottleneck concat buffer [alloc_episodes, bn_dim + portfolio_dim] bf16.
exp_bn_concat: Option<CudaSlice<half::bf16>>,
/// #33 GPU-native saboteur: per-episode params [alloc_episodes, 3] f32.
saboteur_params_buf: CudaSlice<f32>,
/// #33 Saboteur base params [3] f32 (center of perturbation).
saboteur_base_buf: CudaSlice<f32>,
/// #33 Saboteur best output [3] f32 (result of selection).
saboteur_best_buf: CudaSlice<f32>,
/// #33 Saboteur best return [1] f32.
saboteur_best_return_buf: CudaSlice<f32>,
/// #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::<f32>(alloc_episodes * 9)
.map_err(|e| MLError::ModelError(format!("alloc position_histogram: {e}")))?;
// #33 Saboteur GPU buffers
let saboteur_params_buf = stream.alloc_zeros::<f32>(alloc_episodes * 3)
.map_err(|e| MLError::ModelError(format!("alloc saboteur_params: {e}")))?;
let mut saboteur_base_buf = stream.alloc_zeros::<f32>(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::<f32>(3)
.map_err(|e| MLError::ModelError(format!("alloc saboteur_best: {e}")))?;
let saboteur_best_return_buf = stream.alloc_zeros::<f32>(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::<f32>(),
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;

View File

@@ -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
}
}

View File

@@ -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(