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

feat(plan): PORTFOLIO_STRIDE 23→30 — 7 plan slots ps[23-29]

Browse Source 
Plan slots: target_bars, profit_target, stop_loss, scale_aggression,
conviction, asymmetry, counter_plan_q. Zeroed on episode reset.
Fixed hardcoded i*23 references. All CUDA + Rust files updated.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-17 01:13:36 +02:00
parent a1f967db88
commit c7a2e39c1e
3 changed files with 29 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -20,7 +20,7 @@
* [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 experience kernels ([N, PORTFOLIO_STRIDE=23]):
* Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=30]):
* [0] position — current contract position (signed)
* [1] cash — cash balance
* [2] portfolio_value — mark-to-market total value (cash + position * price)
@@ -36,6 +36,13 @@
* [20] intra_trade_max_dd — worst unrealized drawdown during current trade (v8)
* [21] (reserved) — was last_trade_t (clustering stats, slot retained)
* [22] (reserved) — was interval_sum (clustering stats, slot retained)
* [23] plan_target_bars — 0=no plan, >0=active plan max hold bars
* [24] plan_profit_target — raw profit threshold %
* [25] plan_stop_loss — raw stop loss threshold %
* [26] plan_scale_aggression — position ramp speed
* [27] plan_conviction — position size fraction
* [28] plan_asymmetry — profit/stop ratio
* [29] counter_plan_q — opposite direction Q at entry
*
* Branching DQN (Tavakoli et al., 2018) — 4-branch hierarchical:
* 4 independent advantage heads: direction (b0_size=3), magnitude (b1_size=3),
@@ -52,10 +59,10 @@
*/
/* ------------------------------------------------------------------ */
/* Portfolio stride for experience kernels (23 bf16 per episode). */
/* Portfolio stride for experience kernels (30 floats per episode). */
/* portfolio_sim_kernel uses its own stride of 8 — do NOT change it. */
/* ------------------------------------------------------------------ */
#define PORTFOLIO_STRIDE 23
#define PORTFOLIO_STRIDE 30
#define DSR_A_SLOT 3
#define DSR_B_SLOT 4
#define DSR_TRADE_COUNT_SLOT 5
@@ -771,7 +778,7 @@ extern "C" __global__ void experience_action_select(
int in_hold = 0;
float cur_position = 0.0f;
if (min_hold_bars > 0 && portfolio_states != NULL) {
int ps_base = i * 23; /* PORTFOLIO_STRIDE = 23 */
int ps_base = i * PORTFOLIO_STRIDE;
float hold_time_val = portfolio_states[ps_base + 10];
cur_position = portfolio_states[ps_base + 0];
/* Adaptive hold: base + ISV-driven extension (same formula as Layer 2) */
@@ -1032,7 +1039,7 @@ extern "C" __global__ void experience_action_select(
* [2] raw_close — raw close price (for position cost + tx)
* [3] raw_next — raw next-bar close (UNUSED in reward path)
*
* portfolio_states layout: [N, PORTFOLIO_STRIDE=23] (read-write, bf16)
* portfolio_states layout: [N, PORTFOLIO_STRIDE=30] (read-write, float)
* See file header for field definitions.
*/
extern "C" __global__ void experience_env_step(
@@ -1179,7 +1186,7 @@ extern "C" __global__ void experience_env_step(
/* Guard against degenerate prices from data gaps. */
if (raw_close <= 0.0f) raw_close = 1.0f;
/* ---- Read full portfolio state (PORTFOLIO_STRIDE=23) ---- */
/* ---- Read full portfolio state (PORTFOLIO_STRIDE=30) ---- */
/* Portfolio arithmetic uses float accumulators because trade physics
* functions (execute_trade, apply_margin_cap, etc.) in trade_physics.cuh
* are all float. Converting the entire physics engine to bf16 would
@@ -1796,7 +1803,7 @@ extern "C" __global__ void experience_env_step(
reward += position_entropy_weight * combined_entropy;
}
/* ---- Update full portfolio state (PORTFOLIO_STRIDE=23) ---- */
/* ---- Update full portfolio state (PORTFOLIO_STRIDE=30) ---- */
ps[0] = (position);
ps[1] = (cash);
ps[2] = (new_portfolio_value);
@@ -2003,6 +2010,10 @@ extern "C" __global__ void experience_env_step(
ps[17] = 0.0f;
ps[18] = 0.0f;
ps[19] = 0.0f;
ps[20] = 0.0f; /* intra_trade_max_dd */
ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f;
ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f;
ps[29] = 0.0f; /* plan slots zeroed */
current_timesteps[i] = 0;
} else {
/* Soft reset (trade_complete): keep equity, clear trade state only.
@@ -2013,6 +2024,9 @@ extern "C" __global__ void experience_env_step(
ps[12] = 0.0f; /* entry_price — no active trade */
ps[13] = (ps[11]); /* trade_start_pnl = current realized_pnl */
ps[20] = 0.0f; /* intra_trade_max_dd — reset for next trade */
ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f;
ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f;
ps[29] = 0.0f; /* plan slots zeroed on trade complete */
current_timesteps[i] = 0;
}
}
@@ -2393,7 +2407,7 @@ extern "C" __global__ void expert_action_override(
/* Check if current position already matches expert signal — skip if redundant.
* portfolio_states[i * PORTFOLIO_STRIDE + 0] = current position.
* If expert says Long and we're already Long (position > 0.5), skip. */
float current_pos = (portfolio_states[i * 23 + 0]); /* PORTFOLIO_STRIDE=23, pos at idx 0 */
float current_pos = (portfolio_states[i * PORTFOLIO_STRIDE + 0]);
if (expert_dir == 2 && current_pos > 0.5f) return; /* Already Long */
if (expert_dir == 0 && current_pos < -0.5f) return; /* Already Short */

10
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -41,10 +41,10 @@ const MAX_EPISODES_LIMIT: usize = 0x8000;
/// Absolute upper bound for validation — reject configs above this.
const MAX_TIMESTEPS_LIMIT: usize = 1000;
const PORTFOLIO_STATE_SIZE: usize = 8;
/// Portfolio stride for DQN experience kernels (23 floats per episode).
/// Portfolio stride for DQN experience kernels (30 floats per episode).
/// Matches PORTFOLIO_STRIDE in experience_kernels.cu.
/// Do NOT change PORTFOLIO_STATE_SIZE above — it's for the PPO/legacy path.
const PORTFOLIO_STRIDE: usize = 23;
const PORTFOLIO_STRIDE: usize = 30;
/// Number of floats in the trade_stats_reduce output buffer.
/// Layout: [win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns]
@@ -567,7 +567,7 @@ pub struct GpuExperienceCollector {
trade_stats_buf: CudaSlice<f32>,
// Per-episode state buffers [alloc_episodes, ...]
portfolio_states: CudaSlice<f32>, // [alloc_episodes * PORTFOLIO_STRIDE] (20 floats per episode)
portfolio_states: CudaSlice<f32>, // [alloc_episodes * PORTFOLIO_STRIDE] (30 floats per episode)
episode_starts_buf: CudaSlice<i32>,// [alloc_episodes]
// Output buffers [alloc_episodes * alloc_timesteps, ...]
@@ -861,7 +861,7 @@ impl GpuExperienceCollector {
}
// ── Step 6: Allocate per-episode buffers ────────────────────────
// Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=20]
// Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=30]
let mut portfolio_states = stream
.alloc_zeros::<f32>(alloc_episodes * PORTFOLIO_STRIDE)
.map_err(|e| MLError::ModelError(format!("alloc portfolio_states: {e}")))?;
@@ -2428,7 +2428,7 @@ impl GpuExperienceCollector {
_avg_spread: f32,
_cash_reserve_pct: f32,
) -> Result<(), MLError> {
// Reset portfolio states [N, PORTFOLIO_STRIDE=20]
// Reset portfolio states [N, PORTFOLIO_STRIDE=30]
let mut portfolio_init = vec![0.0_f32; self.alloc_episodes * PORTFOLIO_STRIDE];
for i in 0..self.alloc_episodes {
let off = i * PORTFOLIO_STRIDE;

4
crates/ml/src/cuda_pipeline/trade_stats_kernel.cu
View File

@@ -1,7 +1,7 @@
/**
* Trade stats reduction kernel.
*
* Reads portfolio_states[N, PORTFOLIO_STRIDE=23] and sums fields [14:19]
* Reads portfolio_states[N, PORTFOLIO_STRIDE=30] and sums fields [14:19]
* (win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns)
* across all N episodes into a 6-float output buffer.
* Uses shared-memory warp reduction (same pattern as monitoring_reduce).
@@ -9,7 +9,7 @@
* Launch config: grid=(1, 1, 1), block=(256, 1, 1).
*/
#define PORTFOLIO_STRIDE 23
#define PORTFOLIO_STRIDE 30
extern "C" __global__ void trade_stats_reduce(
const float* __restrict__ portfolio_states, // [N * PORTFOLIO_STRIDE]