From e04edc0a84cd96f83b7ad93472a3bfebb93aa537 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 23 Mar 2026 09:11:54 +0100 Subject: [PATCH] =?UTF-8?q?feat:=209-exposure=20CUDA=20kernels=20=E2=80=94?= =?UTF-8?q?=20formula-based=20fraction,=20generalized=20flat=5Fidx?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - common_device_functions.cuh: DQN_NUM_ACTIONS 5→9, TOTAL_ACTIONS 45→81 - experience_kernels.cu: exposure_idx_to_fraction() formula (-1+idx*0.25) - experience_kernels.cu: action_to_exposure() same formula - backtest_env_kernel.cu: switch→formula for 9-level mapping - All Q-gap flat_idx: hardcoded 2 → b0_size/2 (generalized) - c51_loss_kernel.cu: MAX_BRANCH_SIZE 5→9, BRANCH_0_SIZE 5→9 Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/backtest_env_kernel.cu | 13 +++------- .../ml/src/cuda_pipeline/c51_loss_kernel.cu | 4 +-- .../cuda_pipeline/common_device_functions.cuh | 25 +++++++++---------- .../cuda_pipeline/epsilon_greedy_kernel.cu | 10 ++++---- .../src/cuda_pipeline/experience_kernels.cu | 23 +++++------------ 5 files changed, 28 insertions(+), 47 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 73c8bf834..734b1264a 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -78,16 +78,9 @@ 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-4) to target exposure - float target_exposure; - switch (actions[w]) { - case 0: target_exposure = -1.0f; break; // Short100 - case 1: target_exposure = -0.5f; break; // Short50 - case 2: target_exposure = 0.0f; break; // Flat - case 3: target_exposure = 0.5f; break; // Long50 - case 4: target_exposure = 1.0f; break; // Long100 - default: target_exposure = 0.0f; break; - } + // Map action (0-8) to target exposure: -1.0 + idx * 0.25 + int action_val = actions[w]; + float target_exposure = -1.0f + (float)action_val * 0.25f; target_exposure *= max_position; // Execute trade if position changes diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index b1cf2843b..f7baf1581 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -26,7 +26,7 @@ #define NUM_ATOMS 51 #endif #ifndef BRANCH_0_SIZE -#define BRANCH_0_SIZE 5 +#define BRANCH_0_SIZE 9 #endif #ifndef BRANCH_1_SIZE #define BRANCH_1_SIZE 3 @@ -48,7 +48,7 @@ #define BRANCH_2_ATOMS (BRANCH_2_SIZE * NUM_ATOMS) /* Largest branch output size (floats) — used to size shmem regions */ #define MAX_BRANCH_ATOMS (MAX_BRANCH_SIZE * NUM_ATOMS) -#define MAX_BRANCH_SIZE 5 /* max(BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE) */ +#define MAX_BRANCH_SIZE 9 /* max(BRANCH_0_SIZE, BRANCH_1_SIZE, BRANCH_2_SIZE) */ #define NUM_BRANCHES 3 /* Maximum warp count for BLOCK_THREADS = 256 */ #define NUM_WARPS 8 /* 256 / 32 */ diff --git a/crates/ml/src/cuda_pipeline/common_device_functions.cuh b/crates/ml/src/cuda_pipeline/common_device_functions.cuh index 18b8c84dc..29a3b7b7d 100644 --- a/crates/ml/src/cuda_pipeline/common_device_functions.cuh +++ b/crates/ml/src/cuda_pipeline/common_device_functions.cuh @@ -44,17 +44,17 @@ __device__ __forceinline__ __nv_bfloat16 f32_to_bf16(float x) { #ifndef OFI_DIM #define OFI_DIM 0 /* Default: no OFI features (safe fallback for kernels that don't use OFI) */ #endif -#define DQN_NUM_ACTIONS 5 +#define DQN_NUM_ACTIONS 9 #define DQN_ORDER_ACTIONS 3 #define DQN_URGENCY_ACTIONS 3 -#define DQN_TOTAL_ACTIONS 45 /* 5 * 3 * 3 factored action space */ -#define PPO_NUM_ACTIONS 45 +#define DQN_TOTAL_ACTIONS 81 /* 9 * 3 * 3 factored action space */ +#define PPO_NUM_ACTIONS 45 /* PPO unchanged — separate action space */ #define NUM_ACTIONS DQN_NUM_ACTIONS /* default for DQN kernels */ /* Branching DQN head sizes (Tavakoli et al., 2018). - * 3 independent advantage heads: exposure (5), order (3), urgency (3). - * Factored action index = exposure * 9 + order * 3 + urgency (0-44). */ -#define BRANCH_EXPOSURE_ACTIONS 5 + * 3 independent advantage heads: exposure (9), order (3), urgency (3). + * Factored action index = exposure * 9 + order * 3 + urgency (0-80). */ +#define BRANCH_EXPOSURE_ACTIONS 9 #define BRANCH_ORDER_ACTIONS 3 #define BRANCH_URGENCY_ACTIONS 3 #define BRANCH_MAX_ACTIONS BRANCH_EXPOSURE_ACTIONS /* largest branch */ @@ -912,14 +912,13 @@ __device__ __forceinline__ void barrier_reset(float* barrier_state) { /** * Diversity penalty based on Shannon entropy over a sliding window - * of recent actions, tracked at the exposure level (5 categories). + * of recent actions, tracked at the exposure level (9 categories). * - * With DQN_NUM_ACTIONS=5, each action IS an exposure level directly. - * Returns -0.1 if entropy < 1.5 (low diversity penalty), 0.0 otherwise. - * Max entropy for 5 categories = log2(5) = 2.32; threshold 1.5 ≈ 65%. + * With DQN_NUM_ACTIONS=9, each action IS an exposure level directly. + * Max entropy for 9 categories = log2(9) = 3.17. * Caller applies diversity_scale as a multiplier. * - * diversity_window[DIVERSITY_WINDOW]: ring buffer of exposure indices (0-4). + * diversity_window[DIVERSITY_WINDOW]: ring buffer of exposure indices (0-8). * diversity_meta[2]: [write_pos, count]. */ __device__ float diversity_entropy( @@ -927,7 +926,7 @@ __device__ float diversity_entropy( int* diversity_meta, int action_idx ) { - /* With 5-action DQN, action_idx IS the exposure level (0-4) */ + /* With 9-action DQN, action_idx IS the exposure level (0-8) */ int category = action_idx; if (category < 0) category = 0; if (category >= DQN_NUM_ACTIONS) category = DQN_NUM_ACTIONS - 1; @@ -967,7 +966,7 @@ __device__ float diversity_entropy( * Max entropy for 5 categories = log2(5) ≈ 2.32. * Penalty scales linearly with how far entropy is below maximum. * Caller applies diversity_scale as multiplier. */ - float max_entropy = log2f((float)DQN_NUM_ACTIONS); /* 2.32 for 5 actions */ + float max_entropy = log2f((float)DQN_NUM_ACTIONS); /* 3.17 for 9 actions */ if (entropy < max_entropy) { return -((max_entropy - entropy) / max_entropy) * 0.1f; } diff --git a/crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu b/crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu index 6d7a84f36..68763dadf 100644 --- a/crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu +++ b/crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu @@ -50,7 +50,7 @@ extern "C" __global__ void epsilon_greedy_select( } /* Q-gap conviction filter: force flat when conviction is low */ if (q_gap_threshold > 0.0f && num_actions > 2) { - float q_flat = q[2]; /* index 2 = Flat */ + float q_flat = q[num_actions / 2]; /* center = Flat */ if (best_q - q_flat < q_gap_threshold) { action = 2; } @@ -124,9 +124,9 @@ extern "C" __global__ void epsilon_greedy_routed( } /* Q-gap conviction filter */ if (q_gap_threshold > 0.0f && num_actions > 2) { - float q_flat = q[2]; + float q_flat = q[num_actions / 2]; if (best_q - q_flat < q_gap_threshold) { - exposure_idx = 2; + exposure_idx = num_actions / 2; } } } @@ -149,7 +149,7 @@ extern "C" __global__ void epsilon_greedy_routed( ); /* Step 4: Override to Flat if not filled */ - int final_action = filled ? exposure_idx : 2; /* 2 = Flat */ + int final_action = filled ? exposure_idx : (num_actions / 2); /* center = Flat */ rng_states[idx] = rng; actions_out[idx] = (unsigned int)final_action; @@ -200,7 +200,7 @@ extern "C" __global__ void branching_action_select( } /* Q-gap filter: force flat when conviction is low */ if (q_gap_threshold > 0.0f) { - float q_flat = qe[2]; /* index 2 = Flat */ + float q_flat = qe[4]; /* index 4 = Flat (center of 9-action) */ if (best - q_flat < q_gap_threshold) { exposure = 2; } diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 209407949..1f353e539 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -112,14 +112,9 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) { * 4 → 1.0 (Long100) */ __device__ __forceinline__ float exposure_idx_to_fraction(int idx) { - switch (idx) { - case 0: return -1.0f; - case 1: return -0.5f; - case 2: return 0.0f; - case 3: return 0.5f; - case 4: return 1.0f; - default: return 0.0f; - } + /* 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; } /* ================================================================== */ @@ -264,7 +259,7 @@ extern "C" __global__ void experience_action_select( a0 = argmax_n(q_b0, b0_size); /* Q-gap filter: force flat when conviction is low */ if (q_gap_threshold > 0.0f && b0_size > 2) { - int flat_idx = 2; /* exposure_idx_to_fraction(2) = 0.0 */ + int flat_idx = b0_size / 2; /* center = Flat (index 4 for 9-action) */ float q_best = q_b0[a0]; float q_flat = q_b0[flat_idx]; if (q_best - q_flat < q_gap_threshold) { @@ -661,14 +656,8 @@ extern "C" __global__ void experience_env_step( * ══════════════════════════════════════════════════════════════════════════ */ __device__ __forceinline__ float action_to_exposure(int action_idx) { - switch (action_idx) { - case 0: return -1.0f; - case 1: return -0.5f; - case 2: return 0.0f; - case 3: return 0.5f; - case 4: return 1.0f; - default: return 0.0f; - } + /* 9 levels: same formula as exposure_idx_to_fraction */ + return -1.0f + (float)action_idx * 0.25f; } __device__ __forceinline__ float action_to_tx_cost(int action_idx) {