diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
index c13dd3ff5..5fb55d164 100644
--- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
@@ -1431,3 +1431,91 @@ extern "C" __global__ void exposure_aux_grad_kernel(
         }
     }
 }
+
+/* ══════════════════════════════════════════════════════════════════════
+ * POPART REWARD NORMALIZATION KERNEL (v8 D2)
+ *
+ * Normalizes rewards in-place using running mean/variance (Welford).
+ * Updates statistics from the current batch, then normalizes.
+ * Eliminates manual v_min/v_max tuning for C51.
+ *
+ * Launch: grid=(1), block=(256). Single block for atomic stat update.
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void popart_normalize_kernel(
+    float* __restrict__ rewards,
+    float* __restrict__ mean_buf,
+    float* __restrict__ var_buf,
+    float* __restrict__ count_buf,
+    int N, int warmup_count
+) {
+    __shared__ float s_sum;
+    __shared__ float s_sq_sum;
+    __shared__ float s_mu;
+    __shared__ float s_sigma;
+
+    int tid = threadIdx.x;
+
+    /* Phase 1: compute batch sum and sum-of-squares via reduction */
+    float local_sum = 0.0f;
+    float local_sq = 0.0f;
+    for (int i = tid; i < N; i += blockDim.x) {
+        float r = rewards[i];
+        local_sum += r;
+        local_sq += r * r;
+    }
+
+    /* Warp reduction */
+    for (int offset = 16; offset > 0; offset >>= 1) {
+        local_sum += __shfl_xor_sync(0xFFFFFFFF, local_sum, offset);
+        local_sq += __shfl_xor_sync(0xFFFFFFFF, local_sq, offset);
+    }
+
+    __shared__ float warp_sums[8];
+    __shared__ float warp_sqs[8];
+    int warp_id = tid / 32;
+    int lane = tid % 32;
+    if (lane == 0) { warp_sums[warp_id] = local_sum; warp_sqs[warp_id] = local_sq; }
+    __syncthreads();
+
+    if (warp_id == 0) {
+        float val = (lane < blockDim.x / 32) ? warp_sums[lane] : 0.0f;
+        float sq = (lane < blockDim.x / 32) ? warp_sqs[lane] : 0.0f;
+        for (int offset = 16; offset > 0; offset >>= 1) {
+            val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
+            sq += __shfl_xor_sync(0xFFFFFFFF, sq, offset);
+        }
+        if (lane == 0) { s_sum = val; s_sq_sum = sq; }
+    }
+    __syncthreads();
+
+    /* Phase 2: thread 0 updates running statistics */
+    if (tid == 0) {
+        float old_count = *count_buf;
+        float new_count = old_count + (float)N;
+        float batch_mean = s_sum / fmaxf((float)N, 1.0f);
+        float new_mean = (old_count > 0.0f)
+            ? (*mean_buf * old_count + s_sum) / new_count
+            : batch_mean;
+        float new_var = (new_count > 1.0f)
+            ? (s_sq_sum + *var_buf * old_count) / new_count - new_mean * new_mean
+            : 1.0f;
+        new_var = fmaxf(new_var, 0.0001f);
+        *mean_buf = new_mean;
+        *var_buf = new_var;
+        *count_buf = new_count;
+        s_mu = new_mean;
+        s_sigma = sqrtf(new_var);
+    }
+    __syncthreads();
+
+    /* Phase 3: normalize rewards in-place */
+    float mu = s_mu;
+    float sigma = fmaxf(s_sigma, 0.01f);
+    float total = *count_buf;
+
+    if (total >= (float)warmup_count) {
+        for (int i = tid; i < N; i += blockDim.x) {
+            rewards[i] = (rewards[i] - mu) / sigma;
+        }
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 496fb55b9..a0452277d 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -2015,3 +2015,80 @@ extern "C" __global__ void classify_regime_gpu(
 extern "C" __global__ void step_counter_advance(int* counter) {
     counter[0] += 1;
 }
+
+/* ══════════════════════════════════════════════════════════════════════
+ * CURRICULUM DIFFICULTY SCORING KERNEL (v8 D3)
+ *
+ * Per-bar difficulty: inverse of 5-bar directional clarity.
+ * Low score = strong trend (easy). High score = choppy (hard).
+ * Used to order episode starts from easy → hard across training.
+ *
+ * Launch: grid=(ceil(total_bars/256)), block=(256).
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void compute_difficulty_scores(
+    const __nv_bfloat16* __restrict__ targets,
+    float* __restrict__ scores,
+    int total_bars
+) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= total_bars) return;
+    if (i >= total_bars - 5) { scores[i] = 1e6f; return; }
+
+    float close_0 = __bfloat162float(targets[i * 4 + 2]);
+    float close_5 = __bfloat162float(targets[(i + 5) * 4 + 2]);
+    float ret_5 = fabsf(close_5 - close_0) / fmaxf(close_0, 1.0f);
+    scores[i] = 1.0f / (ret_5 + 0.001f);
+}
+
+/* ══════════════════════════════════════════════════════════════════════
+ * HINDSIGHT OPTIMAL EXIT RELABELING KERNEL (v8 D4)
+ *
+ * For a fraction of replay samples, replaces the reward with the
+ * optimal exit PnL within a lookahead window. Teaches the model
+ * "what was the best possible outcome for this entry?"
+ *
+ * Launch: grid=(ceil(batch_size/256)), block=(256).
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void hindsight_relabel_kernel(
+    const __nv_bfloat16* __restrict__ targets,
+    const int*   __restrict__ actions,
+    const int*   __restrict__ bar_indices,
+    float*       __restrict__ rewards,
+    int b0_size, int b1_size, int b2_size,
+    float max_position, int lookahead,
+    int total_bars, float hindsight_fraction,
+    int batch_size
+) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= batch_size) return;
+
+    /* Deterministic mask */
+    unsigned int hash = (unsigned int)(i * 48271 + bar_indices[i] * 7919);
+    hash ^= (hash >> 16); hash *= 0x45d9f3bu; hash ^= (hash >> 16);
+    if ((float)(hash & 0xFFFF) / 65535.0f >= hindsight_fraction) return;
+
+    int bar = bar_indices[i];
+    if (bar < 0 || bar + lookahead >= total_bars) return;
+
+    int exposure_idx = actions[i] / (b1_size * b2_size);
+    float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
+    float fraction = -1.0f + (float)exposure_idx * step_size;
+    float pos = fraction * max_position;
+    if (fabsf(pos) < 0.001f) return;
+
+    float entry_price = __bfloat162float(targets[bar * 4 + 2]);
+    if (entry_price < 1.0f) return;
+
+    float best_pnl = 0.0f;
+    for (int k = 1; k <= lookahead && (bar + k) < total_bars; k++) {
+        float future_price = __bfloat162float(targets[(bar + k) * 4 + 2]);
+        float pnl = pos * (future_price - entry_price);
+        if (pnl > best_pnl) best_pnl = pnl;
+    }
+
+    if (best_pnl > 0.0f) {
+        float opt_return = best_pnl / entry_price;
+        float opt_reward = asymmetric_soft_clamp(10.0f * opt_return / 0.005f);
+        rewards[i] = opt_reward;
+    }
+}