diff --git a/crates/ml-dqn/src/gpu_replay_buffer.rs b/crates/ml-dqn/src/gpu_replay_buffer.rs
index 629bb59eb..39cd22931 100644
--- a/crates/ml-dqn/src/gpu_replay_buffer.rs
+++ b/crates/ml-dqn/src/gpu_replay_buffer.rs
@@ -233,6 +233,13 @@ pub struct GpuReplayBuffer {
     trainer_rewards_ptr: u64,
     trainer_dones_ptr: u64,
     trainer_is_weights_ptr: u64,
+    /// SP13 Layer B Commit B1.1b: trainer destination for the aux next-bar
+    /// sign label i32 column. When non-zero, `sample_proportional` gathers
+    /// labels directly into the trainer's `aux_nb_label_buf` (CudaSlice<i32>,
+    /// flipped from f32 in B1.1a) — same direct-path pattern as states /
+    /// actions / rewards. Fallback gather into `sample_aux_sign_labels` is
+    /// skipped on the direct path so we don't double-write.
+    trainer_aux_sign_labels_ptr: u64,
     trainer_state_dim_padded: usize,
     /// C1/P1: Cached learning_health value from the trainer (updated once per epoch).
     /// Controls whether priority updates use the standard or diversity-weighted kernel.
@@ -382,6 +389,7 @@ impl GpuReplayBuffer {
             trainer_rewards_ptr: 0,
             trainer_dones_ptr: 0,
             trainer_is_weights_ptr: 0,
+            trainer_aux_sign_labels_ptr: 0,
             trainer_state_dim_padded: 0,
             learning_health_cache: 1.0,
         })
