diff --git a/CHANGELOG.md b/CHANGELOG.md
index bbaf491..721f66d 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -9,6 +9,10 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ### Fixed
 
+- `use_avx512_mr12` documented itself as "DEFAULT ON" with a
+  `YSCV_AVX512_SGEMM=0` kill switch. The gate it calls is opt-in on the exact
+  value `1`, so the AVX-512 MR=12×NR=32 GEMM is off unless asked for — the
+  comment claimed the opposite of the code directly below it.
 - `remove_dropout_nodes` deleted nodes by `NodeProto.name`, which ONNX makes
   optional. A single unnamed `Dropout` put the empty string into the delete set
   and took every other unnamed node in the graph with it — on a fully unnamed
@@ -22,6 +26,26 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ### Changed
 
+- The `prfm pldl1keep` helper lived in three copies across `conv/pointwise.rs`,
+  and is now a single `ops::prefetch::prefetch_l1_keep` documenting when a hint
+  is worth adding: hoist the gate out of the K-loop, and only hint operands
+  whose stride actually defeats the hardware prefetcher.
+- Dropped the software prefetch from the AVX `binary_same_shape` loops. The
+  access is unit-stride over three streams, which every hardware prefetcher
+  tracks; measured on Zen 4 the hint is within run-to-run drift (the A/B flips
+  sign depending on run order), so it only added uops to a bandwidth-bound loop.
+
+### Added
+
+- `rebuild_runtime_index` now debug-asserts one execution-plan action per node.
+  With the sequential fallback gone the runner walks the plan and nothing else,
+  so a short plan would silently skip trailing nodes.
+- Benchmarks for the two kernels that had none: `trans_a_m64_k256_n256` (the
+  tracker's `FusedTransposeMatMul` shape, the only caller of the transposed-A
+  tiles) and raw-slice `binary_same_shape_dispatch` add/mul.
+
+### Changed
+
 - Optimizer passes no longer rebuild the runtime index individually. Nine of
   the twelve passes called `rebuild_runtime_index()` on exit, so a single
   `optimize_onnx_graph` re-ran execution-plan construction and weight
diff --git a/crates/yscv-kernels/benches/kernels_cpu_ops.rs b/crates/yscv-kernels/benches/kernels_cpu_ops.rs
index a06254e..300b3ec 100644
--- a/crates/yscv-kernels/benches/kernels_cpu_ops.rs
+++ b/crates/yscv-kernels/benches/kernels_cpu_ops.rs
@@ -2,12 +2,12 @@ use std::num::NonZeroUsize;
 
 use criterion::{Criterion, black_box, criterion_group, criterion_main};
 use yscv_kernels::{
-    Backend, BatchNorm2dParams, LayerNormLastDimParams, ParallelElementwiseConfig,
+    Backend, BatchNorm2dParams, BinaryKind, LayerNormLastDimParams, ParallelElementwiseConfig,
     ParallelMatmulConfig, SeparableConv2dParams, ThreadedCpuBackend, ThreadedCpuBackendConfig, add,
-    avg_pool2d_nhwc, batch_norm2d_nhwc, conv2d_nhwc, conv2d_nhwc_indirect_padded,
-    conv2d_nhwc_padded, depthwise_conv2d_nhwc, layer_norm_last_dim, log_softmax_last_dim,
-    logsumexp_last_dim, matmul_2d, matmul_2d_sequential, max_pool2d_nhwc, relu,
-    separable_conv2d_nhwc, sigmoid, softmax_last_dim,
+    avg_pool2d_nhwc, batch_norm2d_nhwc, binary_same_shape_dispatch, conv2d_nhwc,
+    conv2d_nhwc_indirect_padded, conv2d_nhwc_padded, depthwise_conv2d_nhwc, layer_norm_last_dim,
+    log_softmax_last_dim, logsumexp_last_dim, matmul_2d, matmul_2d_sequential, max_pool2d_nhwc,
+    relu, separable_conv2d_nhwc, sigmoid, softmax_last_dim,
 };
 use yscv_tensor::Tensor;
 
