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

fix: 8 critical bugs — tx cost 10000x, action decode, state mismatch, monitoring

Browse Source 
CATASTROPHIC fixes:
1. TX cost missing *0.0001f bps conversion — trades cost $17K instead of $1
2. Backtest action decode: raw factored int→800% exposure (should decode exposure_idx)
3. State mismatch: training 66 features, backtest 45 — Q-values at eval were garbage
4. Position scaling disabled: CVaR/conviction/Kelly destroyed credit assignment

Monitoring fixes:
5. Q-value labels: 5-element array→9-element for 9-action exposure space
6. Monitoring kernel: add common_device_functions.cuh for DQN_ORDER_ACTIONS defines
7. Backtest metrics: factored action decode for buy/sell/hold classification
8. Backtest gather: add common_device_functions.cuh for MARKET_DIM/PORTFOLIO_DIM

Result: agent now trades 71K times (was 1) with all 9 exposure levels explored.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-24 15:32:39 +01:00
parent ddac49509b
commit 432fcdc7ee
12 changed files with 317 additions and 216 deletions
Download Patch File Download Diff File Expand all files Collapse all files

30
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
View File

@@ -19,7 +19,7 @@ extern "C" __global__ void backtest_env_step(
const int* __restrict__ window_lens, // [n_windows]
// Actions from model for current step
const int* __restrict__ actions, // [n_windows] (0-4: Short100..Long100)
const int* __restrict__ actions, // [n_windows] factored action index
// Portfolio state (read-write, persistent across steps)
float* portfolio_state, // [n_windows * PORTFOLIO_STATE_SIZE]
@@ -36,7 +36,11 @@ extern "C" __global__ void backtest_env_step(
float max_position,
float tx_cost_bps,
float spread_cost,
int current_step
int current_step,
// Branching DQN action space sizes
int b0_size, // exposure branch size (default 9)
int b1_size, // order type branch size (default 3)
int b2_size // urgency branch size (default 3)
) {
// Shared memory tile for coalesced portfolio reads/writes
__shared__ float shmem_pf[256 * PORTFOLIO_STATE_SIZE];
@@ -78,11 +82,29 @@ extern "C" __global__ void backtest_env_step(
float max_equity = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 4];
float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6];
// Map action (0-8) to target exposure: -1.0 + idx * 0.25
// Decode factored action to exposure index (matches experience_env_step).
// Factored action = exposure_idx * b1_size * b2_size + order_idx * b2_size + urgency_idx
int action_val = actions[w];
float target_exposure = -1.0f + (float)action_val * 0.25f;
int exposure_idx;
if (b1_size > 0 && b2_size > 0) {
int denom = b1_size * b2_size;
exposure_idx = action_val / denom;
} else {
exposure_idx = action_val;
}
// Clamp to valid range
if (exposure_idx < 0) exposure_idx = 0;
if (exposure_idx >= b0_size) exposure_idx = b0_size - 1;
// Map exposure index to fraction: -1.0 + idx * (2.0 / (b0_size - 1))
// With b0_size=9: step=0.25, range [-1.0, +1.0]
float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
float target_exposure = -1.0f + (float)exposure_idx * step_size;
target_exposure *= max_position;
// Suppress unused variable warnings
(void)open; (void)high; (void)low;
// Execute trade if position changes
float delta = target_exposure - position;
float trade_cost = 0.0f;

253
crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
View File

@@ -1,9 +1,11 @@
// Gather state vectors from pre-uploaded features + live portfolio state.
// Output: [n_windows, state_dim] tensor for model forward pass.
//
// This kernel eliminates the CPU roundtrip in the old gather_states() path,
// which previously downloaded the full features buffer (n_windows * max_len * feat_dim
// floats) to CPU just to slice out a single step's row per window.
// State layout MATCHES the training experience_state_gather kernel:
// [0 .. market_dim) : market features (42)
// [market_dim .. market_dim+8) : 8 portfolio features
// [market_dim+8 .. market_dim+8+16) : 16 multi-timeframe features
// [market_dim+24 .. state_dim) : zero-pad / OFI / alignment
//
// Portfolio state layout per window [8 floats]:
// [0] value - current portfolio value
@@ -15,17 +17,19 @@
// [6] cum_return - cumulative log return
// [7] step_count - number of completed steps
//
// State output layout per window [state_dim floats]:
// When ofi_dim == 0 (no OFI):
// [0 .. feat_dim) - market features
// [feat_dim + 0..2] - portfolio (value, position, spread)
// [feat_dim + 3 .. state_dim) - zero pad
// 8 portfolio features (matches training experience_state_gather):
// +0: position — raw position [-1, +1]
// +1: unrealized_pnl / equity — how is THIS trade doing?
// +2: drawdown — (peak - equity) / peak [0, 1]
// +3: hold_time / 100.0 — normalized holding duration (step_count as proxy)
// +4: realized_pnl / equity — cumulative return signal
// +5: distance_to_floor — (equity - floor) / equity [0, 1]
// +6: trade_return — unrealized PnL / equity
// +7: cash / equity — available capital ratio
//
// When ofi_dim > 0 (OFI reorder):
// [0 .. market_dim) - market features (feat_dim - ofi_dim)
// [market_dim + 0..2] - portfolio (value, position, spread)
// [market_dim + 3 .. market_dim + 3 + ofi_dim) - OFI features
// [market_dim + 3 + ofi_dim .. state_dim) - zero pad
// 16 multi-timeframe features (4 lookbacks × 4 features):
// Lookback windows: 5, 15, 60, 240 bars
// Per window: return, volatility, volume_trend, momentum
extern "C" __global__ void gather_states(
const float* __restrict__ features, // [n_windows * max_len * feat_dim]
@@ -61,84 +65,165 @@ extern "C" __global__ void gather_states(
int feat_base = (w * max_len + current_step) * feat_dim;
int out_base = w * state_dim;
// Portfolio features (shared between both paths) — read from shared memory tile
float norm_value = (initial_capital > 0.0f)
? shmem_portfolio[local_tid * 8 + 0] / initial_capital
: 0.0f;
float position = shmem_portfolio[local_tid * 8 + 1];
// Read portfolio state from shared memory tile
float value = shmem_portfolio[local_tid * 8 + 0];
float position = shmem_portfolio[local_tid * 8 + 1];
float cash = shmem_portfolio[local_tid * 8 + 2];
float entry_price = shmem_portfolio[local_tid * 8 + 3];
float max_equity = shmem_portfolio[local_tid * 8 + 4];
float step_count = shmem_portfolio[local_tid * 8 + 7];
if (ofi_dim > 0) {
// OFI reorder: storage is [market, OFI], model expects [market, portfolio, OFI, pad]
int market_dim = feat_dim - ofi_dim;
// Determine market_dim: if OFI is present, market features are feat_dim - ofi_dim
int market_dim = (ofi_dim > 0) ? (feat_dim - ofi_dim) : feat_dim;
// 1. Market features [0 .. market_dim) — __ldg + manual 4x unroll
int i = 0;
for (; i + 3 < market_dim; i += 4) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
}
for (; i < market_dim; i++) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
}
// ── 1. Market features [0 .. market_dim) ── //
int i = 0;
for (; i + 3 < market_dim; i += 4) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
}
for (; i < market_dim; i++) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
}
// 2. Portfolio [market_dim .. market_dim + 3)
states_out[out_base + market_dim + 0] = norm_value;
states_out[out_base + market_dim + 1] = position;
states_out[out_base + market_dim + 2] = spread_cost;
// ── 2. Portfolio features [market_dim .. market_dim+8) ── //
// Match training experience_state_gather exactly.
float equity = (value > 1.0f) ? value : 1.0f;
// 3. OFI features [market_dim + 3 .. market_dim + 3 + ofi_dim)
// ofi_dim=8 is always float4-aligned, full 4x unroll
for (i = 0; i + 3 < ofi_dim; i += 4) {
states_out[out_base + market_dim + 3 + i] = __ldg(&features[feat_base + market_dim + i]);
states_out[out_base + market_dim + 3 + i + 1] = __ldg(&features[feat_base + market_dim + i + 1]);
states_out[out_base + market_dim + 3 + i + 2] = __ldg(&features[feat_base + market_dim + i + 2]);
states_out[out_base + market_dim + 3 + i + 3] = __ldg(&features[feat_base + market_dim + i + 3]);
}
for (; i < ofi_dim; i++) {
states_out[out_base + market_dim + 3 + i] = __ldg(&features[feat_base + market_dim + i]);
}
// Drawdown from peak
float drawdown = (max_equity > 1.0f)
? (max_equity - value) / max_equity
: 0.0f;
if (drawdown < 0.0f) drawdown = 0.0f;
// 4. Zero-pad [market_dim + 3 + ofi_dim .. state_dim) — 4x unroll
i = market_dim + 3 + ofi_dim;
for (; i + 3 < state_dim; i += 4) {
states_out[out_base + i] = 0.0f;
states_out[out_base + i + 1] = 0.0f;
states_out[out_base + i + 2] = 0.0f;
states_out[out_base + i + 3] = 0.0f;
}
for (; i < state_dim; i++) {
states_out[out_base + i] = 0.0f;
}
} else {
// No OFI: straight copy [market, portfolio, pad]
// Market features with __ldg and manual 4x unroll
int i = 0;
for (; i + 3 < feat_dim; i += 4) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
}
for (; i < feat_dim; i++) {
states_out[out_base + i] = __ldg(&features[feat_base + i]);
}
// Unrealized PnL: position * (current_close - entry_price)
// Use feature[0] as close price proxy (same as training kernel)
float close_now = __ldg(&features[feat_base]);
float unrealized_pnl = (entry_price > 0.0f && position != 0.0f)
? position * (close_now - entry_price)
: 0.0f;
states_out[out_base + feat_dim + 0] = norm_value;
states_out[out_base + feat_dim + 1] = position;
states_out[out_base + feat_dim + 2] = spread_cost;
// Realized PnL: derive from equity state
// In backtest, realized PnL = value - initial_capital (approximate)
float realized_pnl = value - initial_capital;
// Zero-pad [feat_dim + 3 .. state_dim) — 4x unroll
i = feat_dim + 3;
for (; i + 3 < state_dim; i += 4) {
states_out[out_base + i] = 0.0f;
states_out[out_base + i + 1] = 0.0f;
states_out[out_base + i + 2] = 0.0f;
states_out[out_base + i + 3] = 0.0f;
}
for (; i < state_dim; i++) {
states_out[out_base + i] = 0.0f;
// Trade return (unrealized / equity)
float trade_return = unrealized_pnl / equity;
// Capital floor distance (75% of peak)
float floor = max_equity * 0.75f;
float floor_dist = (equity > floor && equity > 1.0f)
? (equity - floor) / equity
: 0.0f;
int pf_base = market_dim;
if (pf_base + 7 < state_dim) {
states_out[out_base + pf_base + 0] = position; // raw position
states_out[out_base + pf_base + 1] = unrealized_pnl / equity; // trade P&L signal
states_out[out_base + pf_base + 2] = drawdown; // risk: how deep are we?
states_out[out_base + pf_base + 3] = step_count / 100.0f; // how long in trade?
states_out[out_base + pf_base + 4] = realized_pnl / equity; // session P&L
states_out[out_base + pf_base + 5] = floor_dist; // distance to game over
states_out[out_base + pf_base + 6] = trade_return; // this trade's return
states_out[out_base + pf_base + 7] = cash / equity; // available capital
}
// ── 3. Multi-timeframe features [market_dim+8 .. market_dim+8+16) ── //
// 4 lookback windows × 4 features = 16 GPU-native features.
// Matches training experience_state_gather exactly.
const int lookbacks[4] = {5, 15, 60, 240};
int mtf_base = market_dim + 8;
// Global bar index for this window at current step
int window_bar_base = w * max_len;
for (int lb = 0; lb < 4; lb++) {
int N_lb = lookbacks[lb];
int past_step = current_step - N_lb;
int slot = mtf_base + lb * 4;
if (past_step >= 0 && slot + 3 < state_dim) {
int now_feat_off = (window_bar_base + current_step) * feat_dim;
int past_feat_off = (window_bar_base + past_step) * feat_dim;
float close_cur = __ldg(&features[now_feat_off]);
float close_past = __ldg(&features[past_feat_off]);
// Return over N bars
float ret = (close_past > 0.0f) ? (close_cur - close_past) / close_past : 0.0f;
float scaled_ret = ret * 100.0f;
states_out[out_base + slot + 0] = fmaxf(-10.0f, fminf(10.0f, scaled_ret));
// Volatility: scan high/low over window (from close changes)
float max_val = close_cur;
float min_val = close_cur;
float vol_sum = 0.0f;
int vol_count = 0;
for (int j = past_step; j <= current_step; j++) {
int j_off = (window_bar_base + j) * feat_dim;
float v = __ldg(&features[j_off]);
if (v > max_val) max_val = v;
if (v < min_val) min_val = v;
// Volume proxy: feature index 4
if (feat_dim > 4) {
vol_sum += __ldg(&features[j_off + 4]);
vol_count++;
}
}
float range = (close_cur > 0.0f) ? (max_val - min_val) / close_cur : 0.0f;
float scaled_range = range * 100.0f;
states_out[out_base + slot + 1] = fmaxf(0.0f, fminf(10.0f, scaled_range));
// Volume trend: current vs average
float avg_vol = (vol_count > 0) ? vol_sum / (float)vol_count : 1.0f;
float cur_vol = (feat_dim > 4) ? __ldg(&features[now_feat_off + 4]) : 1.0f;
float vol_ratio = (avg_vol > 0.0f) ? cur_vol / avg_vol : 1.0f;
states_out[out_base + slot + 2] = fmaxf(0.0f, fminf(5.0f, vol_ratio));
// Momentum: position within range [0=bottom, 1=top]
float range_size = max_val - min_val;
states_out[out_base + slot + 3] = (range_size > 0.0f)
? (close_cur - min_val) / range_size
: 0.5f;
} else {
// Not enough history — zero pad
if (slot + 3 < state_dim) {
states_out[out_base + slot + 0] = 0.0f;
states_out[out_base + slot + 1] = 0.0f;
states_out[out_base + slot + 2] = 0.0f;
states_out[out_base + slot + 3] = 0.0f;
}
}
}
// ── 4. OFI features (if present) ── //
if (ofi_dim > 0) {
int ofi_out_base = market_dim + 8 + 16;
// OFI features are stored after market features in the feature vector
for (i = 0; i + 3 < ofi_dim; i += 4) {
states_out[out_base + ofi_out_base + i] = __ldg(&features[feat_base + market_dim + i]);
states_out[out_base + ofi_out_base + i + 1] = __ldg(&features[feat_base + market_dim + i + 1]);
states_out[out_base + ofi_out_base + i + 2] = __ldg(&features[feat_base + market_dim + i + 2]);
states_out[out_base + ofi_out_base + i + 3] = __ldg(&features[feat_base + market_dim + i + 3]);
}
for (; i < ofi_dim; i++) {
states_out[out_base + ofi_out_base + i] = __ldg(&features[feat_base + market_dim + i]);
}
}
// ── 5. Zero-pad remaining [filled .. state_dim) ── //
int filled = market_dim + 8 + 16 + ofi_dim;
for (i = filled; i + 3 < state_dim; i += 4) {
states_out[out_base + i] = 0.0f;
states_out[out_base + i + 1] = 0.0f;
states_out[out_base + i + 2] = 0.0f;
states_out[out_base + i + 3] = 0.0f;
}
for (; i < state_dim; i++) {
states_out[out_base + i] = 0.0f;
}
// Suppress unused variable warning
(void)spread_cost;
}

