diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 6e30240e4..123116969 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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;