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

feat: 9-exposure CUDA kernels — formula-based fraction, generalized flat_idx

Browse Source 
- 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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-23 09:11:54 +01:00
parent 1277e8d10e
commit e04edc0a84
5 changed files with 28 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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

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

@@ -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 */

25
crates/ml/src/cuda_pipeline/common_device_functions.cuh
View File

@@ -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;
}

10
crates/ml/src/cuda_pipeline/epsilon_greedy_kernel.cu
View File

@@ -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;
}

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

@@ -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) {