28
crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
View File

@@ -12,10 +12,10 @@
// [7] cvar_95 (mean of returns below VaR — Expected Shortfall)
// [8] calmar_ratio (annualized mean return / max drawdown)
// [9] omega_ratio (sum of gains / sum of losses)
// [10] buy_count (actions 3,4 — Long50, Long100)
// [11] sell_count (actions 0,1 — Short100, Short50)
// [10] buy_count (exposure_idx 5-8 — Long25..Long100)
// [11] sell_count (exposure_idx 0-3 — Short100..Short25)
// [12] hold_count (action 2 — Flat)
// [13] unique_action_count (popcount of 5-bit mask — how many of 0..4 were used)
// [13] unique_action_count (popcount of 9-bit mask — how many exposure levels 0..8 were used)
extern "C" __global__ void compute_backtest_metrics(
const float* __restrict__ step_returns, // [n_windows * max_len]
@@ -65,7 +65,7 @@ extern "C" __global__ void compute_backtest_metrics(
float local_cum = 0.0f, local_peak = 0.0f, local_max_dd = 0.0f;
int local_wins = 0, local_trades = 0;
int local_buys = 0, local_sells = 0, local_holds = 0;
int local_action_mask = 0; // 5-bit bitmask: bit i set if action i was seen
int local_action_mask = 0; // 9-bit bitmask: bit i set if exposure_idx i was seen
int prev_action = -1;
for (int i = tid; i < wlen; i += stride) {
@@ -89,13 +89,17 @@ extern "C" __global__ void compute_backtest_metrics(
if (act != prev_action && i > 0) local_trades++;
prev_action = act;
// Action distribution: 0,1=sell 2=hold 3,4=buy
if (act <= 1) local_sells++;
else if (act == 2) local_holds++;
else local_buys++;
// Action distribution: decode exposure from factored action.
// Factored action = exposure * (b1*b2) + order * b2 + urgency.
// exposure_idx: 0-3=sell, 4=flat, 5-8=buy (for 9-action space)
int exp_idx = act / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
if (exp_idx < DQN_NUM_ACTIONS / 2) local_sells++;
else if (exp_idx == DQN_NUM_ACTIONS / 2) local_holds++;
else local_buys++;
// Unique action tracking: set bit for each distinct action ID (0..4)
if (act >= 0 && act < 5) local_action_mask |= (1 << act);
// Unique action tracking: set bit per exposure index (0..8)
if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS)
local_action_mask |= (1 << exp_idx);
}
// Store ALL local values to shared memory
@@ -148,8 +152,8 @@ extern "C" __global__ void compute_backtest_metrics(
// ── Unique action count via bitmask OR-reduction ─────────────────────────
// Reuses s_sorted memory (safe: OR-reduction completes before bitonic sort).
// Each thread accumulated a 5-bit mask of observed action IDs. OR-reduce
// across the block, then popcount gives unique_action_count (1-5).
// Each thread accumulated a 9-bit mask of observed exposure indices. OR-reduce
// across the block, then popcount gives unique_action_count (1-9).
{
int* s_action_mask = (int*)s_sorted; // alias — sequential use, no conflict
s_action_mask[tid] = local_action_mask;

12
crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
View File

@@ -94,11 +94,15 @@ extern "C" __global__ void curiosity_forward_backward(
input[MARKET_DIM + 1] = 0.0f;
input[MARKET_DIM + 2] = 0.0f;
/* Action to category one-hot (5-action DQN mapping from common_device_functions.cuh) */
/* Action to category one-hot — decode exposure from factored action.
* Factored action = exposure_idx * (b1*b2) + order_idx * b2 + urgency_idx.
* With b0=9, b1=3, b2=3: exposure_idx = action / 9.
* Categories: Short (idx 0-3), Flat (idx 4), Long (idx 5-8). */
int exposure_idx = action_idx / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
int category;
if (action_idx <= 1) category = 0; /* Short100/Short50 */
else if (action_idx == 2) category = 1; /* Flat */
else category = 2; /* Long50/Long100 */
if (exposure_idx < DQN_NUM_ACTIONS / 2) category = 0; /* Short */
else if (exposure_idx == DQN_NUM_ACTIONS / 2) category = 1; /* Flat */
else category = 2; /* Long */
input[MARKET_DIM + category] = 1.0f;
/* Layer 1: pre_act = w1 * input + b1, hidden = LeakyReLU(pre_act, alpha=0.01) */

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

@@ -104,16 +104,21 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) {
}
/**
* Map exposure branch index (0–4) to target exposure fraction.
* 0 → -1.0 (Short100)
* 1 → -0.5 (Short50)
* 2 → 0.0 (Flat)
* 3 → 0.5 (Long50)
* 4 → 1.0 (Long100)
* Map exposure branch index (0–8) to target exposure fraction.
* 0 → -1.00 (S100)
* 1 → -0.75 (S75)
* 2 → -0.50 (S50)
* 3 → -0.25 (S25)
* 4 → 0.00 (Flat)
* 5 → +0.25 (L25)
* 6 → +0.50 (L50)
* 7 → +0.75 (L75)
* 8 → +1.00 (L100)
*
* Formula: -1.0 + idx * (2.0 / (b0_size - 1)).
* With b0_size=9: step = 0.25. Flat = index 4.
*/
__device__ __forceinline__ float exposure_idx_to_fraction(int idx) {
/* 9 levels: -1.0, -0.75, -0.50, -0.25, 0.0, +0.25, +0.50, +0.75, +1.0
* Formula: -1.0 + idx * 0.25. Flat = index 4. */
return -1.0f + (float)idx * 0.25f;
}
@@ -597,84 +602,31 @@ extern "C" __global__ void experience_env_step(
float target_exposure = exposure_idx_to_fraction(exposure_idx);
float target_position = target_exposure * max_position;
/* Apply CVaR risk scaling: reduce position when tail risk is high */
if (cvar_scales != NULL) {
target_position *= cvar_scales[i];
}
/* Apply Q-gap conviction scaling: size UP with conviction.
* conviction = clamp(q_gap / 2.0, 0.25, 1.0)
* Low Q-gap (0.1-0.5) → small position (25-50% of target)
* High Q-gap (2.0+) → full position (100%)
* Combined with CVaR: size = exposure × CVaR_scale × conviction */
if (q_gaps != NULL) {
float conviction = q_gaps[i] / 2.0f;
conviction = (conviction < 0.25f) ? 0.25f : ((conviction > 1.0f) ? 1.0f : conviction);
target_position *= conviction;
}
/* ════════════════════════════════════════════════════════════════
* GPU-NATIVE ENHANCED KELLY CRITERION (full implementation)
/* CVaR, conviction, and Kelly scaling DISABLED during training.
*
* Weighted blend of two Kelly estimators:
* 1. Continuous Kelly: f* = μ / σ² (Gaussian approximation)
* 2. Discrete Kelly: f* = (b*p - q) / b (Bernoulli model)
* Enhanced: f* = 0.7 × continuous + 0.3 × discrete
* These overrides destroy credit assignment: the agent selects an exposure
* level, but the actual position is silently rescaled by CVaR, Q-gap
* conviction, and Kelly fraction. The reward for action A then reflects
* the outcome of a DIFFERENT position size, corrupting Q-value learning.
*
* With confidence scaling from sample size:
* confidence = 1 - 1/√(n_trades) (20 trades = 0.78, 100 = 0.90)
* final_f = f* × confidence × fractional_kelly (0.5 = half-Kelly)
*
* All statistics accumulated in portfolio state ps[14]-ps[19].
* Zero CPU involvement. Pure GPU. Updated only at trade completions.
* ════════════════════════════════════════════════════════════════ */
{
float total_trades = win_count + loss_count;
if (total_trades >= 20.0f) {
/* ── Continuous Kelly: f* = μ / σ² ── */
float mu = sum_returns / total_trades;
float variance = (sum_sq_returns / total_trades) - (mu * mu);
float continuous_kelly = (variance > 1e-10f)
? mu / variance
: 0.0f;
* Stop-loss/take-profit and capital floor are kept — they are environment
* rules (hard constraints), not action overrides. */
(void)cvar_scales;
(void)q_gaps;
/* ── Discrete Kelly: f* = (b*p - q) / b ── */
float p = win_count / total_trades;
float q = 1.0f - p;
float avg_win = sum_wins / fmaxf(win_count, 1.0f);
float avg_loss = sum_losses / fmaxf(loss_count, 1.0f);
float b = avg_win / fmaxf(avg_loss, 1e-6f);
float discrete_kelly = (b > 1e-6f)
? (b * p - q) / b
: 0.0f;
/* Kelly criterion DISABLED during training (same rationale as CVaR/conviction above).
* Kelly statistics (ps[14]-ps[19]) are still accumulated for logging/diagnostics,
* but the position size is NOT rescaled. The agent must learn its own sizing
* through the reward signal, not have it overridden by a running Kelly estimate. */
(void)win_count; (void)loss_count;
(void)sum_wins; (void)sum_losses;
(void)sum_returns; (void)sum_sq_returns;
/* ── Enhanced Kelly: weighted blend ── */
/* NaN guard: if either Kelly produces NaN/Inf, use 0.0 (safe default) */
if (isnan(continuous_kelly) || isinf(continuous_kelly)) continuous_kelly = 0.0f;
if (isnan(discrete_kelly) || isinf(discrete_kelly)) discrete_kelly = 0.0f;
float enhanced_kelly = 0.7f * continuous_kelly + 0.3f * discrete_kelly;
/* ── Confidence scaling: penalize small samples ── */
float confidence = 1.0f - 1.0f / sqrtf(total_trades);
/* ── Fractional Kelly (half-Kelly = 0.5) ── */
float fractional = 0.5f;
float kelly_frac = enhanced_kelly * confidence * fractional;
/* ── Clamp to [0.05, 0.25] ── */
kelly_frac = fmaxf(0.05f, fminf(0.25f, kelly_frac));
/* ── Apply to position (normalize: 0.25 = full size) ── */
target_position *= kelly_frac / 0.25f;
}
/* < 20 trades: no Kelly applied (full position until we have statistics) */
/* Final NaN/Inf guard on target_position after ALL scaling.
* If any scaling (CVaR, conviction, Kelly) produced NaN, zero the position. */
/* Final NaN/Inf guard on target_position.
* If exposure_idx_to_fraction produced NaN, zero the position. */
if (isnan(target_position) || isinf(target_position)) {
target_position = 0.0f;
}
}
/* ---- Position adjustment with volatility-scaled transaction cost ---- */
/* Real spread widens in volatile markets. vol_scale (from CUSUM) is
@@ -707,7 +659,7 @@ extern "C" __global__ void experience_env_step(
: (order_type_idx == 1) ? 0.0002f /* IoC: +2 bps */
: -0.0005f; /* LimitMaker: -5 bps rebate */
tx_cost = fabsf(delta) * raw_close * (tx_cost_multiplier * spread_scale * impact_scale + order_premium);
tx_cost = fabsf(delta) * raw_close * (tx_cost_multiplier * 0.0001f * spread_scale * impact_scale + order_premium);
cash -= delta * raw_close;
cash -= tx_cost;
position = target_position;
@@ -905,7 +857,7 @@ extern "C" __global__ void experience_env_step(
if (auto_stop_loss || auto_take_profit || auto_time_stop) {
/* Force exit — override model's action to flat */
position = 0.0f;
float exit_cost = fabsf(ps[0]) * raw_close * tx_cost_multiplier;
float exit_cost = fabsf(ps[0]) * raw_close * tx_cost_multiplier * 0.0001f;
cash = ps[1] + ps[0] * raw_close - exit_cost;
is_flat = 1.0f;
exiting_trade = 1;

70
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -79,13 +79,25 @@ fn compile_env_ptx(context: &CudaContext) -> Result<Ptx, String> {
}
fn compile_metrics_ptx(context: &CudaContext) -> Result<Ptx, String> {
let defines = "\
#define STATE_DIM 72\n\
#define MARKET_DIM 42\n\
#define PORTFOLIO_DIM 8\n";
let common_src = include_str!("common_device_functions.cuh");
let src = include_str!("backtest_metrics_kernel.cu");
crate::cuda_pipeline::compile_ptx_for_device(src, context)
let full_source = format!("{defines}{common_src}\n{src}");
crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
}
fn compile_gather_ptx(context: &CudaContext) -> Result<Ptx, String> {
let defines = "\
#define STATE_DIM 72\n\
#define MARKET_DIM 42\n\
#define PORTFOLIO_DIM 8\n";
let common_src = include_str!("common_device_functions.cuh");
let src = include_str!("backtest_gather_kernel.cu");
crate::cuda_pipeline::compile_ptx_for_device(src, context)
let full_source = format!("{defines}{common_src}\n{src}");
crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
}
/// Compile PPO backtest forward kernel with common device functions prepended.
@@ -570,11 +582,12 @@ impl GpuBacktestEvaluator {
.alloc_zeros::<f32>(n_windows * 14)
.map_err(|e| MLError::ModelError(format!("metrics alloc: {e}")))?;
// Portfolio dimension is always 3: (normalised value, position, spread_cost).
// state_dim = feature_dim + 3, then 8-aligned for H100 tensor core HMMA dispatch.
// The CUDA gather_states kernel zero-pads positions [feat_dim+3 .. state_dim).
const PORTFOLIO_DIM: usize = 8;
let state_dim = (feature_dim + PORTFOLIO_DIM + 7) & !7;
// State layout matches training experience_state_gather:
// [0..42) market features + [42..50) 8 portfolio + [50..66) 16 multi-timeframe
// Total filled = feat_dim + 8 portfolio + 16 multi-timeframe = feat_dim + 24
// 8-aligned for H100 tensor core HMMA dispatch.
const PORTFOLIO_AND_MTF_DIM: usize = 24; // 8 portfolio + 16 multi-timeframe
let state_dim = (feature_dim + PORTFOLIO_AND_MTF_DIM + 7) & !7;
let states_buf = stream
.alloc_zeros::<f32>(n_windows * state_dim)
.map_err(|e| MLError::ModelError(format!("states_buf alloc: {e}")))?;
@@ -611,7 +624,7 @@ impl GpuBacktestEvaluator {
n_windows,
max_len,
feature_dim,
portfolio_dim: PORTFOLIO_DIM,
portfolio_dim: PORTFOLIO_AND_MTF_DIM,
state_dim,
config,
forward_actions_buf,
@@ -691,8 +704,8 @@ impl GpuBacktestEvaluator {
/// CUDA storage — eliminating the GPU→CPU→GPU roundtrip entirely.
///
/// # Panics
/// `portfolio_dim` must equal `self.portfolio_dim` (always 3). Callers using a
/// different value should be updated — the kernel signature is fixed.
/// `portfolio_dim` must equal `self.portfolio_dim` (24: 8 portfolio + 16 multi-timeframe).
/// This matches the training experience_state_gather layout.
pub fn gather_states(
&self,
step: usize,
@@ -1315,7 +1328,7 @@ impl GpuBacktestEvaluator {
/// * `forward_fn` — model forward: `&CudaSlice<f32>[N * state_dim]` → `CudaSlice<f32>[N]` predictions
/// * `high_threshold_bps` — strong signal threshold (actions 0 and 4)
/// * `low_threshold_bps` — mild signal threshold (actions 1 and 3)
/// * `portfolio_dim` — must be 3
/// * `portfolio_dim` — must be 24 (8 portfolio + 16 multi-timeframe)
pub fn evaluate_supervised<F>(
&mut self,
forward_fn: &F,
@@ -1731,6 +1744,10 @@ impl GpuBacktestEvaluator {
let n_win_i32 = self.n_windows as i32;
let max_len_i32 = self.max_len as i32;
let step_i32 = step as i32;
// Branch sizes for factored action decoding (default 9/3/3)
let b0_i32: i32 = 9;
let b1_i32: i32 = 3;
let b2_i32: i32 = 3;
unsafe {
self.stream
@@ -1749,6 +1766,9 @@ impl GpuBacktestEvaluator {
.arg(&self.config.tx_cost_bps)
.arg(&self.config.spread_cost)
.arg(&step_i32)
.arg(&b0_i32)
.arg(&b1_i32)
.arg(&b2_i32)
.launch(env_cfg)
.map_err(|e| {
MLError::ModelError(format!("backtest_env_step launch step {step}: {e}"))
@@ -1980,13 +2000,14 @@ mod tests {
}
}
/// Verify that gather_states() returns a ConfigError when portfolio_dim != 3.
/// Verify that gather_states() returns a ConfigError when portfolio_dim != 24.
#[test]
fn test_gather_states_portfolio_dim_mismatch() {
// We don't have a CUDA device in unit tests, so test the validation path
// indirectly by checking the struct's portfolio_dim field is always 3.
// The real mismatch check is exercised at runtime when portfolio_dim != 3.
assert_eq!(3_usize, 3_usize, "PORTFOLIO_DIM constant is 3");
// indirectly by checking the struct's portfolio_dim field is always 24.
// The real mismatch check is exercised at runtime when portfolio_dim != 24.
// 24 = 8 portfolio features + 16 multi-timeframe features.
assert_eq!(24_usize, 24_usize, "PORTFOLIO_AND_MTF_DIM constant is 24");
}
/// Verify that the cuBLAS DQN helper kernels (compute_expected_q,
@@ -2031,16 +2052,19 @@ mod tests {
#[test]
fn test_gpu_backtest_evaluator_state_dim_calculation() {
// state_dim must equal feature_dim + PORTFOLIO_DIM (3)
// state_dim = (feature_dim + 24 + 7) & !7
// 24 = 8 portfolio + 16 multi-timeframe (matches training layout)
let feature_dim: usize = 42;
let portfolio_dim: usize = 3;
let state_dim = feature_dim + portfolio_dim;
assert_eq!(state_dim, 45);
let portfolio_and_mtf: usize = 24;
let state_dim = (feature_dim + portfolio_and_mtf + 7) & !7;
// 42 + 24 = 66, aligned to 72
assert_eq!(state_dim, 72);
// With 8-alignment padding (56 features + 3)
let feature_dim_aligned: usize = 53; // 53-feature state
let state_dim_aligned = feature_dim_aligned + portfolio_dim;
assert_eq!(state_dim_aligned, 56);
// With OFI (feature_dim = 50 = 42 market + 8 OFI in feature vector)
let feature_dim_ofi: usize = 50;
let state_dim_ofi = (feature_dim_ofi + portfolio_and_mtf + 7) & !7;
// 50 + 24 = 74, aligned to 80
assert_eq!(state_dim_ofi, 80);
}
/// Verify that CUDA Graph infrastructure compiles and struct fields are present.

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

@@ -34,8 +34,14 @@ pub struct GpuMonitoringReducer {
static MONITORING_PTX: OnceLock<Result<Ptx, String>> = OnceLock::new();
fn compile_monitoring_ptx(context: &CudaContext) -> Result<Ptx, String> {
let defines = "\
#define STATE_DIM 72\n\
#define MARKET_DIM 42\n\
#define PORTFOLIO_DIM 8\n";
let common_src = include_str!("common_device_functions.cuh");
let kernel_src = include_str!("monitoring_kernel.cu");
crate::cuda_pipeline::compile_ptx_for_device(kernel_src, context)
let full_source = format!("{defines}{common_src}\n{kernel_src}");
crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
}
impl GpuMonitoringReducer {

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

@@ -57,8 +57,11 @@ extern "C" __global__ void monitoring_reduce(
local_sq += r * r;
local_min = fminf(local_min, r);
local_max = fmaxf(local_max, r);
/* Decode exposure index from factored action for monitoring.
* Factored action = exposure * (b1*b2) + order * b2 + urgency. */
int a = actions[i];
if (a >= 0 && a < 9) local_counts[a]++;
int exp_idx = a / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS) local_counts[exp_idx]++;
}
// Warp-level reduction before atomic to shared (reduces contention 32x)

2
crates/ml/src/hyperopt/adapters/ppo.rs
View File

@@ -1328,7 +1328,7 @@ impl PPOTrainer {
MLError::ModelError(format!("PPO backtest action upload: {e}"))
})
},
3, // portfolio_dim
24, // portfolio_dim: 8 portfolio + 16 multi-timeframe (matches training state layout)
)?;
let first = metrics.first().ok_or_else(|| {

2
crates/ml/src/trainers/dqn/config.rs
View File

@@ -1566,7 +1566,7 @@ use crate::dqn::dqn::DQNConfig;
pub(crate) fn dqn_default_config() -> DQNConfig {
DQNConfig {
// Network Architecture
state_dim: 48, // 42 market + 3 portfolio = 45, aligned to 48 for tensor cores
state_dim: 48, // Default placeholder; overridden by constructor to (42+8+16+7)&!7=72
num_actions: 9, // 9 exposure levels (25% steps)
hidden_dims: vec![512, 256, 128], // Sufficient capacity without overfitting (Trial 19)
leaky_relu_alpha: 0.01,

7
crates/ml/src/trainers/dqn/trainer/metrics.rs
View File

@@ -241,7 +241,7 @@ impl DQNTrainer {
let row_offset = r * cols;
let mut best_val = f32::NEG_INFINITY;
let mut second_best = f32::NEG_INFINITY;
for c in 0..cols.min(5) {
for c in 0..cols.min(9) {
let v = host_q.get(row_offset + c).copied().unwrap_or(0.0);
if let Some(s) = col_sums.get_mut(c) {
*s += v as f64;
@@ -255,8 +255,9 @@ impl DQNTrainer {
// v <= second_best: no update needed
}
}
// Also check remaining columns (6..total_actions) for best/second_best
for c in 5..cols {
// Also check remaining columns (9..total_actions) for best/second_best
// (order_type + urgency branches, if total_actions > 9)
for c in 9..cols {
let v = host_q.get(row_offset + c).copied().unwrap_or(0.0);
if v > best_val {
second_best = best_val;

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

@@ -1884,7 +1884,7 @@ impl DQNTrainer {
gap_mean, gap_min, gap_max
);
let names = ["S100", "S50", "Flat", "L50", "L100"];
let names = ["S100", "S75", "S50", "S25", "Flat", "L25", "L50", "L75", "L100"];
let parts: Vec<String> = names.iter()
.zip(per_action_avgs.iter())
.map(|(n, q)| format!("{}={:.4}", n, q))