From 2ecf36b94ce2506c36865d165ef4d00a4556f396 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 24 Apr 2026 11:28:31 +0200 Subject: [PATCH] =?UTF-8?q?refactor(dqn):=20Plan=201=20Task=204E=20?= =?UTF-8?q?=E2=80=94=20DIR=5F*=20direction=20sub-index=20named=20constants?= =?UTF-8?q?=20(Invariant=208)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Define DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3, NUM_DIRECTIONS=4 in state_layout.cuh. Migrate all literal dir_idx/raw_dir comparisons to named constants across experience_kernels.cu (15 sites), trade_physics.cuh (3 sites), and backtest_metrics_kernel.cu (2 sites). Raw integer comparisons (== 0/1/2/3) eliminated from all direction-branch consumers. Co-Authored-By: Claude Sonnet 4.6 --- .../cuda_pipeline/backtest_metrics_kernel.cu | 7 ++++--- .../src/cuda_pipeline/experience_kernels.cu | 20 +++++++++---------- crates/ml/src/cuda_pipeline/state_layout.cuh | 10 ++++++++++ crates/ml/src/cuda_pipeline/trade_physics.cuh | 8 +++++--- docs/dqn-named-dims.md | 3 ++- 5 files changed, 31 insertions(+), 17 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 439450ceb..5e7a5d354 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -1,4 +1,5 @@ // Per-window metrics reduction kernel — native F32. +#include "state_layout.cuh" // One block per window. Threads cooperate to reduce step_returns. // // Output per window [14 f32 values written to metrics_out]: @@ -105,9 +106,9 @@ extern "C" __global__ void compute_backtest_metrics( int raw_dir = exp_idx / 3; /* 0=Short, 1=Hold, 2=Long, 3=Flat */ /* Collapse Hold(1) and Flat(3) → no-exposure (-1). * Keep Short(0) and Long(2) as distinct sign states. */ - int signed_dir = (raw_dir == 0) ? -1 : - (raw_dir == 2) ? 1 : - 0; /* Hold/Flat = no exposure */ + int signed_dir = (raw_dir == DIR_SHORT) ? -1 : + (raw_dir == DIR_LONG) ? 1 : + 0; /* Hold/Flat = no exposure */ /* Action-distribution counters: now reflect ACTUAL directions, not * legacy 9-action exposure-band thresholds. Counts: long, short, diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 83f01a6b7..ce1975201 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -925,7 +925,7 @@ extern "C" __global__ void experience_action_select( /* Hold/Flat detection shortcut — when direction=Hold(1) or Flat(3), * magnitude is irrelevant. Set to 0 for deterministic encoding. */ - if (dir_idx == 1 || dir_idx == 3) { + if (dir_idx == DIR_HOLD || dir_idx == DIR_FLAT) { mag_idx = 0; /* Hold/Flat: magnitude masked, set to 0 for deterministic encoding */ } else { /* Branch 1: magnitude — training uses BOLTZMANN SOFTMAX; eval uses @@ -1612,8 +1612,8 @@ extern "C" __global__ void experience_env_step( * consistent with the mirrored price data. * Direction(4): Short(0)↔Long(2), Hold(1)→Hold(1), Flat(3)→Flat(3). */ if (mirror_active) { - if (dir_idx == 0) dir_idx = 2; /* Short → Long */ - else if (dir_idx == 2) dir_idx = 0; /* Long → Short */ + if (dir_idx == DIR_SHORT) dir_idx = DIR_LONG; /* Short → Long */ + else if (dir_idx == DIR_LONG) dir_idx = DIR_SHORT; /* Long → Short */ /* Hold(1) and Flat(3) unchanged */ } @@ -1683,7 +1683,7 @@ extern "C" __global__ void experience_env_step( } /* ── Hold action flag preserved for downstream reward/plan logic. */ - int is_hold_action = (dir_idx == 1); + int is_hold_action = (dir_idx == DIR_HOLD); /* Save pre-trade position + sign + hold_time for reversal detection * and segment-reward patience multiplier that run AFTER the helper. @@ -2477,8 +2477,8 @@ extern "C" __global__ void experience_env_step( /* Short(0)↔Long(2): opposite direction. Hold(1)/Flat(3): no useful mirror. * When dir is Hold or Flat, skip trivial self-mirror — use magnitude CF instead. */ int cf_dir; - if (dir_idx == 0) cf_dir = 2; /* Short→Long */ - else if (dir_idx == 2) cf_dir = 0; /* Long→Short */ + if (dir_idx == DIR_SHORT) cf_dir = DIR_LONG; /* Short→Long */ + else if (dir_idx == DIR_LONG) cf_dir = DIR_SHORT; /* Long→Short */ else { /* Hold/Flat: directional mirror is trivial (same action). * Fall through to magnitude cycle instead for useful counterfactual. */ @@ -2521,7 +2521,7 @@ extern "C" __global__ void experience_env_step( cf_reward = fminf(fmaxf(cf_reward, -10.0f), 10.0f); } else { /* Near-zero reward: magnitude scaling uninformative, use directional mirror */ - int cf_dir = (dir_idx == 0) ? 2 : (dir_idx == 2) ? 0 : dir_idx; /* Short↔Long, Hold/Flat stay */ + int cf_dir = (dir_idx == DIR_SHORT) ? DIR_LONG : (dir_idx == DIR_LONG) ? DIR_SHORT : dir_idx; /* Short↔Long, Hold/Flat stay */ cf_action = cf_dir * b1_size * b2_size * b3_size + mag_idx * b2_size * b3_size + orig_order * b3_size + orig_urgency; @@ -2555,7 +2555,7 @@ extern "C" __global__ void experience_env_step( cf_reward = fminf(fmaxf(cf_reward, -10.0f), 10.0f); } else { /* Flat position: order irrelevant, directional mirror fallback */ - int cf_dir = (dir_idx == 0) ? 2 : (dir_idx == 2) ? 0 : dir_idx; /* Short↔Long, Hold/Flat stay */ + int cf_dir = (dir_idx == DIR_SHORT) ? DIR_LONG : (dir_idx == DIR_LONG) ? DIR_SHORT : dir_idx; /* Short↔Long, Hold/Flat stay */ cf_action = cf_dir * b1_size * b2_size * b3_size + mag_idx * b2_size * b3_size + orig_order * b3_size + orig_urgency; @@ -3060,8 +3060,8 @@ extern "C" __global__ void expert_action_override( * 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 * 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 */ + if (expert_dir == DIR_LONG && current_pos > 0.5f) return; /* Already Long */ + if (expert_dir == DIR_SHORT && current_pos < -0.5f) return; /* Already Short */ /* Override direction+magnitude, keep Q-network's order and urgency */ int current = out_actions[i]; diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index da1fb8726..8fc5a509c 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -105,6 +105,16 @@ #define BRANCH_URG 3 // urgency branch (3 actions) #define NUM_BRANCHES 4 +// ──────────────────────────────────────────────────────────────────────────── +// Direction action sub-indices (dir_idx in [0..NUM_DIRECTIONS)). +// Invariant 8: every direction value has a named constant. +// ──────────────────────────────────────────────────────────────────────────── +#define DIR_SHORT 0 // open/maintain short position +#define DIR_HOLD 1 // keep current position (no-op) +#define DIR_LONG 2 // open/maintain long position +#define DIR_FLAT 3 // close all positions to zero +#define NUM_DIRECTIONS 4 + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index f72d98785..d7f7f134f 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -61,13 +61,15 @@ __device__ __forceinline__ float compute_target_position( * L25 L50 L100 * 0 0 0 (Flat × any = 0) */ +#include "state_layout.cuh" + __device__ __forceinline__ float compute_target_position_4branch( int dir_idx, int mag_idx, float max_position ) { /* Hold sentinel: caller must handle dir_idx==1 before calling. * Returns NaN as safety net — any arithmetic with NaN propagates visibly. */ - if (dir_idx == 1) return __int_as_float(0x7FC00000); /* NaN sentinel */ - float direction = (dir_idx == 0) ? -1.0f : (dir_idx == 2) ? 1.0f : 0.0f; + if (dir_idx == DIR_HOLD) return __int_as_float(0x7FC00000); /* NaN sentinel */ + float direction = (dir_idx == DIR_SHORT) ? -1.0f : (dir_idx == DIR_LONG) ? 1.0f : 0.0f; float magnitude = (mag_idx == 0) ? 0.25f : (mag_idx == 1) ? 0.5f : 1.0f; return direction * magnitude * max_position; } @@ -593,7 +595,7 @@ __device__ __forceinline__ void unified_env_step_core( int mag_idx = decode_magnitude_4b(action_val, b1_size, b2_size, b3_size); int order_type_idx = decode_order_4b(action_val, b2_size, b3_size); - int is_hold_action = (dir_idx == 1); + int is_hold_action = (dir_idx == DIR_HOLD); /* ── Step 2-3: Target position (uses target scale, which training shrinks) ── */ float target_position; diff --git a/docs/dqn-named-dims.md b/docs/dqn-named-dims.md index b7f5d86a3..5ec25d2e9 100644 --- a/docs/dqn-named-dims.md +++ b/docs/dqn-named-dims.md @@ -121,4 +121,5 @@ From the trade_plan MLP output. - Task 4A (ps[0..PS_STRIDE) constants): commit 144c85b85 - Task 4B (plan_isv[0..6) constants): commit 741cb48d5 - Task 4C (plan_params[0..6) constants): commit 0ac83479e -- Task 4D (BRANCH_* constants): commit +- Task 4D (BRANCH_* constants): commit 11755632b +- Task 4E (DIR_* direction sub-indices): commit TBD