diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 6af86fc31..a4a9e0ccf 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -476,15 +476,26 @@ extern "C" __global__ void c51_loss_batched( * When Q-values differentiate, this converges to argmax. * When Q-values are flat, this averages → stable learning signal. */ - /* Compute softmax weights from expected Q per action */ + /* Compute softmax weights from expected Q per action. + * Temperature = max(Q_gap/3, 0.001). When Q-values differentiate, + * tau shrinks → weights sharpen toward argmax. When flat, tau=0.001 + * still produces exp(0.06/0.001)=exp(60) → effectively argmax for + * ANY non-zero Q-difference. This prevents the uniform-averaging + * fixed point where all actions get equal weight → zero gradient + * to push Q-values apart. */ __shared__ float action_weights[MAX_BRANCH_SIZE]; if (tid == 0) { float max_eq = eq_per_action[0]; - for (int a = 1; a < n_d; a++) + float min_eq = eq_per_action[0]; + for (int a = 1; a < n_d; a++) { max_eq = fmaxf(max_eq, eq_per_action[a]); + min_eq = fminf(min_eq, eq_per_action[a]); + } + float q_gap_local = max_eq - min_eq; + float tau = fmaxf(q_gap_local / 3.0f, 0.001f); float sum_exp = 0.0f; for (int a = 0; a < n_d; a++) { - action_weights[a] = expf(eq_per_action[a] - max_eq); + action_weights[a] = expf((eq_per_action[a] - max_eq) / tau); sum_exp += action_weights[a]; } for (int a = 0; a < n_d; a++)