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

feat(cuda): add DSR + n-step to scalar experience kernel

Browse Source 
Adds Differential Sharpe Ratio reward shaping and n-step return
accumulation to dqn_full_experience_kernel (scalar path). DSR replaces
EMA normalization when use_dsr=1; n-step ring buffer accumulates
discounted returns for effective_n>1. Both are reset on episode
boundaries.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-10 13:03:01 +01:00
parent 923fab9919
commit aee3871788
1 changed files with 158 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

182
crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu
View File

@@ -1085,6 +1085,23 @@ extern "C" __global__ void dqn_full_experience_kernel(
const float* __restrict__ rms_v_gamma, /* [VALUE_H] */
const float* __restrict__ rms_a_gamma, /* [ADV_H] */
/* ---- Fill simulation config ---- */
float fill_median_spread,
float fill_median_vol,
float fill_ioc_fill_prob,
float fill_limit_fill_min,
float fill_limit_fill_max,
float fill_spread_cost_frac,
float fill_spread_capture_frac,
int fill_simulation_enabled,
/* ---- DSR (Differential Sharpe Ratio) config ---- */
int use_dsr,
float dsr_eta,
/* ---- N-step returns config ---- */
int n_steps,
/* ---- RNG states [N] ---- */
unsigned int* rng_states,
@@ -1166,6 +1183,18 @@ extern "C" __global__ void dqn_full_experience_kernel(
float ema_var = 1.0f;
int ema_init = 0;
/* ---- DSR accumulators (per-episode) ---- */
float dsr_A = 0.0f;
float dsr_B = 1e-8f;
int dsr_initialized = 0;
/* ---- N-step ring buffer (per-episode) ---- */
float nstep_ring[N_STEPS_MAX];
int nstep_ring_idx = 0;
int nstep_ring_len = 0;
int effective_n = (n_steps > 0 && n_steps <= N_STEPS_MAX) ? n_steps : 1;
nstep_reset(nstep_ring, &nstep_ring_idx, &nstep_ring_len, effective_n);
/* ---- Main episode loop ---- */
for (int t = 0; t < L; t++) {
int global_bar = ep_start + t;
@@ -1309,13 +1338,54 @@ extern "C" __global__ void dqn_full_experience_kernel(
action_counts[action_idx]++;
total_action_count++;
out_actions[out_off] = action_idx;
online_q_selected = q_values[action_idx];
/* ---- Step 5b: Order routing + fill simulation ---- */
/* Route action through spread/vol-dependent order type selection,
* then check fill probability. Unfilled orders override to Flat. */
if (fill_simulation_enabled) {
int order_type, urgency;
route_order(spread, fill_median_spread, 0.0f /* vol from price */, fill_median_vol,
&order_type, &urgency);
float norm_vol = (fill_median_vol > 0.0f) ? (spread / fill_median_spread) : 1.0f;
/* Clamp normalized vol to a reasonable range for fill probability */
norm_vol = fminf(fmaxf(norm_vol, 0.0f), 3.0f);
float cost_adj;
int filled = simulate_fill_check(
order_type, urgency, norm_vol,
spread * 10000.0f, /* convert spread fraction to bps */
global_bar, action_idx,
fill_ioc_fill_prob, fill_limit_fill_min, fill_limit_fill_max,
fill_spread_cost_frac, fill_spread_capture_frac,
&cost_adj
);
if (!filled) {
action_idx = 2; /* Override to Flat — order didn't fill */
}
}
out_actions[out_off] = action_idx;
/* ---- Step 6: Portfolio simulation ---- */
float target_exposure = action_to_exposure(action_idx);
float target_position = target_exposure * max_position;
float tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier;
/* Use order-type-dependent tx cost when fill sim enabled,
* otherwise fall back to fixed market rate. */
float tx_rate;
if (fill_simulation_enabled) {
int order_type_for_cost, urgency_unused;
route_order(spread, fill_median_spread, 0.0f, fill_median_vol,
&order_type_for_cost, &urgency_unused);
float spread_bps = spread * 10000.0f;
tx_rate = fabsf(order_type_tx_cost(order_type_for_cost, spread_bps,
fill_spread_cost_frac, fill_spread_capture_frac));
tx_rate *= tx_cost_multiplier;
} else {
tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier;
}
/* Detect reversal (sign change) */
int is_reversal = (position > 0.0f && target_position < 0.0f) ||
@@ -1488,27 +1558,37 @@ extern "C" __global__ void dqn_full_experience_kernel(
+ diversity_scale * div_penalty
+ curiosity_scale * curiosity_reward;
/* ---- EMA reward normalization (matches CPU RewardNormalizer) ----
* Normalize-then-update: use OLD stats to normalize, then update
* stats with the raw value. Clamp to [-3, 3] like the CPU path. */
float raw_reward = combined_reward;
if (ema_init) {
float std = sqrtf(ema_var);
if (std > 1e-8f)
combined_reward = (combined_reward - ema_mean) / std;
combined_reward = fmaxf(-3.0f, fminf(3.0f, combined_reward));
}
/* Update EMA running stats with the raw (un-normalized) value */
if (!ema_init) {
ema_mean = raw_reward;
ema_var = 1.0f;
ema_init = 1;
/* DSR: replace EMA normalization when enabled */
if (use_dsr) {
combined_reward = dsr_step(combined_reward,
&dsr_A, &dsr_B, &dsr_initialized, dsr_eta);
} else {
ema_mean = reward_norm_alpha * raw_reward
+ (1.0f - reward_norm_alpha) * ema_mean;
float diff = raw_reward - ema_mean;
ema_var = reward_norm_alpha * diff * diff
+ (1.0f - reward_norm_alpha) * ema_var;
/* ---- EMA reward normalization (matches CPU RewardNormalizer) ---- */
float raw_reward = combined_reward;
if (ema_init) {
float std = sqrtf(ema_var);
if (std > 1e-8f)
combined_reward = (combined_reward - ema_mean) / std;
combined_reward = fmaxf(-3.0f, fminf(3.0f, combined_reward));
}
if (!ema_init) {
ema_mean = raw_reward;
ema_var = 1.0f;
ema_init = 1;
} else {
ema_mean = reward_norm_alpha * raw_reward
+ (1.0f - reward_norm_alpha) * ema_mean;
float diff = raw_reward - ema_mean;
ema_var = reward_norm_alpha * diff * diff
+ (1.0f - reward_norm_alpha) * ema_var;
}
}
/* N-step: accumulate discounted return */
if (effective_n > 1) {
combined_reward = nstep_push_and_sum(
nstep_ring, &nstep_ring_idx, &nstep_ring_len,
combined_reward, gamma, effective_n);
}
/* ---- Step 13: Episode done check ---- */
@@ -1565,6 +1645,14 @@ extern "C" __global__ void dqn_full_experience_kernel(
ema_mean = 0.0f;
ema_var = 1.0f;
ema_init = 0;
/* Reset DSR accumulators */
dsr_A = 0.0f;
dsr_B = 1e-8f;
dsr_initialized = 0;
/* Reset n-step ring */
nstep_reset(nstep_ring, &nstep_ring_idx, &nstep_ring_len, effective_n);
}
} /* end if (!skip_data) */
} /* end timestep loop */
@@ -1710,6 +1798,16 @@ extern "C" __global__ void dqn_full_experience_kernel_warp(
const float* __restrict__ rms_v_gamma, /* [VALUE_H] */
const float* __restrict__ rms_a_gamma, /* [ADV_H] */
/* ---- Fill simulation config ---- */
float fill_median_spread,
float fill_median_vol,
float fill_ioc_fill_prob,
float fill_limit_fill_min,
float fill_limit_fill_max,
float fill_spread_cost_frac,
float fill_spread_capture_frac,
int fill_simulation_enabled,
/* ---- RNG states [N] ---- */
unsigned int* rng_states,
@@ -1955,13 +2053,49 @@ extern "C" __global__ void dqn_full_experience_kernel_warp(
action_idx = __shfl_sync(0xFFFFFFFF, action_idx, 0);
if (lane_id == 0) {
out_actions[out_off] = action_idx;
online_q_selected = q_values[action_idx];
/* ---- Step 5b: Order routing + fill simulation ---- */
if (fill_simulation_enabled) {
int order_type, urgency;
route_order(spread, fill_median_spread, 0.0f, fill_median_vol,
&order_type, &urgency);
float norm_vol = (fill_median_vol > 0.0f) ? (spread / fill_median_spread) : 1.0f;
norm_vol = fminf(fmaxf(norm_vol, 0.0f), 3.0f);
float cost_adj;
int filled = simulate_fill_check(
order_type, urgency, norm_vol,
spread * 10000.0f,
global_bar, action_idx,
fill_ioc_fill_prob, fill_limit_fill_min, fill_limit_fill_max,
fill_spread_cost_frac, fill_spread_capture_frac,
&cost_adj
);
if (!filled) {
action_idx = 2; /* Override to Flat */
}
}
out_actions[out_off] = action_idx;
/* ---- Step 6: Portfolio simulation ---- */
float target_exposure = action_to_exposure(action_idx);
float target_position = target_exposure * max_position;
float tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier;
float tx_rate;
if (fill_simulation_enabled) {
int order_type_for_cost, urgency_unused;
route_order(spread, fill_median_spread, 0.0f, fill_median_vol,
&order_type_for_cost, &urgency_unused);
float spread_bps = spread * 10000.0f;
tx_rate = fabsf(order_type_tx_cost(order_type_for_cost, spread_bps,
fill_spread_cost_frac, fill_spread_capture_frac));
tx_rate *= tx_cost_multiplier;
} else {
tx_rate = action_to_tx_cost(action_idx) * tx_cost_multiplier;
}
int is_reversal = (position > 0.0f && target_position < 0.0f) ||
(position < 0.0f && target_position > 0.0f);