@@ -64,6 +64,26 @@ fn bench_matmul_modes(c: &mut Criterion) {
         });
     });
 
+    // Transposed-A 4-row tile at the tracker's FusedTransposeMatMul shape
+    // (correlation head: A^T is [k=256, m=64], B is [256, 256]). This is the
+    // only caller of the trans_a kernels, so it is the shape their prefetch
+    // and tiling choices have to be judged on.
+    let a_kt_ftmm = vec![0.31f32; 256 * 64];
+    let b_ftmm = vec![0.17f32; 256 * 256];
+    let mut out_ftmm = vec![0.0f32; 64 * 256];
+    group.bench_function("trans_a_m64_k256_n256", |b| {
+        b.iter(|| {
+            yscv_kernels::matmul_2d_slices_trans_a(
+                black_box(&a_kt_ftmm),
+                64,
+                256,
+                black_box(&b_ftmm),
+                256,
+                black_box(&mut out_ftmm),
+            );
+        });
+    });
+
     group.bench_function("rect_96x192x64", |b| {
         b.iter(|| {
             let out = matmul_2d(black_box(&lhs_rect), black_box(&rhs_rect)).expect("matmul rect");
@@ -170,6 +190,29 @@ fn bench_elementwise_modes(c: &mut Criterion) {
             black_box(out);
         });
     });
+    // Raw-slice entry point: no Tensor allocation per iteration, so the SIMD
+    // loop itself is what gets timed rather than the surrounding bookkeeping.
+    let mut raw_out = vec![0.0; lhs.data().len()];
+    group.bench_function("add_same_shape_raw_slice", |b| {
+        b.iter(|| {
+            binary_same_shape_dispatch(
+                black_box(lhs.data()),
+                black_box(rhs.data()),
+                black_box(&mut raw_out),
+                BinaryKind::Add,
+            );
+        });
+    });
+    group.bench_function("mul_same_shape_raw_slice", |b| {
+        b.iter(|| {
+            binary_same_shape_dispatch(
+                black_box(lhs.data()),
+                black_box(rhs.data()),
+                black_box(&mut raw_out),
+                BinaryKind::Mul,
+            );
+        });
+    });
     group.bench_function("add_same_shape_threaded_2", |b| {
         b.iter(|| {
             let out = threaded_backend
diff --git a/crates/yscv-kernels/src/ops/conv/pointwise.rs b/crates/yscv-kernels/src/ops/conv/pointwise.rs
index 60292af..ba9d18f 100644
--- a/crates/yscv-kernels/src/ops/conv/pointwise.rs
+++ b/crates/yscv-kernels/src/ops/conv/pointwise.rs
@@ -1239,16 +1239,14 @@ unsafe fn pointwise_nx16_direct_rows_avx2(
     }
 }
 
-/// L1 prefetch hint (`prfm pldl1keep`) for the strided PW-reduce weight stream.
-/// The weight stride is `n*4` bytes (96-448 B for tracker reduce shapes) — too
-/// large for the in-order A53's HW prefetcher stride detector, so a manual hint
-/// several K-iters ahead hides the load latency behind the FMA pipe.
+/// L1 prefetch hint for the strided PW-reduce weight stream. The weight stride
+/// is `n*4` bytes (96-448 B for tracker reduce shapes) — too large for the
+/// in-order A53's HW prefetcher stride detector, so a manual hint several
+/// K-iters ahead hides the load latency behind the FMA pipe. The local
+/// `PREFETCH_AHEAD` of 8 below is tuned for this weight stream and is
+/// deliberately longer than the shared GEMM default.
 #[cfg(target_arch = "aarch64")]
-#[inline(always)]
-#[allow(unsafe_code, unsafe_op_in_unsafe_fn)]
-unsafe fn prefetch_l1_keep(p: *const f32) {
-    core::arch::asm!("prfm pldl1keep, [{p}]", p = in(reg) p, options(nostack, preserves_flags, readonly));
-}
+use super::super::prefetch::prefetch_l1_keep;
 
 #[cfg(target_arch = "aarch64")]
 #[target_feature(enable = "neon")]
diff --git a/crates/yscv-kernels/src/ops/matmul/avx512_mr12.rs b/crates/yscv-kernels/src/ops/matmul/avx512_mr12.rs
index d099e79..232e498 100644
--- a/crates/yscv-kernels/src/ops/matmul/avx512_mr12.rs
+++ b/crates/yscv-kernels/src/ops/matmul/avx512_mr12.rs
@@ -245,7 +245,8 @@ pub(super) fn pack_b_panel_nr32(
 /// - m >= 12 (tail rows handled via tail_tile in gebp)
 /// - k >= 16 (smaller k doesn't amortize the blocked-GEMM setup)
 ///
-/// DEFAULT ON. Kill-switch: `YSCV_AVX512_SGEMM=0` disables.
+/// DEFAULT OFF — opt in with the exact value `YSCV_AVX512_SGEMM=1`. See
+/// `avx512_mr12_enabled` below for why Zen 4 loses to the AVX2 4×24 tile.
 /// All epilogue combinations supported: None/Relu (in asm),
 /// SiLU (ZMM post-store via apply_silu_zmm_tile), residual.
 #[cfg(all(target_arch = "x86_64", any(target_os = "linux", target_os = "macos")))]
diff --git a/crates/yscv-kernels/src/ops/mod.rs b/crates/yscv-kernels/src/ops/mod.rs
index 58cdbbe..bf16ddc 100644
--- a/crates/yscv-kernels/src/ops/mod.rs
+++ b/crates/yscv-kernels/src/ops/mod.rs
@@ -22,6 +22,7 @@ pub mod nchwc_pack;
 mod nchwc_pointwise;
 mod norm;
 mod pool;
+mod prefetch;
 pub mod quantize;
 pub mod rope;
 pub mod simd;
diff --git a/crates/yscv-kernels/src/ops/prefetch.rs b/crates/yscv-kernels/src/ops/prefetch.rs
new file mode 100644
index 0000000..f537de9
--- /dev/null
+++ b/crates/yscv-kernels/src/ops/prefetch.rs
@@ -0,0 +1,34 @@
+//! L1 prefetch hint shared by the kernels that walk a strided operand.
+//!
+//! A kernel that reads its weight or `B` operand down a column at stride
+//! `n * 4` bytes leaves the range a hardware stride detector tracks once `n`
+//! grows. On an in-order core the resulting miss stalls the FMA pipe outright,
+//! so a hint issued several K-iterations ahead covers the latency.
+//!
+//! Two rules the measured results here depend on:
+//!
+//! * Hoist the decision out of the inner loop. Split the K-loop into a
+//!   prefetching range and a plain tail rather than testing a flag per
+//!   iteration — on the Cortex-A53 the per-iteration branch alone costs more
+//!   than the hint saves.
+//! * Only add a hint where the operand stride actually defeats the hardware
+//!   prefetcher. Unit-stride streaming loops do not need one, and out-of-order
+//!   cores cover strided GEMM operands on their own; both cases measure as
+//!   neutral and only add uops.
+
+/// Issue an L1 "keep" prefetch hint for `p`.
+///
+/// `p` is a hint only and is never dereferenced, so it may point at a line the
+/// caller will not read. It must still be a pointer the caller formed without
+/// running off the end of the allocation, since computing an out-of-bounds
+/// pointer is already undefined behaviour regardless of this call.
+#[cfg(target_arch = "aarch64")]
+#[inline(always)]
+#[allow(unsafe_code, unsafe_op_in_unsafe_fn)]
+pub(crate) unsafe fn prefetch_l1_keep(p: *const f32) {
+    core::arch::asm!(
+        "prfm pldl1keep, [{p}]",
+        p = in(reg) p,
+        options(nostack, preserves_flags, readonly),
+    );
+}
diff --git a/crates/yscv-kernels/src/ops/simd/binary.rs b/crates/yscv-kernels/src/ops/simd/binary.rs
index d9de455..f77075c 100644
--- a/crates/yscv-kernels/src/ops/simd/binary.rs
+++ b/crates/yscv-kernels/src/ops/simd/binary.rs
@@ -507,23 +507,12 @@ unsafe fn binary_same_shape_avx(lhs: &[f32], rhs: &[f32], out: &mut [f32], kind:
     let out_ptr = out.as_mut_ptr();
     let mut index = 0usize;
 
-    // 4x unrolled: process 32 floats per iteration with software prefetch.
-    // Matches vDSP throughput by keeping the OoO pipeline fully saturated.
+    // 4x unrolled: process 32 floats per iteration. The access is unit-stride
+    // over three streams, which every hardware prefetcher tracks, so an
+    // explicit hint only adds uops to a bandwidth-bound loop.
     match kind {
         BinaryKind::Add => {
             while index + 32 <= len {
-                #[cfg(target_arch = "x86")]
-                {
-                    use std::arch::x86::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
-                #[cfg(target_arch = "x86_64")]
-                {
-                    use std::arch::x86_64::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
                 let a0 = _mm256_loadu_ps(left_ptr.add(index));
                 let b0 = _mm256_loadu_ps(right_ptr.add(index));
                 let a1 = _mm256_loadu_ps(left_ptr.add(index + 8));
@@ -541,18 +530,6 @@ unsafe fn binary_same_shape_avx(lhs: &[f32], rhs: &[f32], out: &mut [f32], kind:
         }
         BinaryKind::Sub => {
             while index + 32 <= len {
-                #[cfg(target_arch = "x86")]
-                {
-                    use std::arch::x86::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
-                #[cfg(target_arch = "x86_64")]
-                {
-                    use std::arch::x86_64::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
                 let a0 = _mm256_loadu_ps(left_ptr.add(index));
                 let b0 = _mm256_loadu_ps(right_ptr.add(index));
                 let a1 = _mm256_loadu_ps(left_ptr.add(index + 8));
@@ -570,18 +547,6 @@ unsafe fn binary_same_shape_avx(lhs: &[f32], rhs: &[f32], out: &mut [f32], kind:
         }
         BinaryKind::Mul => {
             while index + 32 <= len {
-                #[cfg(target_arch = "x86")]
-                {
-                    use std::arch::x86::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
-                #[cfg(target_arch = "x86_64")]
-                {
-                    use std::arch::x86_64::_mm_prefetch;
-                    _mm_prefetch::<3>(left_ptr.add(index + 32) as *const i8);
-                    _mm_prefetch::<3>(right_ptr.add(index + 32) as *const i8);
-                }
                 let a0 = _mm256_loadu_ps(left_ptr.add(index));
                 let b0 = _mm256_loadu_ps(right_ptr.add(index));
                 let a1 = _mm256_loadu_ps(left_ptr.add(index + 8));
diff --git a/crates/yscv-onnx/src/loader/mod.rs b/crates/yscv-onnx/src/loader/mod.rs
index bf62661..33d0a9d 100644
--- a/crates/yscv-onnx/src/loader/mod.rs
+++ b/crates/yscv-onnx/src/loader/mod.rs
@@ -400,6 +400,15 @@ impl OnnxModel {
             &self.dw_khwc_weights,
             &self.group_khwc_weights,
         );
+        // Execution now has no non-plan fallback: the runner walks the plan and
+        // nothing else. A plan shorter than the node list would silently skip
+        // the trailing nodes, so hold the one-action-per-node invariant here,
+        // where it is established, rather than at the point it would be missed.
+        debug_assert_eq!(
+            self.runtime_index.execution_plan.len(),
+            self.nodes.len(),
+            "execution plan must carry one action per node"
+        );
     }
 }