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

feat: 8-component composite reward CUDA kernel

Browse Source 
Replace raw PnL reward in experience_env_step with GPU-native composite:
- DSR (Moody & Saffell 2001) with pre-update A/B formulation
- Z-scored normalized PnL with running EMA
- Drawdown penalty with peak_equity guard
- Idle penalty (replaces hold_reward)
- Regime-adaptive scaling from ADX/CUSUM features
- Asymmetric loss scaling (prospect theory)
- Position-time decay (stale position rent)
- Transaction cost (unchanged)

PORTFOLIO_STRIDE=12 (was 3). portfolio_sim_kernel stride-8 unchanged.
All division-by-zero guards per spec. ~25 FLOPs overhead (<5%).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-22 20:49:51 +01:00
parent 86863d72e6
commit aa272666a1
1 changed files with 175 additions and 55 deletions
Download Patch File Download Diff File Expand all files Collapse all files

230
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -23,10 +23,19 @@
* [market_dim .. market_dim+3) : portfolio features (value_norm, position, cash_norm)
* [market_dim+3 .. state_dim) : zero-pad for tensor-core alignment
*
* Portfolio state layout used by these kernels ([N, 3]):
* [0] position — current contract position (signed)
* [1] cash — cash balance
* [2] portfolio_value — mark-to-market total value (cash + position * price)
* Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=12]):
* [0] position — current contract position (signed)
* [1] cash — cash balance
* [2] portfolio_value — mark-to-market total value (cash + position * price)
* [3] dsr_A — EMA of returns (Moody & Saffell DSR)
* [4] dsr_B — EMA of squared returns
* [5] pnl_ema — running mean PnL (for z-score normalization)
* [6] pnl_var — running variance PnL
* [7] peak_equity — high-water mark (init to initial_capital)
* [8] flat_counter — consecutive flat steps (float for GPU simplicity)
* [9] prev_equity — equity at previous step (init to initial_capital)
* [10] hold_time — consecutive steps with position
* [11] realized_pnl — cumulative realized PnL
*
* Branching DQN (Tavakoli et al., 2018):
* 3 independent advantage heads: exposure (b0_size=5), order (b1_size=3),
@@ -35,10 +44,16 @@
* Exposure mapping: 0→-1.0, 1→-0.5, 2→0.0, 3→0.5, 4→1.0.
*
* Advanced features deferred to follow-up tasks (Phase 4+):
* DSR, curiosity, barrier tracking, diversity entropy, fill simulation,
* curiosity, barrier tracking, diversity entropy, fill simulation,
* NoisyNet, C51, RMSNorm, N-step returns, action masking.
*/
/* ------------------------------------------------------------------ */
/* Portfolio stride for experience kernels (12 floats per episode). */
/* portfolio_sim_kernel uses its own stride of 8 — do NOT change it. */
/* ------------------------------------------------------------------ */
#define PORTFOLIO_STRIDE 12
/* ------------------------------------------------------------------ */
/* Compile-time constants — overridable via NVRTC #define injection. */
/* ------------------------------------------------------------------ */
@@ -166,7 +181,7 @@ extern "C" __global__ void experience_state_gather(
* denom = max(portfolio_value, 1.0) avoids division by zero on an
* empty account. This mirrors the current_norm / pos_norm pattern
* from the monolithic kernel (spread is omitted in Phase 3). */
const float* ps = portfolio_states + (long long)i * 3;
const float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
float position = ps[0];
float cash = ps[1];
float portfolio_value = ps[2];
@@ -295,51 +310,40 @@ extern "C" __global__ void experience_action_select(
/* ================================================================== */
/**
* Portfolio simulation, reward computation and done detection.
* Portfolio simulation, 8-component composite reward, and done detection.
*
* For each episode i this kernel:
* 1. Reads the current-bar raw prices from targets[bar_idx].
* 2. Decodes the factored action to an exposure branch index.
* 3. Computes target position = exposure_fraction * max_position.
* 4. Applies position adjustment delta with a flat proportional tx cost:
* tx_cost = |delta| * raw_close * tx_cost_multiplier
* cash -= delta * raw_close + tx_cost
* 5. Computes single-step PnL:
* pnl = position_after_trade * (raw_next - raw_close)
* 6. Adds hold_reward when position ≈ 0 (flat exposure).
* 7. Writes (batch_states, action, reward, done) to the output replay
* buffer at column current_t.
* 8. Updates portfolio_states in place.
* 9. Increments current_timesteps[i].
*
* Reversal-in-two-legs, cash-reserve checks, barrier tracking, DSR,
* curiosity, diversity, fill simulation and N-step returns are deferred
* to Phase 4+ follow-up tasks.
* 4. Applies position adjustment delta with a flat proportional tx cost.
* 5. Computes composite reward:
* reward = eff_w_dsr * dsr_t (Sharpe contribution)
* + eff_w_pnl * asym_pnl (asymmetric normalized PnL)
* - eff_w_dd * drawdown_penalty_t (drawdown penalty)
* - w_idle * idle_penalty_t (idle penalty)
* - hold_time_penalty (position staleness rent)
* - tx_cost (transaction friction)
* 6. Writes (batch_states, action, reward, done) to output replay buffer.
* 7. Updates portfolio_states[0..11] in place.
* 8. Increments current_timesteps[i].
*
* Grid: ceil(N / 256), Block: 256. One thread per episode.
*
* targets layout: [total_bars, 4]
* [0] preproc_close — log-return-normalised close (used only by the
* network; env_step uses raw prices for PnL)
* [0] preproc_close — log-return-normalised close (network only)
* [1] preproc_next — log-return-normalised next close (unused here)
* [2] raw_close — raw close price (for position cost + tx)
* [3] raw_next — raw next-bar close (for mark-to-market PnL)
*
* portfolio_states layout: [N, 3] (read-write)
* [0] position
* [1] cash
* [2] portfolio_value
*
* out_states: [N, L, state_dim] — written at column current_t
* out_actions: [N, L]
* out_rewards: [N, L]
* out_dones: [N, L]
* portfolio_states layout: [N, PORTFOLIO_STRIDE=12] (read-write)
* See file header for field definitions.
*
* @param targets [total_bars, 4] price data
* @param episode_starts [N] global bar offset per episode
* @param current_timesteps [N] current episode step (incremented)
* @param actions [N] factored action from action_select
* @param portfolio_states [N, 3] position/cash/portfolio_value (read-write)
* @param portfolio_states [N, 12] full per-episode state (read-write)
* @param out_states [N, L, state_dim] replay output: states
* @param out_actions [N, L] replay output: actions
* @param out_rewards [N, L] replay output: rewards
@@ -347,7 +351,16 @@ extern "C" __global__ void experience_action_select(
* @param batch_states [N, state_dim] current state batch from gather kernel
* @param max_position max allowed contract position (absolute)
* @param tx_cost_multiplier proportional transaction cost rate
* @param hold_reward reward bonus for holding flat
* @param w_dsr DSR (Sharpe) weight
* @param w_pnl normalized PnL weight
* @param w_dd drawdown penalty weight
* @param w_idle idle penalty weight
* @param dd_threshold drawdown tolerance before penalty
* @param loss_aversion asymmetric loss scaling factor
* @param time_decay_rate position staleness rent per step
* @param eta DSR/normalization EMA decay rate
* @param features [total_bars, market_dim] market features for regime detection
* @param market_dim MARKET_DIM (42)
* @param L episode length (replay buffer columns)
* @param N number of episodes
* @param total_bars length of targets/market_features
@@ -370,7 +383,16 @@ extern "C" __global__ void experience_env_step(
const float* __restrict__ batch_states,
float max_position,
float tx_cost_multiplier,
float hold_reward,
float w_dsr,
float w_pnl,
float w_dd,
float w_idle,
float dd_threshold,
float loss_aversion,
float time_decay_rate,
float eta,
const float* __restrict__ features,
int market_dim,
int L,
int N,
int total_bars,
@@ -416,17 +438,27 @@ extern "C" __global__ void experience_env_step(
if (raw_close <= 0.0f) raw_close = 1.0f;
if (raw_next <= 0.0f) raw_next = raw_close;
/* ---- Read portfolio state ---- */
float* ps = portfolio_states + (long long)i * 3;
float position = ps[0];
float cash = ps[1];
/* ---- Read full portfolio state (PORTFOLIO_STRIDE=12) ---- */
float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
float position = ps[0];
float cash = ps[1];
/* ps[2] = portfolio_value (updated at end) */
float dsr_A = ps[3];
float dsr_B = ps[4];
float pnl_ema = ps[5];
float pnl_var = ps[6];
float peak_equity = ps[7];
float flat_counter = ps[8];
float prev_equity = ps[9];
float hold_time = ps[10];
/* ps[11] = realized_pnl (cumulative, updated at end) */
/* ---- Decode exposure index from factored action ---- */
int exposure_idx;
if (b1_size > 0 && b2_size > 0) {
/* Branching: exposure branch = action / (b1_size * b2_size) */
int denom = b1_size * b2_size;
exposure_idx = action_idx / denom;
int denom_act = b1_size * b2_size;
exposure_idx = action_idx / denom_act;
} else {
/* Flat: action IS the exposure index */
exposure_idx = action_idx;
@@ -443,36 +475,124 @@ extern "C" __global__ void experience_env_step(
float tx_cost = 0.0f;
if (delta != 0.0f) {
/* Transaction cost: proportional to trade size × price × rate. */
tx_cost = fabsf(delta) * raw_close * tx_cost_multiplier;
/* Cash flows: receive/pay for position change + pay tx cost. */
cash -= delta * raw_close;
cash -= tx_cost;
position = target_position;
}
/* ---- Mark-to-market PnL for this timestep ---- */
float pnl = position * (raw_next - raw_close);
float raw_pnl = position * (raw_next - raw_close);
/* ---- Reward: PnL minus transaction costs ---- */
float reward = pnl - tx_cost;
/* ==== Component 1: Differential Sharpe Ratio (Moody & Saffell 2001) ==== */
/* Return: guard division by zero on prev_equity */
float equity_denom = (prev_equity > 1.0f) ? prev_equity : 1.0f;
float R_t = raw_pnl / equity_denom;
/* Small bonus for holding flat, matching the monolithic kernel's
* hold_reward term that prevents Flat from being structurally
* disadvantaged (Flat always produces ~0 PnL). */
if (fabsf(position) < 0.001f) {
reward += hold_reward;
/* Compute deltas BEFORE updating EMA (pre-update values for numerator+denominator) */
float delta_A = R_t - dsr_A;
float delta_B = R_t * R_t - dsr_B;
float variance = dsr_B - dsr_A * dsr_A;
float abs_var = (variance >= 0.0f) ? variance : -variance;
float var_32 = abs_var * sqrtf(abs_var + 1e-8f); /* |variance|^(3/2) */
float dsr_denom = (var_32 > 1e-12f) ? var_32 : 1e-12f;
/* DSR from pre-update state (Moody & Saffell correct formulation) */
float dsr_t = (dsr_B * delta_A - 0.5f * dsr_A * delta_B) / dsr_denom;
/* NOW update EMA */
dsr_A = dsr_A + eta * delta_A;
dsr_B = dsr_B + eta * delta_B;
/* ==== Component 2: Normalized PnL (z-scored) ==== */
pnl_ema = pnl_ema + eta * (raw_pnl - pnl_ema);
float pnl_diff = raw_pnl - pnl_ema;
pnl_var = pnl_var + eta * (pnl_diff * pnl_diff - pnl_var);
float pnl_std = sqrtf(pnl_var + 1e-8f);
float normalized_pnl = pnl_diff / pnl_std;
/* ==== Component 6: Asymmetric profit/loss scaling (prospect theory) ==== */
float asym_pnl = (normalized_pnl >= 0.0f)
? normalized_pnl
: normalized_pnl * loss_aversion;
/* ==== Component 3: Drawdown penalty ==== */
float equity = cash + position * raw_next;
peak_equity = (equity > peak_equity) ? equity : peak_equity;
float drawdown = 0.0f;
if (peak_equity > 0.0f) {
drawdown = (peak_equity - equity) / peak_equity;
}
float drawdown_penalty_t = (drawdown > dd_threshold)
? (drawdown - dd_threshold)
: 0.0f;
/* ==== Component 4: Idle penalty (replaces hold_reward) ==== */
float is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
if (is_flat > 0.5f) {
flat_counter += 1.0f;
} else {
flat_counter = 0.0f;
}
/* max idle penalty capped at 0.05 */
float idle_penalty_t = flat_counter * 0.001f;
idle_penalty_t = (idle_penalty_t < 0.05f) ? idle_penalty_t : 0.05f;
/* ==== Component 7: Position-time decay (stale position rent) ==== */
float hold_time_penalty = 0.0f;
if (is_flat < 0.5f) {
hold_time += 1.0f;
hold_time_penalty = hold_time * time_decay_rate;
} else {
hold_time = 0.0f;
}
/* ==== Component 5: Regime-adaptive weight scaling ==== */
/* Read ADX and CUSUM from market features (already in GPU memory) */
float adx = 0.0f;
float cusum = 0.0f;
if (bar_idx < total_bars && market_dim >= 42) {
long long feat_off = (long long)bar_idx * market_dim;
adx = features[feat_off + 40];
cusum = features[feat_off + 41];
}
/* trend_scale = clamp(adx / 25.0, 0.5, 2.0) */
float trend_scale = adx / 25.0f;
trend_scale = (trend_scale < 0.5f) ? 0.5f : ((trend_scale > 2.0f) ? 2.0f : trend_scale);
/* vol_scale = clamp(cusum / 0.5, 0.5, 2.0) */
float vol_scale = cusum / 0.5f;
vol_scale = (vol_scale < 0.5f) ? 0.5f : ((vol_scale > 2.0f) ? 2.0f : vol_scale);
float eff_w_pnl = w_pnl * trend_scale;
float eff_w_dsr = w_dsr * (2.0f - trend_scale);
float eff_w_dd = w_dd * vol_scale;
/* ==== Composite reward ==== */
float reward = eff_w_dsr * dsr_t
+ eff_w_pnl * asym_pnl
- eff_w_dd * drawdown_penalty_t
- w_idle * idle_penalty_t
- hold_time_penalty
- tx_cost;
/* ---- Done detection ---- */
int next_bar = bar_idx + 1;
int done = (next_bar >= total_bars) ? 1 : 0;
/* ---- Update portfolio state ---- */
/* ---- Update full portfolio state (PORTFOLIO_STRIDE=12) ---- */
float new_portfolio_value = cash + position * raw_next;
ps[0] = position;
ps[1] = cash;
ps[2] = new_portfolio_value;
ps[0] = position;
ps[1] = cash;
ps[2] = new_portfolio_value;
ps[3] = dsr_A;
ps[4] = dsr_B;
ps[5] = pnl_ema;
ps[6] = pnl_var;
ps[7] = peak_equity;
ps[8] = flat_counter;
ps[9] = new_portfolio_value; /* prev_equity = current equity for next step */
ps[10] = hold_time;
ps[11] = ps[11] + raw_pnl; /* realized_pnl accumulates */
/* ---- Write reward and done flag ---- */
out_rewards[out_off] = reward;