diff --git a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
index d174a31e7..552868255 100644
--- a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
@@ -12,12 +12,19 @@ use crate::MLError;
 static MONITORING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/monitoring_kernel.cubin"));
 
 /// Compact monitoring summary from GPU reduction (48-byte GPU transfer, 24 × bf16).
+///
+/// Reward statistics (mean, std, min, max, sharpe) are computed **per trade**
+/// (non-zero rewards only), not per bar. With sparse trade-completion rewards
+/// (98% of bars have reward=0), per-bar mean is always ~0 and uninformative.
 #[derive(Debug, Clone, Default)]
 pub struct MonitoringSummary {
+    /// Mean reward per trade completion (non-zero rewards only).
     pub mean_reward: f32,
+    /// Std dev of trade rewards.
     pub reward_std: f32,
     pub min_reward: f32,
     pub max_reward: f32,
+    /// Sharpe estimate: mean_reward / reward_std (per-trade).
     pub sharpe_estimate: f32,
     /// Per-exposure action counts (9 levels: S100..L100).
     pub action_counts: [usize; 9],
@@ -25,7 +32,10 @@ pub struct MonitoringSummary {
     pub order_counts: [usize; 3],
     /// Per-urgency counts (3: Patient, Normal, Aggressive).
     pub urgency_counts: [usize; 3],
+    /// Total bars (experiences) in the epoch.
     pub total_experiences: usize,
+    /// Number of trade completions (non-zero rewards). mean_reward = sum / total_trades.
+    pub total_trades: usize,
 }
 
 /// GPU monitoring reducer.
@@ -108,6 +118,7 @@ impl GpuMonitoringReducer {
             order_counts: [raw[14] as usize, raw[15] as usize, raw[16] as usize],
             urgency_counts: [raw[17] as usize, raw[18] as usize, raw[19] as usize],
             total_experiences: raw[20] as usize,
+            total_trades: raw[21] as usize,
         })
     }
 }
