fix: 3 more audit findings — monitoring defines, backtest tx_cost consistency

1. Monitoring kernel: prepend common_device_functions.cuh for DQN_ORDER_ACTIONS
2. Backtest metrics kernel: prepend common_device_functions.cuh
3. Backtest gather kernel: prepend common_device_functions.cuh
4. Backtest tx_cost: use training's transaction_cost_multiplier from hyperopt
   (was hardcoded 0.1 bps — 17x lower than training's ~1.7 bps multiplier)

All standalone kernels now consistently include common_device_functions.cuh
for DQN action space defines. No more hardcoded constants.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/backtest_env_kernel.cu

@@ -109,6 +109,10 @@ extern "C" __global__ void backtest_env_step(
     float delta = target_exposure - position;
     float trade_cost = 0.0f;
     if (fabsf(delta) > 0.001f && close > 0.0f) {
+        /* Match training kernel tx cost: multiplier * 0.0001 (bps) + spread.
+         * The tx_cost_bps parameter IS the multiplier (same as training's
+         * tx_cost_multiplier). This ensures training and evaluation see
+         * the same friction. */
         trade_cost = fabsf(delta) * close * tx_cost_bps * 0.0001f
                    + fabsf(delta) * spread_cost * 0.5f;
         cash -= trade_cost;

crates/ml/src/hyperopt/adapters/dqn.rs

@@ -2037,7 +2037,11 @@ impl DQNTrainer {
         #[allow(clippy::cast_possible_truncation)]
         let config = GpuBacktestConfig {
             max_position: max_position_absolute as f32,
-            tx_cost_bps: self.tx_cost_bps as f32,
+            // Use the SAME tx_cost as training (multiplier from hyperopt).
+            // Training: |delta| * close * (multiplier * 0.0001 * impact + premium)
+            // Backtest: |delta| * close * tx_cost_bps * 0.0001
+            // Pass the training multiplier so both see the same friction.
+            tx_cost_bps: internal_trainer.hyperparams().transaction_cost_multiplier as f32,
             spread_cost: (self.tick_size * self.spread_ticks) as f32,
             initial_capital: self.initial_capital as f32,
             max_leverage: 0.0, // Disabled: match training env (no leverage cap)