From 39158ecab2f3f1fa5ca1ea2b5edb0a0d45bbeb46 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 28 Feb 2026 12:10:36 +0100 Subject: [PATCH] =?UTF-8?q?fix(cuda):=20spec=20compliance=20=E2=80=94=20ba?= =?UTF-8?q?rrier=5Fdone,=20step=5Fin=5Fepisode,=20diversity=20penalty?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three fixes from spec review: 1. diversity_entropy returns -0.1 penalty (threshold < 1.0) instead of raw entropy 2. barrier_done included in episode termination (barrier hit ends episode) 3. step_in_episode counter resets properly on mid-loop episode reset Co-Authored-By: Claude Opus 4.6 --- .../cuda_pipeline/dqn_experience_kernel.cu | 29 ++++++++++++------- 1 file changed, 19 insertions(+), 10 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu b/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu index 5e3a3c23a..e12ecb3ac 100644 --- a/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu +++ b/crates/ml/src/cuda_pipeline/dqn_experience_kernel.cu @@ -268,8 +268,11 @@ __device__ __forceinline__ void barrier_reset(float* barrier_state) { } /** - * Shannon entropy over a sliding window of recent actions, - * mapped to 3 categories: Short (0), Flat (1), Long (2). + * Diversity penalty based on Shannon entropy over a sliding window + * of recent actions, mapped to 3 categories: Short (0), Flat (1), Long (2). + * + * Returns -0.1 if entropy < 1.0 (low diversity penalty), 0.0 otherwise. + * Caller applies diversity_scale as a multiplier. * * diversity_window[DIVERSITY_WINDOW]: ring buffer of category indices. * diversity_meta[2]: [write_pos, count]. @@ -316,7 +319,10 @@ __device__ float diversity_entropy( entropy -= p * log2f(p); } } - return entropy; /* max = log2(3) ~ 1.585 */ + + /* Penalty for low diversity: -0.1 if entropy < 1.0, else 0.0. + * Caller applies diversity_scale as multiplier. */ + return (entropy < 1.0f) ? -0.1f : 0.0f; } /** @@ -503,6 +509,7 @@ extern "C" __global__ void dqn_full_experience_kernel( float cur_scratch[CUR_HIDDEN]; float state[STATE_DIM]; + int step_in_episode = 0; /* ---- Main episode loop ---- */ for (int t = 0; t < L; t++) { @@ -649,8 +656,8 @@ extern "C" __global__ void dqn_full_experience_kernel( barrier_reset(barrier_st); } - /* ---- Step 8: Diversity entropy ---- */ - float div_ent = diversity_entropy(div_window, div_meta, action_idx); + /* ---- Step 8: Diversity penalty ---- */ + float div_penalty = diversity_entropy(div_window, div_meta, action_idx); /* ---- Step 9: Mark-to-market -> next_state ---- */ float next_price = (next_close_raw != 0.0f) ? next_close_raw : next_close; @@ -718,14 +725,15 @@ extern "C" __global__ void dqn_full_experience_kernel( } float combined_reward = pnl_reward * barrier_mult - + diversity_scale * div_ent + + diversity_scale * div_penalty + curiosity_scale * curiosity_reward; /* ---- Step 13: Episode done check ---- */ - int step_in_episode = t + 1; - int time_done = (step_in_episode >= episode_length) ? 1 : 0; - int data_done = (next_bar >= total_bars) ? 1 : 0; - int done = (time_done || data_done) ? 1 : 0; + step_in_episode++; + int time_done = (step_in_episode >= episode_length) ? 1 : 0; + int barrier_done = (barrier_label != 0) ? 1 : 0; + int data_done = (next_bar >= total_bars) ? 1 : 0; + int done = (time_done || barrier_done || data_done) ? 1 : 0; /* ---- Step 14: TD error ---- */ float td_target = combined_reward + gamma * max_target_q * (1.0f - (float)done); @@ -744,6 +752,7 @@ extern "C" __global__ void dqn_full_experience_kernel( entry_price = 0.0f; cum_costs = 0.0f; last_price = 0.0f; + step_in_episode = 0; barrier_reset(barrier_st); /* Reset diversity window */ for (int i = 0; i < DIVERSITY_WINDOW; i++)