gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix: monitoring reducer computes per-trade reward stats instead of per-bar

Browse Source 
Sparse trade-completion rewards (98% of bars = 0.0) made mean_reward
always ~0.0, hiding the actual learning signal. Now the GPU monitoring
kernel only accumulates non-zero rewards, giving meaningful per-trade
mean, std, sharpe, and trade count in the epoch summary log.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-02 13:58:53 +02:00
parent 408188e045
commit 32ff00b0cb
3 changed files with 45 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
crates/ml/src/cuda_pipeline/gpu_monitoring.rs
View File

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

44
crates/ml/src/cuda_pipeline/monitoring_kernel.cu
View File

@@ -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();
}
}

4
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

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