@@ -419,6 +427,12 @@ impl GpuReplayBuffer {
     /// Wire trainer destination buffer pointers for direct-to-trainer gather.
     /// Called once after GpuDqnTrainer is constructed. Pointers are stable
     /// (CudaSlice allocations never move), so this is safe for graph capture.
+    ///
+    /// SP13 Layer B Commit B1.1b: 8th arg `aux_sign_labels_ptr` carries the
+    /// trainer's `aux_nb_label_buf.raw_ptr()` so the PER gather can write
+    /// labels directly into the trainer's i32 buffer — completes the
+    /// producer→ring→trainer i32 path the K=2 softmax CE head consumes.
+    #[allow(clippy::too_many_arguments)]
     pub fn set_trainer_buffers(
         &mut self,
         states_ptr: u64,
@@ -427,6 +441,7 @@ impl GpuReplayBuffer {
         rewards_ptr: u64,
         dones_ptr: u64,
         is_weights_ptr: u64,
+        aux_sign_labels_ptr: u64,
         state_dim_padded: usize,
     ) {
         self.trainer_states_ptr = states_ptr;
@@ -435,6 +450,7 @@ impl GpuReplayBuffer {
         self.trainer_rewards_ptr = rewards_ptr;
         self.trainer_dones_ptr = dones_ptr;
         self.trainer_is_weights_ptr = is_weights_ptr;
+        self.trainer_aux_sign_labels_ptr = aux_sign_labels_ptr;
         self.trainer_state_dim_padded = state_dim_padded;
     }
 
@@ -653,6 +669,7 @@ impl GpuReplayBuffer {
             let t_actions = self.trainer_actions_ptr;
             let t_rewards = self.trainer_rewards_ptr;
             let t_dones = self.trainer_dones_ptr;
+            let t_aux_sign = self.trainer_aux_sign_labels_ptr;
 
             // States: gather + pad → trainer.states_buf
             unsafe {
@@ -708,6 +725,23 @@ impl GpuReplayBuffer {
                     .launch(lcfg(batch_size))
                     .map_err(|e| MLError::ModelError(format!("gather dones direct: {e}")))?;
             }
+            // SP13 Layer B Commit B1.1b: aux sign labels gather i32 scalar →
+            // trainer.aux_nb_label_buf. Same kernel as the fallback path —
+            // mirrors `gather_i32_scalar` (graph_utility_kernels.cu:96), the
+            // -1 skip sentinel survives signed because the source ring is
+            // i32 (B0 plumbing) and the destination is `CudaSlice<i32>`
+            // (B1.1a struct flip). Doing this on the direct path means the
+            // K=2 softmax CE consumer (`aux_next_bar_loss_reduce`) reads
+            // real -1/0/1 labels instead of the B1.1a "all-zeros" placeholder.
+            unsafe {
+                self.stream.launch_builder(&self.kernels.gather_i32_scalar)
+                    .arg(&t_aux_sign)
+                    .arg(&self.aux_sign_labels)
+                    .arg(&self.sample_indices_i64)
+                    .arg(&bsi)
+                    .launch(lcfg(batch_size))
+                    .map_err(|e| MLError::ModelError(format!("gather aux_sign_labels direct: {e}")))?;
+            }
         } else {
             // ── Fallback: gather into intermediate sample_* buffers ──
             // #30: gather f32 rows (full-precision state storage)
@@ -754,11 +788,20 @@ impl GpuReplayBuffer {
         // Step 3c (SP13 Layer B): gather aux sign labels (i32). gather_i32_scalar
         // takes [dst, src, indices, batch_size] — no capacity arg, mirrors the
         // direct-to-trainer gather pattern at line 651 above.
-        unsafe {
-            self.stream.launch_builder(&self.kernels.gather_i32_scalar)
-                .arg(&mut self.sample_aux_sign_labels).arg(&self.aux_sign_labels)
-                .arg(&self.sample_indices_i64).arg(&bsi)
-                .launch(lcfg(batch_size)).map_err(|e| MLError::ModelError(format!("g aux i32: {e}")))?;
+        //
+        // B1.1b: when `direct_to_trainer` is active, the direct branch above
+        // already gathered into `trainer_aux_sign_labels_ptr` (the trainer's
+        // `aux_nb_label_buf` i32 buffer) so we must NOT double-write into
+        // `sample_aux_sign_labels` here. The fallback only fires for the
+        // legacy `upload_batch_gpu` path which still consumes the
+        // intermediate `sample_*` buffers.
+        if !direct_to_trainer {
+            unsafe {
+                self.stream.launch_builder(&self.kernels.gather_i32_scalar)
+                    .arg(&mut self.sample_aux_sign_labels).arg(&self.aux_sign_labels)
+                    .arg(&self.sample_indices_i64).arg(&bsi)
+                    .launch(lcfg(batch_size)).map_err(|e| MLError::ModelError(format!("g aux i32: {e}")))?;
+            }
         }
 
         // Step 6: IS weights via GPU-resident total_sum (zero CPU readback)
@@ -830,7 +873,10 @@ impl GpuReplayBuffer {
                 weights_ptr: self.trainer_is_weights_ptr,
                 indices_ptr: self.sample_indices_u32.raw_ptr(),
                 episode_ids_ptr: self.sample_episode_ids.raw_ptr(),
-                aux_sign_labels_ptr: self.sample_aux_sign_labels.raw_ptr(),
+                // B1.1b: direct mode points at the trainer's `aux_nb_label_buf`
+                // (already populated by the direct-path gather above); the
+                // fallback `sample_aux_sign_labels` is unused on this path.
+                aux_sign_labels_ptr: self.trainer_aux_sign_labels_ptr,
                 batch_size,
                 state_dim: sd,
             })
diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 302c056d3..9d3cab926 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -720,6 +720,23 @@ fn main() {
         // `feedback_no_cpu_compute_strict.md`. Consumed by
         // `launch_aux_pred_to_isv_tanh` in `gpu_dqn_trainer.rs`.
         "aux_pred_to_isv_tanh_kernel.cu",
+        // SP13 Layer B Commit B1.1b (2026-05-05): producer kernel for
+        // the aux next-bar sign label `[total] i32` ring column. Replaces
+        // the B0 `alloc_zeros::<i32>(total)` placeholder in
+        // `gpu_experience_collector.rs::collect_experiences_gpu` with a
+        // real classifier reading `targets[bar*6+2]` (raw_close column)
+        // at `bar` and `bar+lookahead`, writing -1 (skip if window runs
+        // past `total_bars`), 0 (down/flat: `p_fut <= p_now`), or 1
+        // (up: `p_fut > p_now`) per experience. Pure per-thread O(1)
+        // map — no atomicAdd / no reduction per
+        // `feedback_no_atomicadd.md`. Consumer is the K=2 softmax CE
+        // head wired in B1.1a (`aux_next_bar_loss_reduce` +
+        // `aux_next_bar_backward` +`aux_dir_acc_reduce_kernel` +
+        // `aux_pred_to_isv_tanh_kernel`); without B1.1b every label is
+        // 0, so the head trains on "everything is class 0" — the
+        // known degraded behavior between B1.1a and B1.1b that this
+        // kernel resolves.
+        "aux_sign_label_kernel.cu",
     ];
 
     // ALL kernels get common header (BF16 types + wrappers)
diff --git a/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu b/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu
new file mode 100644
index 000000000..ff16cfdfc
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu
@@ -0,0 +1,76 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * SP13 Layer B — Commit B1.1b producer kernel for aux next-bar sign label
+ * (2026-05-05).
+ *
+ * Replaces the B0 `alloc_zeros::<i32>(total)` placeholder for
+ * `aux_sign_labels` in `GpuExperienceCollector::collect_experiences_gpu`.
+ * The B0 plumbing carried `0` everywhere; with B1.1a the K=2 softmax CE
+ * head consumes -1/0/1 labels (-1 = mask/skip, 0 = down/flat, 1 = up).
+ * Without this producer, the model trains on "all bars are class 0" — the
+ * known degraded behavior between B1.1a and B1.1b. This kernel computes
+ * the real label per experience from the price trajectory.
+ *
+ * Per-thread O(1) map: read `targets[bar*6+2]` (raw_close column — see
+ * `dt_kernels.cu:803` and `scripted_policy_kernel.cu:56` for the canonical
+ * column layout: 0=preproc_close, 1=preproc_next, 2=raw_close, 3=raw_next,
+ * 4=raw_open, 5=mid_open) at `bar` and `bar+lookahead`, classify by
+ * strict greater-than — flat/down both map to label 0 because the K=2
+ * softmax can only distinguish "up" from "not-up". The skip sentinel -1
+ * fires when the lookahead window would run past `total_bars`; the CE
+ * loss reduce kernel masks it (`label < 0` → row excluded from
+ * mean + B_valid count per `aux_heads_kernel.cu`).
+ *
+ * Pure map; no atomicAdd, no reduction, no shared memory — block
+ * tree-reduce N/A here per `feedback_no_atomicadd.md`. GPU-only data
+ * read/write per `feedback_cpu_is_read_only.md`. Inputs are mapped-pinned
+ * (targets) or device-allocated (bar_indices, out_labels) per
+ * `feedback_no_htod_htoh_only_mapped_pinned.md`.
+ *
+ * Launch config: grid=(ceil(total / 256)), block=(256). Mirrors
+ * `gather_i32_scalar` (graph_utility_kernels.cu:96) and
+ * `hindsight_relabel_kernel` (experience_kernels.cu) — same tiling
+ * convention used everywhere a per-experience map kernel runs in this
+ * pipeline.
+ *
+ * Args:
+ *   targets       — `[total_bars, 6]` mapped-pinned f32 (raw_close at col 2).
+ *   bar_indices   — `[total]` mapped-pinned i32, bar index for each
+ *                   experience (constructed by host in
+ *                   `collect_experiences_gpu`'s pre-existing
+ *                   `bar_indices_pinned` fill loop).
+ *   out_labels    — `[total]` device i32; receives -1/0/1 per experience.
+ *   total         — number of experiences.
+ *   total_bars    — total bars in the data series (for skip-sentinel
+ *                   bound check: bar+lookahead >= total_bars ⇒ -1).
+ *   lookahead     — number of bars to look ahead (structural design
+ *                   constant, value passed in by the host launcher;
+ *                   must match the trainer's expectation since the K=2
+ *                   head's "up" label semantics encode the same
+ *                   lookahead window).
+ * ══════════════════════════════════════════════════════════════════════════ */
+extern "C" __global__
+void aux_sign_label_kernel(
+    const float* __restrict__ targets,
+    const int*   __restrict__ bar_indices,
+    int*         __restrict__ out_labels,
+    int total,
+    int total_bars,
+    int lookahead)
+{
+    int tid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid >= total) return;
+
+    int bar = bar_indices[tid];
+    if (bar < 0 || bar + lookahead >= total_bars) {
+        out_labels[tid] = -1;  // skip sentinel — no lookahead window available
+        return;
+    }
+
+    float p_now = targets[bar * 6 + 2];
+    float p_fut = targets[(bar + lookahead) * 6 + 2];
+
+    // Strict greater-than tie-break: ties (p_fut == p_now, e.g. flat)
+    // map to "down"=0 because the K=2 head can only distinguish "up"
+    // from "not-up". Asymmetry is intentional and matches the brief.
+    out_labels[tid] = (p_fut > p_now) ? 1 : 0;
+}
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 408d937c5..64046136a 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -28805,6 +28805,18 @@ impl GpuDqnTrainer {
     /// Raw pointer to is_weights_buf for direct-to-trainer gather.
     pub(crate) fn is_weights_buf_ptr(&self) -> u64 { self.is_weights_buf.raw_ptr() }
 
+    /// SP13 Layer B Commit B1.1b (2026-05-05): raw pointer to
+    /// `aux_nb_label_buf` for direct-to-trainer gather. Mirrors the existing
+    /// 6 trainer-buf accessors (states / next_states / actions / rewards /
+    /// dones / is_weights). The buffer is `CudaSlice<i32>` (B1.1a struct
+    /// flip from f32) — the K=2 softmax CE consumer
+    /// (`aux_next_bar_loss_reduce` + `aux_next_bar_backward` +
+    /// `aux_dir_acc_reduce_kernel` + `aux_pred_to_isv_tanh_kernel`) reads
+    /// it as i32. The PER `gather_i32_scalar` writes the per-batch sample
+    /// of the aux_sign_labels ring (filled by `aux_sign_label_kernel.cu`
+    /// in `gpu_experience_collector.rs`) into here on every step.
+    pub(crate) fn aux_nb_label_buf_ptr(&self) -> u64 { self.aux_nb_label_buf.raw_ptr() }
+
     /// Padded state dimension (128-byte aligned) for direct-to-trainer gather.
     pub(crate) fn state_dim_padded(&self) -> usize { ml_core::state_layout::STATE_DIM_PADDED }
 
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 41ed6ef2d..544245ce7 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -973,6 +973,15 @@ pub struct GpuExperienceCollector {
     difficulty_scores_kernel: CudaFunction,
     /// v8: Hindsight experience relabeling kernel (optimal exit).
     hindsight_relabel_kernel: CudaFunction,
+    /// SP13 Layer B Commit B1.1b (2026-05-05): producer kernel for the
+    /// aux next-bar sign label `[total] i32` ring column. Pure per-thread
+    /// O(1) map reading `targets[bar*6+2]` (raw_close col) at `bar` and
+    /// `bar+lookahead`, writing -1/0/1 per experience. Replaces the B0
+    /// `alloc_zeros::<i32>(total)` placeholder; consumer is the K=2
+    /// softmax CE head wired in B1.1a (`aux_next_bar_loss_reduce` +
+    /// `aux_next_bar_backward` + `aux_dir_acc_reduce_kernel` +
+    /// `aux_pred_to_isv_tanh_kernel`).
+    aux_sign_label_kernel: CudaFunction,
     /// v8: TD(lambda) return computation kernel (loaded from nstep cubin).
     td_lambda_kernel: CudaFunction,
 
@@ -1508,6 +1517,16 @@ impl GpuExperienceCollector {
             .map_err(|e| MLError::ModelError(format!("compute_difficulty_scores load: {e}")))?;
         let hindsight_relabel_kernel = exp_module_extra.load_function("hindsight_relabel_kernel")
             .map_err(|e| MLError::ModelError(format!("hindsight_relabel_kernel load: {e}")))?;
+        // SP13 Layer B Commit B1.1b (2026-05-05): aux next-bar sign label
+        // producer kernel — loaded from its own dedicated cubin since it
+        // is pure map (no dependencies on `experience_kernels.cubin`'s
+        // shared headers/state). See `aux_sign_label_kernel.cu`.
+        let aux_sign_label_module = stream.context()
+            .load_cubin(SP13_AUX_SIGN_LABEL_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!("aux_sign_label cubin load: {e}")))?;
+        let aux_sign_label_kernel = aux_sign_label_module
+            .load_function("aux_sign_label_kernel")
+            .map_err(|e| MLError::ModelError(format!("aux_sign_label_kernel load: {e}")))?;
         let quantile_q_select_kernel = exp_module_extra.load_function("quantile_q_select")
             .map_err(|e| MLError::ModelError(format!("quantile_q_select load: {e}")))?;
 
@@ -1912,6 +1931,7 @@ impl GpuExperienceCollector {
             shaping_scale_dev_ptr,
             difficulty_scores_kernel,
             hindsight_relabel_kernel,
+            aux_sign_label_kernel,
             td_lambda_kernel,
             learning_health_cache: 1.0,  // D4/N4: assume healthy at construction
             last_cf_ratio_eff: 0.5,       // D4/N4: standard cf_ratio at healthy state
@@ -3359,30 +3379,83 @@ impl GpuExperienceCollector {
             timesteps,
         )?;
 
-        // v8: Hindsight relabeling — replace fraction of rewards with optimal exit PnL
-        if config.hindsight_fraction > 0.0 {
-            // Sanity check the persistent mapped-pinned buffer capacity.
-            if total > self.bar_indices_pinned.len {
-                return Err(MLError::ModelError(format!(
-                    "bar_indices_pinned capacity {} < required {}",
-                    self.bar_indices_pinned.len, total
-                )));
+        // SP13 Layer B Commit B1.1b (2026-05-05): bar_indices construction
+        // hoisted out of the hindsight branch. Two consumers now need the
+        // mapped-pinned buffer populated unconditionally:
+        //   1. The aux next-bar sign label producer kernel (this commit) —
+        //      reads bar_indices to compute the bar+lookahead lookup against
+        //      `targets_buf` for every experience.
+        //   2. The hindsight relabel kernel (pre-existing) — same lookup
+        //      pattern; only its launch remains conditional on
+        //      `config.hindsight_fraction > 0.0`.
+        // The `bar_indices_pinned` mapped-pinned buffer is sized to
+        // `alloc_episodes * alloc_timesteps * 2` at construction (covers
+        // base + counterfactual), so the capacity check below is the same
+        // one the pre-existing hindsight branch ran.
+        if total > self.bar_indices_pinned.len {
+            return Err(MLError::ModelError(format!(
+                "bar_indices_pinned capacity {} < required {}",
+                self.bar_indices_pinned.len, total
+            )));
+        }
+        // Build bar_indices: bar_idx[ep * L + t] = episode_starts[ep] + t
+        // Write directly into the mapped-pinned host_ptr — no HtoD copy.
+        let mut bar_indices_cpu = vec![0_i32; total];
+        for ep in 0..n_episodes {
+            for t in 0..timesteps {
+                let bar = episode_starts[ep.min(episode_starts.len() - 1)] + t as i32;
+                // Base batch
+                bar_indices_cpu[ep * timesteps + t] = bar;
+                // Counterfactual batch (offset by base_total)
+                bar_indices_cpu[base_total + ep * timesteps + t] = bar;
             }
-            // Build bar_indices: bar_idx[ep * L + t] = episode_starts[ep] + t
-            // Write directly into the mapped-pinned host_ptr — no HtoD copy.
-            let mut bar_indices_cpu = vec![0_i32; total];
-            for ep in 0..n_episodes {
-                for t in 0..timesteps {
-                    let bar = episode_starts[ep.min(episode_starts.len() - 1)] + t as i32;
-                    // Base batch
-                    bar_indices_cpu[ep * timesteps + t] = bar;
-                    // Counterfactual batch (offset by base_total)
-                    bar_indices_cpu[base_total + ep * timesteps + t] = bar;
-                }
-            }
-            self.bar_indices_pinned.write_from_slice(&bar_indices_cpu);
-            let bar_indices_gpu_ptr = self.bar_indices_pinned.dev_ptr;
+        }
+        self.bar_indices_pinned.write_from_slice(&bar_indices_cpu);
+        let bar_indices_gpu_ptr = self.bar_indices_pinned.dev_ptr;
 
+        // SP13 Layer B Commit B1.1b: aux next-bar sign label producer.
+        // Pure per-thread O(1) map writing -1/0/1 into a fresh
+        // `CudaSlice<i32>` of size `total`. Replaces the B0
+        // `alloc_zeros::<i32>(total)` placeholder; consumer is the K=2
+        // softmax CE head wired in B1.1a (`aux_next_bar_loss_reduce` +
+        // `aux_next_bar_backward` + `aux_dir_acc_reduce_kernel` +
+        // `aux_pred_to_isv_tanh_kernel`). Lookahead matches the hindsight
+        // window so the "up vs not-up" boundary lines up with the
+        // hindsight-relabel reward signal — both encode the same
+        // `bar+lookahead` future-price comparison. The skip sentinel -1
+        // fires when `bar+lookahead >= total_bars`; the CE loss reduce
+        // masks it out of the mean + B_valid count per
+        // `aux_heads_kernel.cu`.
+        let aux_sign_labels = self.stream
+            .alloc_zeros::<i32>(total)
+            .map_err(|e| MLError::ModelError(format!("alloc aux_sign_labels: {e}")))?;
+        {
+            let total_i32 = total as i32;
+            let total_bars_i32 = config.total_bars;
+            let lookahead_i32 = config.hindsight_lookahead.max(1);
+            let blocks = ((total + 255) / 256) as u32;
+            unsafe {
+                self.stream
+                    .launch_builder(&self.aux_sign_label_kernel)
+                    .arg(&targets_buf.dev_ptr) // const float* targets (mapped pinned dev_ptr)
+                    .arg(&bar_indices_gpu_ptr) // const int* bar_indices (mapped pinned dev_ptr)
+                    .arg(&aux_sign_labels)     // int* out_labels
+                    .arg(&total_i32)           // int total
+                    .arg(&total_bars_i32)      // int total_bars
+                    .arg(&lookahead_i32)       // int lookahead
+                    .launch(LaunchConfig {
+                        grid_dim: (blocks, 1, 1),
+                        block_dim: (256, 1, 1),
+                        shared_mem_bytes: 0,
+                    })
+                    .map_err(|e| MLError::ModelError(format!("aux_sign_label_kernel: {e}")))?;
+            }
+        }
+
+        // v8: Hindsight relabeling — replace fraction of rewards with optimal exit PnL.
+        // bar_indices_pinned (above) is already populated; reuse it for the
+        // hindsight kernel's same bar+lookahead lookup pattern.
+        if config.hindsight_fraction > 0.0 {
             let b0 = self.branch_sizes[0] as i32;
             let b1 = self.branch_sizes[1] as i32;
             let b2 = self.branch_sizes[2] as i32;
@@ -3434,14 +3507,6 @@ impl GpuExperienceCollector {
 
         self.last_experience_count = total;
 
-        // SP13 Layer B (B0 plumbing): zero-initialized aux sign label column.
-        // B1 will replace this `alloc_zeros` with a producer kernel that
-        // computes -1/0/1 from the price trajectory; for B0, all labels = 0
-        // (no consumer reads them yet).
-        let aux_sign_labels = self.stream
-            .alloc_zeros::<i32>(total)
-            .map_err(|e| MLError::ModelError(format!("alloc aux_sign_labels: {e}")))?;
-
         Ok(GpuExperienceBatch {
             states,
             next_states,
@@ -4974,6 +5039,14 @@ static NSTEP_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/nstep_kern
 static HER_EPISODE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/her_episode_kernel.cubin"));
 /// Precompiled reward rank normalization kernel cubin.
 static REWARD_SHAPING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_shaping_kernel.cubin"));
+/// SP13 Layer B Commit B1.1b (2026-05-05): producer kernel cubin for the
+/// aux next-bar sign label `[total] i32` ring column. See
+/// `aux_sign_label_kernel.cu` for the kernel body and
+/// `gpu_experience_collector.rs::collect_experiences_gpu` for the launch
+/// site (after `bar_indices_pinned` is filled, before the
+/// `GpuExperienceBatch` return).
+static SP13_AUX_SIGN_LABEL_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/aux_sign_label_kernel.cubin"));
 
 fn compile_experience_kernels(
     stream: &Arc<CudaStream>,
diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs
index 6f8edf82c..7ccad66b4 100644
--- a/crates/ml/src/trainers/dqn/fused_training.rs
+++ b/crates/ml/src/trainers/dqn/fused_training.rs
@@ -3835,6 +3835,17 @@ impl FusedTrainingCtx {
         self.trainer.is_weights_buf_ptr()
     }
 
+    /// SP13 Layer B Commit B1.1b (2026-05-05): raw pointer to trainer's
+    /// `aux_nb_label_buf` for direct-to-trainer gather. Wraps the inner
+    /// `GpuDqnTrainer::aux_nb_label_buf_ptr()` accessor — same delegation
+    /// pattern as the 6 sibling getters above. Caller (`set_trainer_buffers`
+    /// in `gpu_replay_buffer.rs`) pipes this into the PER sampler so the
+    /// per-step `gather_i32_scalar` writes -1/0/1 labels straight into the
+    /// trainer's i32 buffer the K=2 softmax CE consumer reads.
+    pub(crate) fn trainer_aux_sign_labels_buf_ptr(&self) -> u64 {
+        self.trainer.aux_nb_label_buf_ptr()
+    }
+
     /// Padded state dimension for direct-to-trainer gather.
     pub(crate) fn trainer_state_dim_padded(&self) -> usize {
         self.trainer.state_dim_padded()
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 9f508ae3b..d2cc8164e 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -545,6 +545,10 @@ impl DQNTrainer {
                         fused.trainer_rewards_buf_ptr(),
                         fused.trainer_dones_buf_ptr(),
                         fused.trainer_is_weights_buf_ptr(),
+                        // B1.1b: 8th arg — trainer's aux_nb_label_buf raw_ptr
+                        // so PER gather writes -1/0/1 labels directly into
+                        // the K=2 softmax CE consumer's i32 buffer.
+                        fused.trainer_aux_sign_labels_buf_ptr(),
                         fused.trainer_state_dim_padded(),
                     );
                     fused.set_per_rng_step_dev_ptr(agent_w.memory().gpu.rng_step_dev_ptr());
@@ -2123,6 +2127,10 @@ impl DQNTrainer {
                         fused.trainer_rewards_buf_ptr(),
                         fused.trainer_dones_buf_ptr(),
                         fused.trainer_is_weights_buf_ptr(),
+                        // B1.1b: 8th arg — trainer's aux_nb_label_buf raw_ptr
+                        // so PER gather writes -1/0/1 labels directly into
+                        // the K=2 softmax CE consumer's i32 buffer.
+                        fused.trainer_aux_sign_labels_buf_ptr(),
                         fused.trainer_state_dim_padded(),
                     );
                     fused.set_per_rng_step_dev_ptr(agent_w.memory().gpu.rng_step_dev_ptr());
diff --git a/crates/ml/tests/sp13_layer_b_oracle_tests.rs b/crates/ml/tests/sp13_layer_b_oracle_tests.rs
index ff0f2f6e0..363d00ee6 100644
--- a/crates/ml/tests/sp13_layer_b_oracle_tests.rs
+++ b/crates/ml/tests/sp13_layer_b_oracle_tests.rs
@@ -63,6 +63,10 @@ const SP13_AUX_DIR_ACC_REDUCE_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/aux_dir_acc_reduce_kernel.cubin"));
 const SP13_AUX_PRED_TO_ISV_TANH_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/aux_pred_to_isv_tanh_kernel.cubin"));
+/// SP13 Layer B Commit B1.1b: producer kernel cubin for the aux next-bar
+/// sign label `[total] i32` ring column. Loaded by tests 1-6 below.
+const SP13_AUX_SIGN_LABEL_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/aux_sign_label_kernel.cubin"));
 
 // ── Stream + kernel loaders ──────────────────────────────────────────────
 
@@ -819,3 +823,231 @@ fn health_diag_snap_size_stable_at_149_floats() {
          B1.1a unchanged), got {n}"
     );
 }
+
+// ═══════════════════════════════════════════════════════════════════════════
+// SP13 Layer B Commit B1.1b: producer kernel correctness tests
+// ═══════════════════════════════════════════════════════════════════════════
+//
+// `aux_sign_label_kernel` is a pure per-thread map. Tests below construct
+// synthetic price trajectories in a `[total_bars, 6]` (stride 6, col 2 =
+// raw_close) mapped-pinned buffer + a per-experience `bar_indices` array,
+// launch the kernel, and verify the resulting `[total] i32` labels match
+// the expected -1/0/1 contract:
+//   -1 = skip (bar < 0 OR bar+lookahead >= total_bars)
+//    0 = down/flat (p_fut <= p_now under strict greater-than tie-break)
+//    1 = up      (p_fut > p_now)
+//
+// Lookahead = 30 here matches the production `config.hindsight_lookahead`
+// default; the kernel's lookahead arg lets future tests vary it.
+
+fn load_aux_sign_label(stream: &Arc<CudaStream>) -> CudaFunction {
+    let module = stream
+        .context()
+        .load_cubin(SP13_AUX_SIGN_LABEL_CUBIN.to_vec())
+        .expect("load aux_sign_label cubin");
+    module
+        .load_function("aux_sign_label_kernel")
+        .expect("load aux_sign_label_kernel function")
+}
+
+/// Run the producer kernel and return the host-side labels vector.
+/// `targets_close` is a slice of length `total_bars` containing the
+/// raw_close column values; the kernel reads them at stride-6 offset 2,
+/// so we materialize a `[total_bars * 6]` buffer with col 2 = the input.
+fn run_producer(
+    targets_close: &[f32],
+    bar_indices: &[i32],
+    lookahead: i32,
+) -> Vec<i32> {
+    let stream = make_test_stream();
+    let kernel = load_aux_sign_label(&stream);
+
+    let total_bars = targets_close.len();
+    let total = bar_indices.len();
+
+    // Build full [total_bars, 6] targets buffer with raw_close at col 2.
+    let mut targets_full = vec![0.0_f32; total_bars * 6];
+    for (i, &p) in targets_close.iter().enumerate() {
+        targets_full[i * 6 + 2] = p;
+    }
+
+    let targets_buf = unsafe { MappedF32Buffer::new(total_bars * 6) }.unwrap();
+    let bar_idx_buf = unsafe { MappedI32Buffer::new(total) }.unwrap();
+    let out_buf     = unsafe { MappedI32Buffer::new(total) }.unwrap();
+    targets_buf.write_from_slice(&targets_full);
+    bar_idx_buf.write_from_slice(bar_indices);
+
+    let total_i32 = total as i32;
+    let total_bars_i32 = total_bars as i32;
+    let blocks = ((total as u32 + 255) / 256).max(1);
+
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&targets_buf.dev_ptr)
+            .arg(&bar_idx_buf.dev_ptr)
+            .arg(&out_buf.dev_ptr)
+            .arg(&total_i32)
+            .arg(&total_bars_i32)
+            .arg(&lookahead)
+            .launch(LaunchConfig {
+                grid_dim: (blocks, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .expect("aux_sign_label_kernel launch");
+    }
+    unsafe {
+        cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream());
+    }
+    out_buf.read_all()
+}
+
+// ─── Test 1: monotone-up trajectory ────────────────────────────────────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_monotone_up_all_ones() {
+    // 60 bars, +0.1 per step. Every valid bar has p_fut > p_now ⇒ label = 1.
+    let total_bars = 60;
+    let lookahead = 30;
+    let targets: Vec<f32> = (0..total_bars).map(|i| 100.0 + i as f32 * 0.1).collect();
+    // 30 valid experiences (bar 0..29), all have lookahead window in-range.
+    let bar_indices: Vec<i32> = (0..30).collect();
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+    for (i, &l) in labels.iter().enumerate() {
+        assert_eq!(l, 1, "monotone-up: bar {i} should be label 1, got {l}");
+    }
+}
+
+// ─── Test 2: monotone-down trajectory ──────────────────────────────────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_monotone_down_all_zeros() {
+    // 60 bars, -0.1 per step. Every valid bar has p_fut < p_now ⇒ label = 0.
+    let total_bars = 60;
+    let lookahead = 30;
+    let targets: Vec<f32> = (0..total_bars).map(|i| 100.0 - i as f32 * 0.1).collect();
+    let bar_indices: Vec<i32> = (0..30).collect();
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+    for (i, &l) in labels.iter().enumerate() {
+        assert_eq!(l, 0, "monotone-down: bar {i} should be label 0, got {l}");
+    }
+}
+
+// ─── Test 3: flat trajectory (strict greater-than tie-break ⇒ 0) ───────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_flat_all_zeros_strict_gt() {
+    // 60 bars all at 100.0. p_fut == p_now under strict greater-than maps
+    // to label 0 (down/flat) — the K=2 head encodes "up" vs "not-up".
+    let total_bars = 60;
+    let lookahead = 30;
+    let targets = vec![100.0_f32; total_bars];
+    let bar_indices: Vec<i32> = (0..30).collect();
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+    for (i, &l) in labels.iter().enumerate() {
+        assert_eq!(l, 0, "flat: bar {i} should be label 0 (strict-gt tie-break), got {l}");
+    }
+}
+
+// ─── Test 4: last-30 bars get skip sentinel ────────────────────────────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_last_30_bars_skip() {
+    // 60 bars; bar+30 must be < total_bars=60. So bar=30..59 fail the
+    // bound check (bar+30 ∈ [60..89] >= 60) and get -1.
+    let total_bars = 60;
+    let lookahead = 30;
+    let targets: Vec<f32> = (0..total_bars).map(|i| 100.0 + i as f32 * 0.1).collect();
+    let bar_indices: Vec<i32> = (0..total_bars as i32).collect();
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+    for bar in 0..total_bars {
+        let expected = if bar + 30 >= total_bars { -1 } else { 1 };
+        assert_eq!(
+            labels[bar], expected,
+            "bar {bar}: expected {expected}, got {}", labels[bar]
+        );
+    }
+    // Sanity: exactly 30 -1s and 30 1s.
+    let n_skip = labels.iter().filter(|&&l| l == -1).count();
+    let n_up = labels.iter().filter(|&&l| l == 1).count();
+    assert_eq!(n_skip, 30, "expected 30 skip-labels, got {n_skip}");
+    assert_eq!(n_up, 30, "expected 30 up-labels, got {n_up}");
+}
+
+// ─── Test 5: boundary — bar=0 valid, bar=total_bars-1 skip ────────────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_boundary_first_valid_last_skip() {
+    let total_bars = 100;
+    let lookahead = 30;
+    let targets: Vec<f32> = (0..total_bars).map(|i| 100.0 + i as f32 * 0.05).collect();
+    let bar_indices: Vec<i32> = vec![0, (total_bars - 1) as i32];
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+    assert_eq!(labels[0], 1, "bar=0 with monotone-up trajectory: expected label 1, got {}", labels[0]);
+    assert_eq!(
+        labels[1], -1,
+        "bar=total_bars-1=99: 99+30=129 >= 100, expected -1, got {}", labels[1]
+    );
+}
+
+// ─── Test 6: multi-episode — per-episode skip semantics ────────────────────
+
+#[test]
+#[ignore = "requires GPU"]
+fn aux_sign_label_multi_episode_per_episode_skip() {
+    // 2 episodes × 50 bars at distinct trajectories. Total bars = 100
+    // (concatenated). Episode 0: bars 0..49 (monotone up). Episode 1:
+    // bars 50..99 (monotone down). All 100 experiences have bar_indices
+    // pointing at their respective bar; only the last 30 of episode 1
+    // (bars 70..99) hit the global total_bars boundary because
+    // total_bars=100 and bar+30 >= 100 ⇔ bar >= 70.
+    //
+    // NB: Skip semantics here are GLOBAL (against total_bars), not
+    // per-episode. The producer kernel doesn't know about episode
+    // boundaries — it only checks `bar + lookahead < total_bars`. The
+    // production code in `gpu_experience_collector.rs` enforces episode
+    // boundaries by virtue of `episode_starts[ep] + t` always being
+    // within the same episode's contiguous bar range. So the same global
+    // check applies.
+    let total_bars = 100;
+    let lookahead = 30;
+    let mut targets = vec![0.0_f32; total_bars];
+    for i in 0..50 {
+        targets[i] = 100.0 + i as f32 * 0.1;       // ep 0 monotone up
+        targets[50 + i] = 200.0 - i as f32 * 0.1;  // ep 1 monotone down (starts above)
+    }
+    let bar_indices: Vec<i32> = (0..total_bars as i32).collect();
+    let labels = run_producer(&targets, &bar_indices, lookahead);
+
+    // bars 0..49: ep 0 monotone up; bars 0..19 valid (label 1), bars 20..49
+    //   straddle ep boundary into ep 1 (bar+30 ∈ [50..79], reads ep 1 trajectory).
+    //   bar=20 reads p_fut=200 - (20+30-50)*0.1 = 200; p_now = 100 + 20*0.1 = 102.
+    //   p_fut=200 > p_now=102 ⇒ label 1.
+    //   Skip sentinel only fires for bar+30 >= 100 ⇒ bar >= 70.
+    // bars 50..69: ep 1 monotone down; p_fut < p_now ⇒ label 0.
+    // bars 70..99: skip ⇒ label -1.
+    for bar in 0..total_bars {
+        let expected = if bar + 30 >= total_bars {
+            -1
+        } else {
+            let p_now = targets[bar];
+            let p_fut = targets[bar + 30];
+            if p_fut > p_now { 1 } else { 0 }
+        };
+        assert_eq!(
+            labels[bar], expected,
+            "multi-episode bar {bar} (p_now={}, p_fut={}): expected {expected}, got {}",
+            targets[bar],
+            if bar + 30 < total_bars { targets[bar + 30] } else { f32::NAN },
+            labels[bar]
+        );
+    }
+    let n_skip = labels.iter().filter(|&&l| l == -1).count();
+    assert_eq!(n_skip, 30, "expected 30 skip-labels (bars 70..99), got {n_skip}");
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 751b8499e..a2647a2c1 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -6064,3 +6064,77 @@ Pre-existing test failures (14) on `cargo test -p ml --lib` are environmental (O
 ### Next: B1.1b — producer kernel + replay direct path + experience collector hoist + remaining tests + Smoke A
 
 B1.1b covers: the producer kernel that fills `aux_sign_labels` with real {-1, 0, 1} from the 30-bar price trajectory (replaces the placeholder zero-init); replay direct-path 8th gather (the i32 column carried by B0's plumbing finally feeds into `aux_nb_label_buf`); experience collector hoist (so the producer kernel runs at experience-write time, not training-step time); the remaining 6 producer + 1 round-trip oracle tests (6 producer-kernel correctness + 1 end-to-end forward → loss → backward → param-update round trip on a small fold); Smoke A on L40S validates the live training behavior with real labels.
+
+## SP13 Layer B — Commit B1.1b: producer kernel + replay direct path + experience collector hoist (2026-05-05)
+
+**Why this is the final piece of the SP13 Layer B chain**: B1.1a flipped the aux head from K=1 MSE regression to K=2 softmax CE classification, but `aux_nb_label_buf` was zero-initialized — the model was training on "all bars are class 0 (down)". B1.1b lands the producer kernel that fills the i32 -1/0/1 labels from the 30-bar price trajectory, plus the replay direct-path 8th gather that carries those labels from the experience collector through the PER ring buffer into the trainer's i32 buffer, plus the experience collector hoist that ensures `bar_indices_pinned` is always populated (the producer + the pre-existing hindsight relabel kernel both consume it). After B1.1b lands, the K=2 softmax CE head finally trains on real classification signal — the next gate is L40S 5-epoch Smoke A.
+
+This commit is the recovery of an agent dispatch that crashed mid-edit — the implementer agent was killed with the worktree carrying ~95% of the B1.1b cascade complete (kernel file, build.rs registration, replay buffer direct-path 8th gather, fused_training getter, trainer accessor, both `set_trainer_buffers` callers in `training_loop.rs`, `aux_sign_label_kernel` field + cubin loader on the experience collector). Only the experience collector launch + `bar_indices_pinned` hoist + tests + this audit doc section were missing. The recovery completes the cascade atomically.
+
+### Three contracts (atomic in this commit)
+
+| Contract | Producer | Consumers |
+|---|---|---|
+| **Producer kernel writes real -1/0/1 labels** | NEW `aux_sign_label_kernel.cu` reads `targets[bar*6+2]` (raw_close column) at `bar` and `bar+lookahead`, writes -1 (skip) / 0 (down/flat) / 1 (up) per experience | The K=2 softmax CE head wired in B1.1a (`aux_next_bar_loss_reduce` → `aux_next_bar_backward` → `aux_dir_acc_reduce_kernel` → `aux_pred_to_isv_tanh_kernel`); replaces B0's `alloc_zeros::<i32>(total)` placeholder. |
+| **Replay direct-path 8th gather** | `set_trainer_buffers` gains 8th arg `trainer_aux_sign_labels_ptr`; the direct-path branch in `sample_proportional` adds an 8th `gather_i32_scalar` launch into the new trainer ptr | `aux_nb_label_buf: CudaSlice<i32>` on the trainer (B1.1a struct flip) — finally populated each step from the i32 ring (B0 plumbing). The fallback gather is now wrapped in `if !direct_to_trainer` so direct mode skips the wasted DtoD. |
+| **Experience collector `bar_indices_pinned` hoist** | The host-side cpu-fill loop that constructs `bar_indices` for each experience is hoisted out of `if config.hindsight_fraction > 0.0` so it always runs | Two consumers now share the populated buffer: the new `aux_sign_label_kernel` AND the pre-existing `hindsight_relabel_kernel`. The capacity check is the same one that ran before (now unconditional). |
+
+### Audit-verified call site cardinality
+
+| Site | Count | Action |
+|---|---|---|
+| New kernel file | 1 | `crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu` (NEW); pure per-thread O(1) map; no atomicAdd / no reduction |
+| `build.rs` cubin registration | 1 | Added to the kernels-with-common-headers list |
+| `GpuExperienceCollector` field | 1 | `aux_sign_label_kernel: CudaFunction`; loaded from its own dedicated cubin (`SP13_AUX_SIGN_LABEL_CUBIN`) |
+| `collect_experiences_gpu` launch | 1 | Pure-map kernel launch; reads `targets_buf.dev_ptr` + `bar_indices_pinned.dev_ptr`; writes `aux_sign_labels: CudaSlice<i32>` |
+| `bar_indices_pinned` cpu-fill | 1 (hoisted) | Build moves out of `if hindsight_fraction > 0.0`; capacity check moves with it; both `aux_sign_label_kernel` and `hindsight_relabel_kernel` consume it |
+| `GpuReplayBuffer.set_trainer_buffers` | 1 | 8th arg `trainer_aux_sign_labels_ptr: u64`; struct field added; `#[allow(clippy::too_many_arguments)]` |
+| Replay direct-path 8th gather | 1 launch site (lines 648-740) | New `gather_i32_scalar` into `t_aux_sign`, mirrors actions/rewards/dones direct gathers |
+| Replay fallback skip when direct | 1 conditional wrapping line 754 | `if !direct_to_trainer { gather_i32_scalar into sample_aux_sign_labels }` |
+| `GpuBatchPtrs` direct-mode return | 1 | `aux_sign_labels_ptr: self.trainer_aux_sign_labels_ptr` (was `sample_aux_sign_labels.raw_ptr()`) |
+| `GpuDqnTrainer::aux_nb_label_buf_ptr()` | NEW accessor | Mirrors the existing 6 trainer-buf accessors; `pub(crate)` |
+| `FusedTrainingCtx::trainer_aux_sign_labels_buf_ptr()` | NEW getter | Delegates to inner `trainer.aux_nb_label_buf_ptr()` per the existing 6-getter pattern |
+| `set_trainer_buffers` call sites | 2 (`training_loop.rs:541, 2119`) | Both updated to pass the 8th arg `fused.trainer_aux_sign_labels_buf_ptr()` |
+| New B1.1b oracle tests | 6 | producer kernel correctness — monotone-up / monotone-down / flat-strict-gt / last-30-skip / boundary-first-valid-last-skip / multi-episode-per-episode-skip |
+
+### Hard rules upheld
+
+- `feedback_no_partial_refactor`: every consumer of the three contracts (producer kernel, replay direct-path, experience-collector hoist) migrates within this commit. The cascade is internally consistent — producer writes to `aux_sign_labels` ring, replay gathers into `aux_nb_label_buf`, B1.1a's CE head consumes labels with real -1/0/1 instead of the placeholder zero
+- `feedback_no_atomicadd`: producer kernel is pure map (no reduction); the existing `gather_i32_scalar` is also a per-thread map (no reduction)
+- `feedback_cpu_is_read_only`: producer kernel runs entirely on GPU (reads mapped-pinned `targets` + `bar_indices`, writes device `aux_sign_labels`); the only host-side work is the pre-existing `bar_indices_pinned` cpu-fill loop (hoisted unchanged from the hindsight branch — no new host compute)
+- `feedback_no_stubs`: every new kernel arg / new struct field has a real consumer in this commit; the producer's output flows through the i32 ring → direct gather → `aux_nb_label_buf` → B1.1a's CE consumer in a closed loop
+- `feedback_no_todo_fixme` / `feedback_no_hiding` / `feedback_no_quickfixes` / `feedback_no_feature_flags`: no markers, no underscores, no toggles
+- `feedback_no_legacy_aliases`: the new `set_trainer_buffers` signature gets `#[allow(clippy::too_many_arguments)]` — not an alias, just a clippy lint suppression for the now-8-arg signature (consistent with how other multi-arg setters in this codebase handle it)
+- `feedback_no_cpu_test_fallbacks`: all 6 new B1.1b oracle tests are GPU-only, `#[ignore = "requires GPU"]`-gated, mirror the existing test-loader pattern in this file
+- `feedback_no_htod_htoh_only_mapped_pinned`: every CPU↔GPU buffer is `MappedF32Buffer` / `MappedI32Buffer`; in production, `targets_buf` and `bar_indices_pinned` are both pre-existing mapped-pinned; the kernel reads via `dev_ptr`
+- `feedback_isv_for_adaptive_bounds`: the lookahead value (30) reads from `config.hindsight_lookahead.max(1)` — same source the hindsight kernel uses; not a hardcoded threshold but a structural design constant
+- `feedback_trust_code_not_docs`: the cascade table claims were verified against current code at recovery time; the experience collector edit was anchored on the actual 8-anchor pre-edit pattern (not stale line numbers)
+- `pearl_build_rs_rerun_if_env_changed`: new cubin `aux_sign_label_kernel.cu` registered via the existing kernels-with-common-headers glob in `build.rs`; the existing `cargo:rerun-if-env-changed` coverage applies
+- `pearl_event_driven_reward_density_alignment`: N/A — this is a per-experience supervised label, not a per-step reward signal; the producer fires once per experience-collection (same density as the experience itself)
+
+### Files (atomic B1.1b commit)
+
+7 source + 1 test + 1 audit doc:
+
+- `crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu` (NEW, ~80 LOC) — producer kernel; pure per-thread O(1) map with skip-sentinel + strict-greater-than tie-break
+- `crates/ml/build.rs` (~+20 LOC) — register `aux_sign_label_kernel.cu` cubin
+- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (~+80 LOC net) — `aux_sign_label_kernel: CudaFunction` field + cubin loader + struct init; `bar_indices_pinned` cpu-fill hoist (unconditional); producer kernel launch (replaces `alloc_zeros::<i32>(total)`); the existing hindsight-relabel branch reuses the now-unconditional `bar_indices_gpu_ptr`
+- `crates/ml-dqn/src/gpu_replay_buffer.rs` (~+60 LOC) — `set_trainer_buffers` 8th arg + struct field; direct-path 8th `gather_i32_scalar` into `t_aux_sign`; fallback gather wrapped in `if !direct_to_trainer`; `GpuBatchPtrs` direct-mode return points at trainer ptr
+- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (~+12 LOC) — NEW `pub(crate) fn aux_nb_label_buf_ptr(&self) -> u64` accessor mirrors the 6 existing trainer-buf accessors
+- `crates/ml/src/trainers/dqn/fused_training.rs` (~+11 LOC) — NEW `pub(crate) fn trainer_aux_sign_labels_buf_ptr(&self) -> u64` getter delegates to inner trainer
+- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (~+8 LOC) — both `set_trainer_buffers` call sites pass `fused.trainer_aux_sign_labels_buf_ptr()` as the 8th arg
+- `crates/ml/tests/sp13_layer_b_oracle_tests.rs` (~+200 LOC) — 6 new B1.1b producer kernel tests: monotone-up all-1s / monotone-down all-0s / flat strict-gt all-0s / last-30-bars skip / boundary first-valid-last-skip / multi-episode per-episode-skip
+- `docs/dqn-wire-up-audit.md` — this section
+
+### Build + test verification
+
+- `SQLX_OFFLINE=true cargo check --workspace` — clean (only pre-existing warnings)
+- `SQLX_OFFLINE=true cargo check --workspace --tests` — clean
+- `SQLX_OFFLINE=true cargo test -p ml --lib snapshot_size_is_stable` — passes (149*4=596 bytes; B1.0/B1.1a/B1.1b stable)
+- `SQLX_OFFLINE=true cargo test -p ml --test sp13_layer_b_oracle_tests fingerprint_bumped_from_pre_b1_1a` — passes (B1.1a fingerprint still bumped vs pre-B1.1a; B1.1b adds no new fingerprint changes)
+- `SQLX_OFFLINE=true cargo test -p ml --test sp13_layer_b_oracle_tests health_diag_snap_size_stable_at_149_floats` — passes
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp13_layer_b_oracle_tests --features cuda -- --ignored --nocapture` — 15 GPU tests pass on local RTX 3050 Ti (9 from B1.1a + 6 new from B1.1b)
+
+### Next: Smoke A — L40S 5-epoch validation of full SP13 stack
+
+Smoke A runs the full SP13 chain (P0a + P0b + B0 + B0.1 + B1.0 + B1.1a + B1.1b) on L40S. Expected behavior post-B1.1b: aux head trains on real K=2 softmax CE with -1/0/1 labels; `aux_dir_acc_short_ema` should rise meaningfully above 0.5 within the first epoch (vs B1.1a's degraded "predict 0 everywhere" baseline that sat at 0.5 because half the labels would have been class 1 if real). HEALTH_DIAG `aux_b1_diag` line emits per-epoch with `n_down/n_up/n_skip/mask_frac` showing label distribution. If `aux_dir_acc_short_ema > 0.55` by epoch 5, B1.1b is validated and the chain merges to main. If not, root-cause investigation against the producer kernel correctness (the 6 B1.1b oracle tests should cover the failure modes; if a real-data pathology emerges, that's a bug in the kernel logic itself).