diff --git a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
index 142bd79ef..35750957e 100644
--- a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
@@ -47,9 +47,12 @@ extern "C" __global__ void monitoring_reduce(
     int stride = blockDim.x;
     int b1b2 = order_actions * urgency_actions;
 
+    __shared__ int s_nonzero;
+
     if (tid == 0) {
         s_sum = bf16_zero(); s_sq_sum = bf16_zero();
         s_min = bf16(1e30f); s_max = bf16(-1e30f);
+        s_nonzero = 0;
         for (int i = 0; i < 9; i++) s_exp[i] = 0;
         for (int i = 0; i < 3; i++) { s_ord[i] = 0; s_urg[i] = 0; }
     }
@@ -59,16 +62,24 @@ extern "C" __global__ void monitoring_reduce(
     __nv_bfloat16 local_sq = bf16_zero();
     __nv_bfloat16 local_min = bf16(1e30f);
     __nv_bfloat16 local_max = bf16(-1e30f);
+    int local_nonzero = 0;
     int lc_exp[9] = {0,0,0,0,0,0,0,0,0};
     int lc_ord[3] = {0,0,0};
     int lc_urg[3] = {0,0,0};
 
     for (int i = tid; i < N; i += stride) {
-        __nv_bfloat16 r = bf16(rewards[i]);
-        local_sum = local_sum + r;
-        local_sq = local_sq + r * r;
-        local_min = bf16_fmin(local_min, r);
-        local_max = bf16_fmax(local_max, r);
+        float rf = rewards[i];
+        __nv_bfloat16 r = bf16(rf);
+        /* Only accumulate non-zero rewards (trade completions).
+         * Sparse reward design: reward=0.0 during hold/flat, non-zero at trade exit.
+         * Computing mean/std/sharpe over all bars masks the signal in 98% zeros. */
+        if (rf != 0.0f) {
+            local_sum = local_sum + r;
+            local_sq = local_sq + r * r;
+            local_min = bf16_fmin(local_min, r);
+            local_max = bf16_fmax(local_max, r);
+            local_nonzero++;
+        }
         /* Decode all 3 branches from factored action:
          * action = exp * (b1*b2) + ord * b2 + urg */
         int a = actions[i];
@@ -87,6 +98,7 @@ extern "C" __global__ void monitoring_reduce(
         local_sq  = local_sq + bf16_shfl_xor(0xFFFFFFFF, local_sq, mask);
         local_min = bf16_fmin(local_min, bf16_shfl_xor(0xFFFFFFFF, local_min, mask));
         local_max = bf16_fmax(local_max, bf16_shfl_xor(0xFFFFFFFF, local_max, mask));
+        local_nonzero += __shfl_xor_sync(0xFFFFFFFF, local_nonzero, mask);
         for (int i = 0; i < 9; i++)
             lc_exp[i] += __shfl_xor_sync(0xFFFFFFFF, lc_exp[i], mask);
         for (int i = 0; i < 3; i++) {
@@ -99,25 +111,33 @@ extern "C" __global__ void monitoring_reduce(
         atomicAddBF16(&s_sq_sum, local_sq);
         atomicMinBF16(&s_min, local_min);
         atomicMaxBF16(&s_max, local_max);
+        atomicAdd(&s_nonzero, local_nonzero);
         for (int i = 0; i < 9; i++) atomicAdd(&s_exp[i], lc_exp[i]);
         for (int i = 0; i < 3; i++) { atomicAdd(&s_ord[i], lc_ord[i]); atomicAdd(&s_urg[i], lc_urg[i]); }
     }
     __syncthreads();
 
     if (tid == 0) {
-        __nv_bfloat16 n_bf = bf16((float)N);
-        __nv_bfloat16 mean = s_sum / n_bf;
-        __nv_bfloat16 var = s_sq_sum / n_bf - mean * mean;
+        /* Per-trade statistics: divide by number of trade completions (non-zero rewards),
+         * not total bars. With sparse rewards (98% zero), per-bar mean ≈ 0 always.
+         * Per-trade mean gives a meaningful signal for monitoring. */
+        int n_trades = s_nonzero;
+        __nv_bfloat16 denom = (n_trades > 0) ? bf16((float)n_trades) : bf16(1.0f);
+        __nv_bfloat16 mean = s_sum / denom;
+        __nv_bfloat16 var = (n_trades > 1)
+            ? s_sq_sum / denom - mean * mean
+            : bf16_zero();
         __nv_bfloat16 std_val = bf16_sqrt(bf16_fmax(var, bf16_zero()));
         summary[0] = mean;
         summary[1] = std_val;
-        summary[2] = s_min;
-        summary[3] = s_max;
+        summary[2] = (n_trades > 0) ? s_min : bf16_zero();
+        summary[3] = (n_trades > 0) ? s_max : bf16_zero();
         summary[4] = (std_val > bf16(1e-8f)) ? mean / std_val : bf16_zero();
         for (int i = 0; i < 9; i++) summary[5 + i] = bf16((float)s_exp[i]);
         for (int i = 0; i < 3; i++) summary[14 + i] = bf16((float)s_ord[i]);
         for (int i = 0; i < 3; i++) summary[17 + i] = bf16((float)s_urg[i]);
-        summary[20] = n_bf;
-        for (int i = 21; i < 24; i++) summary[i] = bf16_zero();
+        summary[20] = bf16((float)N);
+        summary[21] = bf16((float)n_trades);
+        for (int i = 22; i < 24; i++) summary[i] = bf16_zero();
     }
 }
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 316bc7c60..f0e62f99d 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -1864,9 +1864,9 @@ impl DQNTrainer {
             if let Ok(summary) = mon.download_summary() {
                 if summary.total_experiences > 0 {
                     info!(
-                        "GPU epoch summary: mean_reward={:.6}, std={:.6}, sharpe={:.3}, actions={:?}",
+                        "GPU epoch summary: mean_reward={:.6}, std={:.6}, sharpe={:.3}, trades={}, actions={:?}",
                         summary.mean_reward, summary.reward_std, summary.sharpe_estimate,
-                        summary.action_counts
+                        summary.total_trades, summary.action_counts
                     );
                     monitor.track_reward(summary.mean_reward);
                     // Feed all 3 branch distributions from GPU into the monitor.