feat: CVaR objective for Bellman target -- risk-sensitive Expected SARSA

Replace E[Q] with CVaR_alpha[Q] in c51_loss_kernel softmax weights.
alpha = 0.5 - 0.4 * iqn_readiness: starts balanced, becomes
risk-averse as IQN converges. Uses pinned device-mapped pointer
for graph-safe iqn_readiness updates.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/c51_loss_kernel.cu

@@ -222,6 +222,9 @@ extern "C" __global__ void c51_loss_batched(
     /* ── Online-target Q-divergence for adaptive tau ── */
     float* __restrict__ q_divergence,             /* [1] atomicAdd accumulator: sum of (E[Q_online]-E[Q_target])^2 */
 
+    /* ── CVaR alpha from IQN readiness ── */
+    const float* __restrict__ iqn_readiness_ptr,   /* [1] pinned device-mapped for CVaR alpha */
+
     /* ── Adam step counter for stochastic Expected SARSA seed ── */
     const int* __restrict__ step_counter           /* [1] device-mapped, increments every training step */
 ) {
@@ -464,7 +467,30 @@ extern "C" __global__ void c51_loss_batched(
                 block_log_softmax_f(shmem_lp, shmem_lp, shmem_reduce, tid, num_atoms);
                 eq = block_expected_q_f(shmem_lp, shmem_support, shmem_reduce, tid, num_atoms);
             }
-            if (tid == 0) eq_per_action[a] = eq;
+            /* CVaR: mean of atoms in bottom alpha-percentile.
+             * alpha = 0.5 - 0.4 * iqn_readiness (balanced->risk-averse). */
+            if (tid == 0) {
+                float alpha = 0.5f - 0.4f * fminf(fmaxf(iqn_readiness_ptr[0], 0.0f), 1.0f);
+                /* shmem_lp contains log-probs; convert to probs for CDF */
+                float cum_prob = 0.0f;
+                float cvar_sum = 0.0f;
+                float cvar_weight = 0.0f;
+                for (int j = 0; j < num_atoms; j++) {
+                    float prob = expf(shmem_lp[j]);
+                    float z_val = shmem_support[j];
+                    float new_cum = cum_prob + prob;
+                    if (cum_prob < alpha) {
+                        /* This atom is at least partially in the alpha tail */
+                        float contrib = fminf(prob, alpha - cum_prob);
+                        cvar_sum += contrib * z_val;
+                        cvar_weight += contrib;
+                    }
+                    cum_prob = new_cum;
+                }
+                /* CVaR = conditional mean of bottom alpha fraction */
+                float cvar = (cvar_weight > 1e-8f) ? (cvar_sum / cvar_weight) : eq;
+                eq_per_action[a] = cvar;
+            }
             __syncthreads();
         }
 

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -748,6 +748,9 @@ pub struct GpuDqnTrainer {
     eval_ema_initialized: bool,
     /// Adaptive IQN lambda readiness: 0=uncertain (suppress gradient), 1=converged (full weight).
     iqn_readiness: f32,
+    /// IQN readiness — pinned device-mapped for CUDA graph. GPU reads via dev_ptr.
+    iqn_readiness_pinned: *mut f32,
+    iqn_readiness_dev_ptr: u64,
     iqn_loss_ema: f32,
     iqn_loss_initial: f32,
 
@@ -1287,6 +1290,9 @@ impl Drop for GpuDqnTrainer {
         if !self.q_divergence_pinned.is_null() {
             let _ = unsafe { cudarc::driver::result::free_host(self.q_divergence_pinned.cast()) };
         }
+        if !self.iqn_readiness_pinned.is_null() {
+            let _ = unsafe { cudarc::driver::result::free_host(self.iqn_readiness_pinned.cast()) };
+        }
     }
 }
 
@@ -1664,6 +1670,9 @@ impl GpuDqnTrainer {
             let improvement = (self.iqn_loss_initial - self.iqn_loss_ema) / self.iqn_loss_initial.max(1e-6);
             self.iqn_readiness = improvement.clamp(0.0, 1.0);
         }
+        unsafe {
+            *self.iqn_readiness_pinned = self.iqn_readiness;
+        }
     }
 
     /// Raw pointer to the pinned host readback buffer [16 × f32].
@@ -2933,6 +2942,19 @@ impl GpuDqnTrainer {
             cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, q_divergence_pinned.cast(), 0);
             dp
         };
+        // iqn_readiness — pinned device-mapped (CVaR alpha for c51_loss_kernel).
+        let iqn_readiness_pinned: *mut f32 = unsafe {
+            let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP;
+            cudarc::driver::result::malloc_host(std::mem::size_of::<f32>(), flags)
+                .map_err(|e| MLError::ModelError(format!("pinned iqn_readiness alloc: {e}")))?
+                as *mut f32
+        };
+        unsafe { *iqn_readiness_pinned = 0.0; }
+        let iqn_readiness_dev_ptr = unsafe {
+            let mut dp = 0u64;
+            cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, iqn_readiness_pinned.cast(), 0);
+            dp
+        };
         let total_actions = config.branch_0_size + config.branch_1_size + config.branch_2_size + config.branch_3_size;
         let q_out_buf = alloc_f32(&stream, b * total_actions, "q_out")?;
         let eval_td_snapshot = alloc_f32(&stream, b, "eval_td_snapshot")?;
@@ -4058,6 +4080,8 @@ impl GpuDqnTrainer {
             eval_q_std_ema: 0.0,
             eval_ema_initialized: false,
             iqn_readiness: 0.0,
+            iqn_readiness_pinned,
+            iqn_readiness_dev_ptr,
             iqn_loss_ema: 0.0,
             iqn_loss_initial: 0.0,
             lr_pinned,
@@ -6729,6 +6753,8 @@ impl GpuDqnTrainer {
                 .arg(&self.config.spectral_decoupling_lambda)
                 // ── Q-divergence accumulator (1) ──
                 .arg(&self.q_divergence_dev_ptr)
+                // ── CVaR alpha from IQN readiness ──
+                .arg(&self.iqn_readiness_dev_ptr)
                 // ── Adam step counter for stochastic Expected SARSA ──
                 .arg(&self.ptrs.t_buf)
                 .launch(LaunchConfig {