diff --git a/crates/backend-uzu/src/backends/common/gpu_types/gemm.rs b/crates/backend-uzu/src/backends/common/gpu_types/gemm.rs
index 2c9377dd4..5d1ebf676 100644
--- a/crates/backend-uzu/src/backends/common/gpu_types/gemm.rs
+++ b/crates/backend-uzu/src/backends/common/gpu_types/gemm.rs
@@ -18,6 +18,7 @@ bitflags! {
         const ACCUMULATE = 1 << 1;
         const BIAS       = 1 << 2;
         const RHT        = 1 << 3;
+        const SOFT_CAP   = 1 << 4;
     }
 }
 
diff --git a/crates/backend-uzu/src/backends/common/kernel/matmul/arguments.rs b/crates/backend-uzu/src/backends/common/kernel/matmul/arguments.rs
index 30ce849fc..925812443 100644
--- a/crates/backend-uzu/src/backends/common/kernel/matmul/arguments.rs
+++ b/crates/backend-uzu/src/backends/common/kernel/matmul/arguments.rs
@@ -9,6 +9,7 @@ pub struct MatmulArguments<'a, 'b, 'd, B: Backend, TB: BufferArg<'b, B> = &'b Al
     pub b_transpose: bool,
     pub d: &'d mut Allocation<B>,
     pub d_transform: MatmulDOps<'d, B>,
+    pub gather_indices: Option<&'a Allocation<B>>,
     pub m: u32,
     pub n: u32,
     pub k: u32,
diff --git a/crates/backend-uzu/src/backends/common/kernel/matmul/d_ops.rs b/crates/backend-uzu/src/backends/common/kernel/matmul/d_ops.rs
index 5ea89c7bd..55448769e 100644
--- a/crates/backend-uzu/src/backends/common/kernel/matmul/d_ops.rs
+++ b/crates/backend-uzu/src/backends/common/kernel/matmul/d_ops.rs
@@ -5,6 +5,7 @@ pub struct MatmulDOps<'a, B: Backend> {
     pub accumulate: bool,
     pub bias: Option<&'a Allocation<B>>,
     pub rht_factors: Option<&'a Allocation<B>>,
+    pub soft_cap: Option<f32>,
 }
 
 impl<'a, B: Backend> MatmulDOps<'a, B> {
@@ -14,6 +15,7 @@ impl<'a, B: Backend> MatmulDOps<'a, B> {
             accumulate: false,
             bias: None,
             rht_factors: None,
+            soft_cap: None,
         }
     }
 
@@ -31,6 +33,9 @@ impl<'a, B: Backend> MatmulDOps<'a, B> {
         if self.rht_factors.is_some() {
             m |= GemmDTransform::RHT;
         }
+        if self.soft_cap.is_some() {
+            m |= GemmDTransform::SOFT_CAP;
+        }
         m
     }
 
@@ -59,6 +64,11 @@ impl<'a, B: Backend> MatmulDOps<'a, B> {
             } else {
                 self.rht_factors
             },
+            soft_cap: if bits.contains(GemmDTransform::SOFT_CAP) {
+                None
+            } else {
+                self.soft_cap
+            },
         }
     }
 }
diff --git a/crates/backend-uzu/src/backends/cpu/kernel/matmul/kernel.rs b/crates/backend-uzu/src/backends/cpu/kernel/matmul/kernel.rs
index bfd00c048..c0341077c 100644
--- a/crates/backend-uzu/src/backends/cpu/kernel/matmul/kernel.rs
+++ b/crates/backend-uzu/src/backends/cpu/kernel/matmul/kernel.rs
@@ -58,6 +58,7 @@ impl MatmulKernel for MatmulCpuKernel {
         let accumulate = arguments.d_transform.accumulate;
         let bias_alloc = arguments.d_transform.bias;
         let post_rht = arguments.d_transform.rht_factors;
+        let soft_cap = arguments.d_transform.soft_cap;
 
         let MatmulArguments {
             a,
@@ -69,6 +70,7 @@ impl MatmulKernel for MatmulCpuKernel {
             m,
             n,
             k,
+            gather_indices,
             ..
         } = arguments;
 
@@ -86,6 +88,10 @@ impl MatmulKernel for MatmulCpuKernel {
             let r = bias.as_buffer_range_ref();
             SendPtr(unsafe { &*r.buffer().get() }.as_ptr().wrapping_byte_add(r.range().start))
         });
+        let gather_ptr = gather_indices.map(|indices| {
+            let r = indices.as_buffer_range_ref();
+            SendPtr(unsafe { &*r.buffer().get() }.as_ptr().wrapping_byte_add(r.range().start) as *const u32)
+        });
         let d_buffer_range = d.as_buffer_range_mut();
         let d_ptr = SendPtrMut(unsafe {
             (&*d_buffer_range.buffer().get()).as_ptr().wrapping_byte_add(d_buffer_range.range().start) as *mut u8
@@ -122,6 +128,11 @@ impl MatmulKernel for MatmulCpuKernel {
             unsafe {
                 for row in 0..m_u {
                     for col in 0..n_u {
+                        // Gather remaps output column `col` to B-row `gather_indices[row * n + col]`.
+                        let b_col = match gather_ptr {
+                            Some(g) => *g.as_ptr().add(row * n_u + col) as usize,
+                            None => col,
+                        };
                         let mut accumulator = 0.0f32;
                         for inner in 0..k_u {
                             let a_value = read_f32(a_ptr.as_ptr(), input_data_type, row * k_u + inner);
@@ -132,9 +143,9 @@ impl MatmulKernel for MatmulCpuKernel {
                                     transpose,
                                 } => {
                                     let index = if *transpose {
-                                        col * leading_dimension + inner
+                                        b_col * leading_dimension + inner
                                     } else {
-                                        inner * leading_dimension + col
+                                        inner * leading_dimension + b_col
                                     };
                                     read_f32(ptr.as_ptr(), weights_data_type, index)
                                 },
@@ -147,7 +158,7 @@ impl MatmulKernel for MatmulCpuKernel {
                                     group_size,
                                 } => {
                                     let (num_groups_k, zero_point_stride, pack_factor) = quant_layout.unwrap();
-                                    let weight_linear_index = col * k_u + inner;
+                                    let weight_linear_index = b_col * k_u + inner;
                                     let quantized_value = if *bits == 4 {
                                         let word_index = weight_linear_index / pack_factor;
                                         let bit_offset = (weight_linear_index % pack_factor) * 4;
@@ -160,11 +171,14 @@ impl MatmulKernel for MatmulCpuKernel {
                                         ((w.add(word_index).read_unaligned() >> bit_offset) & 0xFF) as f32
                                     };
                                     let group_index = inner / group_size;
-                                    let scale =
-                                        read_f32(scales.as_ptr(), weights_data_type, col * num_groups_k + group_index);
+                                    let scale = read_f32(
+                                        scales.as_ptr(),
+                                        weights_data_type,
+                                        b_col * num_groups_k + group_index,
+                                    );
                                     let bias_term = if let Some(zp) = zero_points {
                                         let zero_point = if *bits == 4 {
-                                            let byte_index = col * zero_point_stride + (group_index >> 1);
+                                            let byte_index = b_col * zero_point_stride + (group_index >> 1);
                                             let byte_value = *zp.as_ptr().add(byte_index);
                                             if (group_index & 1) == 0 {
                                                 (byte_value & 0x0F) as f32
@@ -172,11 +186,11 @@ impl MatmulKernel for MatmulCpuKernel {
                                                 ((byte_value >> 4) & 0x0F) as f32
                                             }
                                         } else {
-                                            *zp.as_ptr().add(col * zero_point_stride + group_index) as f32
+                                            *zp.as_ptr().add(b_col * zero_point_stride + group_index) as f32
                                         };
                                         -scale * zero_point
                                     } else if let Some(b) = biases {
-                                        read_f32(b.as_ptr(), weights_data_type, col * num_groups_k + group_index)
+                                        read_f32(b.as_ptr(), weights_data_type, b_col * num_groups_k + group_index)
                                     } else {
                                         let midpoint = (1u32 << (bits - 1)) as f32;
                                         -scale * midpoint
@@ -195,6 +209,9 @@ impl MatmulKernel for MatmulCpuKernel {
                         if let Some(bias) = bias_ptr {
                             value += read_f32(bias.as_ptr(), weights_data_type, col);
                         }
+                        if let Some(cap) = soft_cap {
+                            value = cap * (value / cap).tanh();
+                        }
                         write_f32(d_ptr.as_ptr(), output_data_type, output_index, value);
                     }
                 }
diff --git a/crates/backend-uzu/src/backends/metal/kernel/common/soft_cap.h b/crates/backend-uzu/src/backends/metal/kernel/common/soft_cap.h
new file mode 100644
index 000000000..9a8c798c9
--- /dev/null
+++ b/crates/backend-uzu/src/backends/metal/kernel/common/soft_cap.h
@@ -0,0 +1,14 @@
+#pragma once
+
+#include <metal_stdlib>
+
+#include "defines.h"
+
+namespace uzu {
+
+template <typename T>
+METAL_FUNC T apply_soft_cap(T value, float cap) {
+  return static_cast<T>(cap * metal::fast::tanh(static_cast<float>(value) / cap));
+}
+
+} // namespace uzu
diff --git a/crates/backend-uzu/src/backends/metal/kernel/generated/gemm.h b/crates/backend-uzu/src/backends/metal/kernel/generated/gemm.h
index d20ebd2b4..cdf2e6208 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/generated/gemm.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/generated/gemm.h
@@ -20,6 +20,7 @@ struct GemmDTransform {
   static constant constexpr uint ACCUMULATE = 1 << 1;
   static constant constexpr uint BIAS = 1 << 2;
   static constant constexpr uint RHT = 1 << 3;
+  static constant constexpr uint SOFT_CAP = 1 << 4;
   constexpr bool contains(uint flag) const thread { return (raw_value & flag) != 0; }
   constexpr bool contains(uint flag) const constant { return (raw_value & flag) != 0; }
   constexpr uint bits() const thread { return raw_value; }
diff --git a/crates/backend-uzu/src/backends/metal/kernel/logit_soft_cap/logit_soft_cap.metal b/crates/backend-uzu/src/backends/metal/kernel/logit_soft_cap/logit_soft_cap.metal
index e3e57cb75..833be0209 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/logit_soft_cap/logit_soft_cap.metal
+++ b/crates/backend-uzu/src/backends/metal/kernel/logit_soft_cap/logit_soft_cap.metal
@@ -1,5 +1,6 @@
 #include <metal_stdlib>
 #include "../common/dsl.h"
+#include "../common/soft_cap.h"
 
 template <typename T>
 VARIANTS(T, float, bfloat)
@@ -9,6 +10,5 @@ PUBLIC KERNEL(LogitSoftCap)(
     constant float& soft_cap,
     const uint position AXIS(length, 256)
 ) {
-  const float value = float(logits[position]);
-  logits[position] = T(fast::tanh(value / soft_cap) * soft_cap);
+  logits[position] = uzu::apply_soft_cap<T>(logits[position], soft_cap);
 }
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemm/kernel.rs b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemm/kernel.rs
index 92a479edc..b99f152c4 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemm/kernel.rs
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemm/kernel.rs
@@ -111,6 +111,10 @@ impl GemmKernel {
         &self,
         arguments: &MatmulArguments<'a, 'b, 'd, Metal, TB>,
     ) -> bool {
+        if arguments.gather_indices.is_some() {
+            // TODO: gathered GEMM
+            return false;
+        }
         match (
             arguments.m,
             arguments.n == arguments.k,
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/b_source.h b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/b_source.h
index b552bac54..5a185b6e8 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/b_source.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/b_source.h
@@ -25,7 +25,10 @@ struct BSource {
       const device uint8_t* zero_points,
       const device BT* biases,
       const device AT* a,
+      const device uint* gather_indices,
+      bool gathered,
       uint in_vec_size,
+      uint out_vec_size,
       uint out_row,
       uint batch_idx,
       uint simd_lane,
@@ -36,7 +39,10 @@ struct BSource {
           result,
           b,
           a,
+          gather_indices,
+          gathered,
           in_vec_size,
+          out_vec_size,
           out_row,
           batch_idx,
           simd_lane,
@@ -50,7 +56,10 @@ struct BSource {
           zero_points,
           biases,
           a,
+          gather_indices,
+          gathered,
           in_vec_size,
+          out_vec_size,
           out_row,
           batch_idx,
           simd_lane
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/epilogue.h b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/epilogue.h
index 393784ea9..b51a3f82e 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/epilogue.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/epilogue.h
@@ -1,5 +1,6 @@
 #pragma once
 
+#include "../../../common/soft_cap.h"
 #include "../../../hadamard_transform/hadamard_transform.h"
 
 namespace uzu {
@@ -14,6 +15,7 @@ struct Epilogue {
       const device int32_t* hadamard_factors,
       threadgroup U* shared_results,
       float ab_scale,
+      float soft_cap,
       GemmDTransform output_transform,
       uint out_row,
       uint out_vec_size,
@@ -25,6 +27,7 @@ struct Epilogue {
     const bool is_scale = output_transform.contains(GemmDTransform::SCALE);
     const bool is_accumulate = output_transform.contains(GemmDTransform::ACCUMULATE);
     const bool is_bias = output_transform.contains(GemmDTransform::BIAS);
+    const bool is_soft_cap = output_transform.contains(GemmDTransform::SOFT_CAP);
     const bool use_hadamard = output_transform.contains(GemmDTransform::RHT);
 
     if (writer && simd_lane == 0) {
@@ -41,6 +44,9 @@ struct Epilogue {
         if (is_bias && global_row < out_vec_size) {
           value += static_cast<U>(output_bias[global_row]);
         }
+        if (is_soft_cap) {
+          value = apply_soft_cap(value, soft_cap);
+        }
         result[row] = value;
       }
     }
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/full_precision_b_source.h b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/full_precision_b_source.h
index b406681cb..9305379a5 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/full_precision_b_source.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/full_precision_b_source.h
@@ -11,7 +11,10 @@ struct FullPrecisionBSource {
       thread U (&result)[RESULTS_PER_SIMDGROUP],
       const device uint32_t* b,
       const device AT* a,
+      const device uint* gather_indices,
+      bool gathered,
       uint in_vec_size,
+      uint out_vec_size,
       uint out_row,
       uint batch_idx,
       uint simd_lane,
@@ -24,19 +27,28 @@ struct FullPrecisionBSource {
 
     const uint k_stride = K_SPLIT * block_size;
     const uint k_start = k_slice * block_size;
+    const uint thread_k = simd_lane * values_per_thread + k_start;
+    const device AT* input = a + batch_idx * in_vec_size + thread_k;
+
+    // One advancing weight pointer per output row; base row = out_row (dense) or gather index.
+    const uint base_row = gathered ? 0u : out_row;
     const device BT* weights = reinterpret_cast<const device BT*>(b);
-    weights += out_row * in_vec_size + simd_lane * values_per_thread + k_start;
-    const device AT* input = a + batch_idx * in_vec_size + simd_lane * values_per_thread + k_start;
+    weights += base_row * in_vec_size + thread_k;
+    thread const device BT* weight_rows[RESULTS_PER_SIMDGROUP];
+    METAL_PRAGMA_UNROLL
+    for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
+      const uint addr_row = gathered ? gather_indices[batch_idx * out_vec_size + out_row + row] : row;
+      weight_rows[row] = weights + addr_row * in_vec_size;
+    }
 
     uint k = k_start;
     for (; k + block_size <= in_vec_size; k += k_stride) {
       float4 input_values = static_cast<float4>(*reinterpret_cast<const device I4*>(input));
       METAL_PRAGMA_UNROLL
       for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
-        const device BT* weight_row = weights + row * in_vec_size;
-        result[row] += dot(static_cast<float4>(*reinterpret_cast<const device W4*>(weight_row)), input_values);
+        result[row] += dot(static_cast<float4>(*reinterpret_cast<const device W4*>(weight_rows[row])), input_values);
+        weight_rows[row] += k_stride;
       }
-      weights += k_stride;
       input += k_stride;
     }
 
@@ -48,9 +60,8 @@ struct FullPrecisionBSource {
       if (remaining > 0) {
         METAL_PRAGMA_UNROLL
         for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
-          const device BT* weight_row = weights + row * in_vec_size;
           for (int index = 0; index < remaining; index++) {
-            result[row] += static_cast<U>(weight_row[index]) * static_cast<U>(input[index]);
+            result[row] += static_cast<U>(weight_rows[row][index]) * static_cast<U>(input[index]);
           }
         }
       }
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_b_source.h b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_b_source.h
index 24723ff43..3aa1639e0 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_b_source.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_b_source.h
@@ -24,7 +24,10 @@ struct QuantizedBSource {
       const device uint8_t* zero_points,
       const device BT* biases,
       const device AT* a,
+      const device uint* gather_indices,
+      bool gathered,
       uint in_vec_size,
+      uint out_vec_size,
       uint out_row,
       uint batch_idx,
       uint simd_lane
@@ -41,10 +44,14 @@ struct QuantizedBSource {
     const uint weights_row_stride = in_vec_size * bytes_per_pack / pack_factor;
     const uint group_count = (in_vec_size + GROUP_SIZE - 1) / GROUP_SIZE;
     const uint group_offset = simd_lane / scale_step_per_thread;
+
+    // Base row = out_row (dense) or 0 (gather); rows indexed relative to it.
+    const uint base_row = gathered ? 0u : out_row;
+
     const device uint8_t* weights = reinterpret_cast<const device uint8_t*>(b);
-    weights += out_row * weights_row_stride + simd_lane * packs_per_thread * bytes_per_pack;
+    weights += base_row * weights_row_stride + simd_lane * packs_per_thread * bytes_per_pack;
 
-    RowState row_state(scales, zero_points, biases, out_row, group_count, group_offset);
+    RowState row_state(scales, zero_points, biases, base_row, group_count, group_offset);
 
     const device AT* input = a + batch_idx * in_vec_size + simd_lane * values_per_thread;
     thread U input_values[values_per_thread];
@@ -54,10 +61,11 @@ struct QuantizedBSource {
       U input_sum = load_vector<AT, U, values_per_thread, BITS>(input, input_values);
 
       RowParams row_params;
-      row_state.load(row_params);
+      row_state.load(row_params, gather_indices, gathered, batch_idx, out_vec_size, out_row);
       METAL_PRAGMA_UNROLL
       for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
-        const device uint8_t* weight_row = weights + row * weights_row_stride;
+        const uint addr_row = gathered ? gather_indices[batch_idx * out_vec_size + out_row + row] : row;
+        const device uint8_t* weight_row = weights + addr_row * weights_row_stride;
         result[row] += qdot<U, values_per_thread, BITS>(
             weight_row,
             input_values,
@@ -81,10 +89,11 @@ struct QuantizedBSource {
         U input_sum = load_vector_safe<AT, U, values_per_thread>(input, input_values, remaining);
 
         RowParams row_params;
-        row_state.load(row_params);
+        row_state.load(row_params, gather_indices, gathered, batch_idx, out_vec_size, out_row);
         METAL_PRAGMA_UNROLL
         for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
-          const device uint8_t* weight_row = weights + row * weights_row_stride;
+          const uint addr_row = gathered ? gather_indices[batch_idx * out_vec_size + out_row + row] : row;
+          const device uint8_t* weight_row = weights + addr_row * weights_row_stride;
           result[row] += qdot_safe<U, values_per_thread, BITS>(
               weight_row,
               input_values,
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_row_state.h b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_row_state.h
index 1467610fc..e62b01d62 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_row_state.h
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/common/quantized_row_state.h
@@ -18,8 +18,8 @@ struct QuantizedGroupRows {
   const device T* values;
   uint group_stride;
 
-  QuantizedGroupRows(const device T* values_base, uint out_row, uint group_count, uint group_offset)
-      : values(values_base + out_row * group_count + group_offset), group_stride(group_count) {}
+  QuantizedGroupRows(const device T* values_base, uint base_row, uint group_count, uint group_offset)
+      : values(values_base + base_row * group_count + group_offset), group_stride(group_count) {}
 
   METAL_FUNC U value(uint row) const { return static_cast<U>(values[row * group_stride]); }
 
@@ -36,11 +36,11 @@ struct QuantizedOffsetState<BT, U, GemmBPrologueKind::ScaleBiasDequant, BITS> {
   QuantizedOffsetState(
       const device uint8_t*,
       const device BT* biases_base,
-      uint out_row,
+      uint base_row,
       uint group_count,
       uint group_offset
   )
-      : biases(biases_base, out_row, group_count, group_offset) {}
+      : biases(biases_base, base_row, group_count, group_offset) {}
 
   METAL_FUNC U value(uint row, U) const { return biases.value(row); }
 
@@ -56,17 +56,17 @@ struct QuantizedOffsetState<BT, U, GemmBPrologueKind::ScaleZeroPointDequant, BIT
   QuantizedOffsetState(
       const device uint8_t* zero_points_base,
       const device BT*,
-      uint out_row,
+      uint base_row,
       uint group_count,
       uint group_offset
   ) {
     if constexpr (BITS == 4) {
       zero_point_stride = (group_count + 1) / 2;
-      zero_points = zero_points_base + out_row * zero_point_stride + group_offset / 2;
+      zero_points = zero_points_base + base_row * zero_point_stride + group_offset / 2;
       zero_point_shift = (group_offset & 1) * 4u;
     } else {
       zero_point_stride = group_count;
-      zero_points = zero_points_base + out_row * zero_point_stride + group_offset;
+      zero_points = zero_points_base + base_row * zero_point_stride + group_offset;
     }
   }
 
@@ -107,19 +107,28 @@ struct QuantizedRowState {
       const device BT* scales_base,
       const device uint8_t* zero_points_base,
       const device BT* biases_base,
-      uint out_row,
+      uint base_row,
       uint group_count,
       uint group_offset
   )
-      : scale_rows(scales_base, out_row, group_count, group_offset),
-        offset_state(zero_points_base, biases_base, out_row, group_count, group_offset) {}
-
-  void load(thread Params& params) const {
+      : scale_rows(scales_base, base_row, group_count, group_offset),
+        offset_state(zero_points_base, biases_base, base_row, group_count, group_offset) {}
+
+  // Row index: dense = local row, gather = absolute gather index.
+  void load(
+      thread Params& params,
+      const device uint* gather_indices,
+      bool gathered,
+      uint batch_idx,
+      uint out_vec_size,
+      uint out_row
+  ) const {
     METAL_PRAGMA_UNROLL
     for (uint row = 0; row < RESULTS_PER_SIMDGROUP; row++) {
-      const U scale = scale_rows.value(row);
+      const uint addr_row = gathered ? gather_indices[batch_idx * out_vec_size + out_row + row] : row;
+      const U scale = scale_rows.value(addr_row);
       params.scale[row] = scale;
-      params.offset[row] = offset_state.value(row, scale);
+      params.offset[row] = offset_state.value(addr_row, scale);
     }
   }
 
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/gemv.metal b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/gemv.metal
index a51f9a4fb..e2d454d56 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/gemv.metal
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/gemv.metal
@@ -62,12 +62,16 @@ KERNEL(Gemv)(
         OPTIONAL(output_transform.contains(GemmDTransform::BIAS)),
     const device int32_t* hadamard_factors
         OPTIONAL(output_transform.contains(GemmDTransform::RHT)),
+    const device uint* gather_indices OPTIONAL(gathered),
     const constant uint& in_vec_size,
     const constant uint& out_vec_size,
     const constant uint& batch_size,
     const constant float& ab_scale,
     const constant uint& group_count_x,
+    const constant float& soft_cap
+        OPTIONAL(output_transform.contains(GemmDTransform::SOFT_CAP)),
     const GemmDTransform output_transform SPECIALIZE,
+    const bool gathered SPECIALIZE,
     threadgroup float shared_results[NUM_SIMDGROUPS * RESULTS_PER_SIMDGROUP],
     const uint batch_idx GROUPS(batch_size),
     const uint out_block_idx GROUPS(group_count_x),
@@ -88,7 +92,10 @@ KERNEL(Gemv)(
       zero_points,
       biases,
       a,
+      gather_indices,
+      gathered,
       in_vec_size,
+      out_vec_size,
       tile.out_row,
       batch_idx,
       simd_lane,
@@ -111,6 +118,7 @@ KERNEL(Gemv)(
       hadamard_factors,
       shared_results,
       ab_scale,
+      soft_cap,
       output_transform,
       tile.out_row,
       out_vec_size,
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/kernel.rs b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/kernel.rs
index fdaf77b73..6b3b6a85e 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/kernel.rs
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/gemv/kernel.rs
@@ -35,6 +35,7 @@ pub(crate) struct GemvSpecialization {
     k_split: u32,
     results_per_simdgroup: u32,
     num_simdgroups: u32,
+    gathered: bool,
 }
 
 impl GemvSpecialization {
@@ -49,6 +50,7 @@ impl GemvSpecialization {
             return None;
         }
         let is_quant = !matches!(args.b, MatmulB::FullPrecision { .. });
+        let gathered = args.gather_indices.is_some();
         let bad_leading_dimension = if is_quant {
             args.b_leading_dimension.is_some()
         } else {
@@ -104,6 +106,7 @@ impl GemvSpecialization {
             k_split: tile.k_split,
             results_per_simdgroup: tile.results_per_simdgroup,
             num_simdgroups: tile.num_simdgroups,
+            gathered,
         })
     }
 }
@@ -158,6 +161,7 @@ impl GemvDispatch {
                     specialization.results_per_simdgroup,
                     specialization.num_simdgroups,
                     specialization.output_transform,
+                    specialization.gathered,
                 )
                 .map_err(MatmulError::BackendError)?;
                 Ok(entry.insert(kernel))
@@ -174,6 +178,7 @@ impl GemvDispatch {
         let ab_scale = arguments.d_transform.ab_scale;
         let output_bias = arguments.d_transform.bias;
         let rht_factors = arguments.d_transform.rht_factors;
+        let soft_cap = arguments.d_transform.soft_cap;
 
         let MatmulArguments {
             a,
@@ -183,6 +188,7 @@ impl GemvDispatch {
             m,
             n,
             k,
+            gather_indices,
             ..
         } = arguments;
 
@@ -208,11 +214,13 @@ impl GemvDispatch {
                     &mut *d,
                     output_bias,
                     rht_factors,
+                    gather_indices,
                     k,
                     n,
                     m,
                     ab_scale,
                     group_count_x,
+                    soft_cap,
                     encoder,
                 );
             },
@@ -256,11 +264,13 @@ impl GemvDispatch {
                     &mut *d,
                     output_bias,
                     rht_factors,
+                    gather_indices,
                     k,
                     n,
                     m,
                     ab_scale,
                     group_count_x,
+                    soft_cap,
                     encoder,
                 );
             },
diff --git a/crates/backend-uzu/src/backends/metal/kernel/matmul/mod.rs b/crates/backend-uzu/src/backends/metal/kernel/matmul/mod.rs
index 574e1afaa..deed61a7e 100644
--- a/crates/backend-uzu/src/backends/metal/kernel/matmul/mod.rs
+++ b/crates/backend-uzu/src/backends/metal/kernel/matmul/mod.rs
@@ -66,6 +66,17 @@ impl MatmulKernel for MatmulMetalKernel {
         {
             return self.gemv.encode(arguments, gemv, encoder).map_err(MetalError::from);
         }
+
+        // TODO: remove after GatherGEMM is supported
+        if arguments.gather_indices.is_some() {
+            return Err(MetalError::KernelDispatchFailed(
+                format!(
+                    "gathered readout requires the GEMV path, but shape (m={}, n={}) routes to GEMM",
+                    arguments.m, arguments.n
+                )
+                .into(),
+            ));
+        }
         self.gemm.encode(arguments, encoder)
     }
 }
diff --git a/crates/backend-uzu/src/encodable_block/embedding.rs b/crates/backend-uzu/src/encodable_block/embedding.rs
index 9f62c86f4..96ff54878 100644
--- a/crates/backend-uzu/src/encodable_block/embedding.rs
+++ b/crates/backend-uzu/src/encodable_block/embedding.rs
@@ -545,6 +545,7 @@ impl<B: Backend> Embedding<B> {
                             b_transpose: true,
                             d: &mut output_allocation,
                             d_transform: MatmulDOps::none(),
+                            gather_indices: None,
                             m: batch_dim as u32,
                             n: self.vocab_size,
                             k: self.model_dim,
@@ -612,6 +613,7 @@ impl<B: Backend> Embedding<B> {
                             b_transpose: true,
                             d: &mut output_allocation,
                             d_transform: MatmulDOps::none(),
+                            gather_indices: None,
                             m: batch_dim as u32,
                             n: self.vocab_size,
                             k: self.model_dim,
@@ -633,6 +635,75 @@ impl<B: Backend> Embedding<B> {
 
         Ok(output_allocation)
     }
+
+    /// Per-row candidate readout via the GEMV B-row gather: `out[r][j] == dense[r][token_ids[r][j]]`,
+    /// soft-capped when configured, one dispatch. Full-precision weights; caller guarantees `token_ids < vocab_size`.
+    #[allow(dead_code)]
+    pub(crate) fn encode_readout_sparse(
+        &self,
+        input: &Allocation<B>,
+        token_ids: &Allocation<B>,
+        rows: usize,
+        ids_per_row: usize,
+        encoder: &mut Encoder<B>,
+    ) -> Result<Allocation<B>, EmbeddingError<B>> {
+        assert!(rows > 0 && ids_per_row > 0);
+        let (weights, readout) = match &self.tying {
+            EmbeddingTying::Tied {
+                ty:
+                    TiedEmbeddingType::FullPrecision {
+                        weights,
+                        readout,
+                        ..
+                    },
+            } => (weights, readout),
+            EmbeddingTying::Untied {
+                output_ty:
+                    UntiedEmbeddingReadoutType::FullPrecision {
+                        weights,
+                        readout,
+                    },
+                ..
+            } => (weights, readout),
+            _ => {
+                return Err(EmbeddingError::UnsupportedConfiguration(
+                    "fused sparse readout requires full-precision embedding weights".to_string(),
+                ));
+            },
+        };
+
+        let mut output = encoder
+            .allocate_scratch(size_for_shape(&[rows, ids_per_row], self.data_type))
+            .map_err(EmbeddingError::BackendError)?;
+
+        let soft_cap = self.logit_soft_cap.as_ref().map(|cap| cap.value);
+        readout
+            .lock()
+            .encode(
+                MatmulArguments {
+                    a: input,
+                    a_offset: 0,
+                    b: MatmulB::FullPrecision {
+                        b: weights,
+                    },
+                    b_leading_dimension: None,
+                    b_transpose: true,
+                    d: &mut output,
+                    d_transform: MatmulDOps {
+                        soft_cap,
+                        ..MatmulDOps::none()
+                    },
+                    gather_indices: Some(token_ids),
+                    m: rows as u32,
+                    n: ids_per_row as u32,
+                    k: self.model_dim,
+                },
+                encoder,
+            )
+            .map_err(EmbeddingError::BackendError)?;
+
+        Ok(output)
+    }
 }
 
 fn input_quantization_from_spec<B: Backend>(
diff --git a/crates/backend-uzu/src/encodable_block/linear/matmul.rs b/crates/backend-uzu/src/encodable_block/linear/matmul.rs
index 9ddacb2f6..beafbb42a 100644
--- a/crates/backend-uzu/src/encodable_block/linear/matmul.rs
+++ b/crates/backend-uzu/src/encodable_block/linear/matmul.rs
@@ -263,10 +263,9 @@ impl<B: Backend> Linear<B> for LinearMatmul<B> {
             _ => None,
         };
         let d_transform = MatmulDOps {
-            ab_scale: 1.0,
-            accumulate: false,
             bias: self.biases.as_ref(),
             rht_factors,
+            ..MatmulDOps::none()
         };
 
         self.kernel.lock().encode(
@@ -278,6 +277,7 @@ impl<B: Backend> Linear<B> for LinearMatmul<B> {
                 b_transpose: true,
                 d: &mut output,
                 d_transform,
+                gather_indices: None,
                 m: batch_dim as u32,
                 n: self.output_dim as u32,
                 k: self.input_dim as u32,
diff --git a/crates/backend-uzu/src/encodable_block/linear/qlora_wrapper.rs b/crates/backend-uzu/src/encodable_block/linear/qlora_wrapper.rs
index b194b5d9b..5f49743de 100644
--- a/crates/backend-uzu/src/encodable_block/linear/qlora_wrapper.rs
+++ b/crates/backend-uzu/src/encodable_block/linear/qlora_wrapper.rs
@@ -184,6 +184,7 @@ impl<B: Backend> Linear<B> for QLoRALinearWrapper<B> {
                     b_transpose: true,
                     d: &mut intermediate,
                     d_transform: MatmulDOps::none(),
+                    gather_indices: None,
                     m: batch_dim as u32,
                     n: self.lora_rank as u32,
                     k: self.input_dim as u32,
@@ -223,11 +224,10 @@ impl<B: Backend> Linear<B> for QLoRALinearWrapper<B> {
                     b_transpose: true,
                     d: &mut output,
                     d_transform: MatmulDOps {
-                        ab_scale: 1.0,
                         accumulate: true,
-                        bias: None,
-                        rht_factors: None,
+                        ..MatmulDOps::none()
                     },
+                    gather_indices: None,
                     m: batch_dim as u32,
                     n: self.output_dim as u32,
                     k: self.lora_rank as u32,
diff --git a/crates/backend-uzu/src/encodable_block/mixer/attention/core/fallback.rs b/crates/backend-uzu/src/encodable_block/mixer/attention/core/fallback.rs
index 38a0db774..be55a128e 100644
--- a/crates/backend-uzu/src/encodable_block/mixer/attention/core/fallback.rs
+++ b/crates/backend-uzu/src/encodable_block/mixer/attention/core/fallback.rs
@@ -101,10 +101,9 @@ impl<B: Backend> AttentionFallbackCore<B> {
                     d: &mut group_scores,
                     d_transform: MatmulDOps {
                         ab_scale: scale,
-                        accumulate: false,
-                        bias: None,
-                        rht_factors: None,
+                        ..MatmulDOps::none()
                     },
+                    gather_indices: None,
                     m: (gqa_factor * arguments.suffix_length) as u32,
                     n: sequence_length as u32,
                     k: self.head_dim as u32,
@@ -150,6 +149,7 @@ impl<B: Backend> AttentionFallbackCore<B> {
                     b_transpose: false,
                     d: &mut group_output,
                     d_transform: MatmulDOps::none(),
+                    gather_indices: None,
                     m: (gqa_factor * arguments.suffix_length) as u32,
                     n: self.head_dim as u32,
                     k: sequence_length as u32,
diff --git a/crates/backend-uzu/src/tests/matmul/harness.rs b/crates/backend-uzu/src/tests/matmul/harness.rs
index a89023859..afa46974d 100644
--- a/crates/backend-uzu/src/tests/matmul/harness.rs
+++ b/crates/backend-uzu/src/tests/matmul/harness.rs
@@ -150,6 +150,7 @@ fn run<B: Backend, T: ArrayElement + Float>(
         accumulate: input.case.accumulate,
         bias: bias_allocation.as_ref(),
         rht_factors: rht_allocation.as_ref(),
+        soft_cap: None,
     };
 
     let mut encoder = Encoder::new(context).expect("encoder");
@@ -165,6 +166,7 @@ fn run<B: Backend, T: ArrayElement + Float>(
             b_transpose: input.case.b_transpose,
             d: &mut d_allocation,
             d_transform,
+            gather_indices: None,
             m: m as u32,
             n: n as u32,
             k: k as u32,
diff --git a/crates/backend-uzu/src/tests/matmul/quant.rs b/crates/backend-uzu/src/tests/matmul/quant.rs
index 61096e670..af49f65b5 100644
--- a/crates/backend-uzu/src/tests/matmul/quant.rs
+++ b/crates/backend-uzu/src/tests/matmul/quant.rs
@@ -161,6 +161,7 @@ pub fn quant_arguments<'a, B: Backend, T: ArrayElement + Float>(
         b_transpose: true,
         d: &mut buffers.y,
         d_transform: MatmulDOps::none(),
+        gather_indices: None,
         m: input.m,
         n: input.n,
         k: input.k,
diff --git a/crates/backend-uzu/unit/backends/common/kernel/attention/attention_fallback_test.rs b/crates/backend-uzu/unit/backends/common/kernel/attention/attention_fallback_test.rs
index 5298f0743..4e64548b9 100644
--- a/crates/backend-uzu/unit/backends/common/kernel/attention/attention_fallback_test.rs
+++ b/crates/backend-uzu/unit/backends/common/kernel/attention/attention_fallback_test.rs
@@ -142,10 +142,9 @@ fn pipeline_output<B: Backend>(
                     d: &mut grp_s,
                     d_transform: MatmulDOps {
                         ab_scale: scale,
-                        accumulate: false,
-                        bias: None,
-                        rht_factors: None,
+                        ..MatmulDOps::none()
                     },
+                    gather_indices: None,
                     m: GQA * SUFFIX,
                     n: SEQ,
                     k: HEAD_DIM,
@@ -181,6 +180,7 @@ fn pipeline_output<B: Backend>(
                     b_transpose: false,
                     d: &mut grp_o,
                     d_transform: MatmulDOps::none(),
+                    gather_indices: None,
                     m: GQA * SUFFIX,
                     n: HEAD_DIM,
                     k: SEQ,
diff --git a/crates/backend-uzu/unit/backends/common/kernel/matmul/gemm_bench.rs b/crates/backend-uzu/unit/backends/common/kernel/matmul/gemm_bench.rs
index 49d66ab7d..d5d430814 100644
--- a/crates/backend-uzu/unit/backends/common/kernel/matmul/gemm_bench.rs
+++ b/crates/backend-uzu/unit/backends/common/kernel/matmul/gemm_bench.rs
@@ -65,6 +65,7 @@ fn bench_gemm(c: &mut Criterion) {
                                 b_transpose: true,
                                 d: &mut d,
                                 d_transform: MatmulDOps::none(),
+                                gather_indices: None,
                                 m: m as u32,
                                 n: n as u32,
                                 k: k as u32,
diff --git a/crates/backend-uzu/unit/backends/common/kernel/matmul/gemv_test.rs b/crates/backend-uzu/unit/backends/common/kernel/matmul/gemv_test.rs
index d383f063d..932448051 100644
--- a/crates/backend-uzu/unit/backends/common/kernel/matmul/gemv_test.rs
+++ b/crates/backend-uzu/unit/backends/common/kernel/matmul/gemv_test.rs
@@ -10,7 +10,8 @@ use crate::{
     array::ArrayElement,
     backends::{
         common::{
-            Backend, Context, Encoder,
+            Allocation, Backend, Context, Encoder,
+            gpu_types::QuantizationMethod,
             kernel::{
                 Kernels,
                 matmul::{MatmulArguments, MatmulB, MatmulDOps, MatmulKernel},
@@ -21,6 +22,7 @@ use crate::{
     tests::{
         assert::assert_eq_float,
         helpers::{alloc_allocation, alloc_allocation_with_data, allocation_to_vec},
+        matmul::{QuantBuffers, QuantInput},
     },
 };
 
@@ -30,6 +32,8 @@ struct Input<T: ArrayElement + Float> {
     m: usize,
     k: usize,
     n: usize,
+    ids: Option<Box<[u32]>>,
+    soft_cap: Option<f32>,
 }
 
 fn get_test_data<T: ArrayElement + Float>(
@@ -46,49 +50,72 @@ fn get_test_data<T: ArrayElement + Float>(
         m,
         k,
         n,
+        ids: None,
+        soft_cap: None,
     };
 
     let expected = get_output::<T, Cpu>(&input);
     (input, expected)
 }
 
-fn get_output<T: ArrayElement + Float, B: Backend>(input: &Input<T>) -> Vec<T> {
-    let context = B::Context::new().expect("Failed to create Context");
-
-    let m = input.m as u32;
-    let k = input.k as u32;
-    let n = input.n as u32;
-
-    let b_allocation = alloc_allocation_with_data::<B, T>(&context, &input.b);
-    let a_allocation = alloc_allocation_with_data::<B, T>(&context, &input.a);
-    let mut d_allocation = alloc_allocation::<B, T>(&context, input.m * input.n);
-
+// Encode one GEMV (dense, or a per-row B-row gather when `gather_indices` is set) and copy out.
+fn run_gemv<'a, B: Backend, T: ArrayElement + Float>(
+    context: &B::Context,
+    a: &'a Allocation<B>,
+    b: MatmulB<'a, B>,
+    gather_indices: Option<&'a Allocation<B>>,
+    m: usize,
+    n_out: usize,
+    k: usize,
+    soft_cap: Option<f32>,
+) -> Vec<T> {
+    let mut d = alloc_allocation::<B, T>(context, m * n_out);
     let mut kernel =
-        <B::Kernels as Kernels>::MatmulKernel::new(&context, T::data_type(), T::data_type(), T::data_type())
-            .expect("Failed to create MatmulKernel");
-
-    let mut encoder = Encoder::new(context.as_ref()).expect("Failed to create encoder");
+        <B::Kernels as Kernels>::MatmulKernel::new(context, T::data_type(), T::data_type(), T::data_type())
+            .expect("MatmulKernel");
+    let mut encoder = Encoder::new(context).expect("encoder");
     kernel
         .encode(
             MatmulArguments {
-                a: &a_allocation,
+                a,
                 a_offset: 0,
-                b: MatmulB::FullPrecision {
-                    b: &b_allocation,
-                },
+                b,
                 b_leading_dimension: None,
                 b_transpose: true,
-                d: &mut d_allocation,
-                d_transform: MatmulDOps::none(),
-                m,
-                n,
-                k,
+                d: &mut d,
+                d_transform: MatmulDOps {
+                    soft_cap,
+                    ..MatmulDOps::none()
+                },
+                gather_indices,
+                m: m as u32,
+                n: n_out as u32,
+                k: k as u32,
             },
             &mut encoder,
         )
         .expect("encode failed");
     encoder.end_encoding().submit().wait_until_completed().unwrap();
-    allocation_to_vec::<B, T>(&d_allocation)
+    allocation_to_vec::<B, T>(&d)
+}
+
+fn get_output<T: ArrayElement + Float, B: Backend>(input: &Input<T>) -> Vec<T> {
+    let context = B::Context::new().expect("Failed to create Context");
+    let a = alloc_allocation_with_data::<B, T>(&context, &input.a);
+    let weights = alloc_allocation_with_data::<B, T>(&context, &input.b);
+    let ids = input.ids.as_ref().map(|ids| alloc_allocation_with_data::<B, u32>(&context, ids));
+    run_gemv::<B, T>(
+        &context,
+        &a,
+        MatmulB::FullPrecision {
+            b: &weights,
+        },
+        ids.as_ref(),
+        input.m,
+        input.n,
+        input.k,
+        input.soft_cap,
+    )
 }
 
 fn test<T: ArrayElement + Float + Debug + Display>(
@@ -137,3 +164,116 @@ fn gemv_f32(
 ) {
     test::<f32>(m, k, n, 0.01);
 }
+
+fn assert_gather<T: ArrayElement + Float + Debug + Display>(
+    dense: &[T],
+    gather: &[T],
+    ids: &[u32],
+    m: usize,
+    vocab: usize,
+    ids_per_row: usize,
+    eps: f32,
+    name: &str,
+) {
+    let mut expected = vec![T::from(0.0).unwrap(); m * ids_per_row];
+    for r in 0..m {
+        for c in 0..ids_per_row {
+            expected[r * ids_per_row + c] = dense[r * vocab + ids[r * ids_per_row + c] as usize];
+        }
+    }
+    assert_eq_float(&expected, gather, eps, &format!("gather vs dense ({name})"));
+}
+
+macro_rules! check_gather {
+    ($m:expr, $vocab:expr, $ids:expr, $ids_per_row:expr, $eps:expr, |$B:ident| $run:block) => {{
+        {
+            #[allow(non_camel_case_types)]
+            type $B = Cpu;
+            let (dense, gather) = $run;
+            assert_gather(&dense, &gather, &$ids, $m, $vocab, $ids_per_row, $eps, "Cpu");
+        }
+        for_each_non_cpu_backend!(|$B| {
+            let (dense, gather) = $run;
+            assert_gather(&dense, &gather, &$ids, $m, $vocab, $ids_per_row, $eps, std::any::type_name::<$B>());
+        });
+    }};
+}
+
+fn fp_gather_case<T: ArrayElement + Float + Debug + Display>(
+    soft_cap: Option<f32>,
+    eps: f32,
+) {
+    let (m, k, vocab, ids_per_row) = (4usize, 128usize, 256usize, 8usize);
+    let a: Vec<T> = (0..m * k).map(|i| T::from(((i % 13) as f32) * 0.1 - 0.6).unwrap()).collect();
+    let weights: Vec<T> = (0..vocab * k).map(|i| T::from(((i % 17) as f32) * 0.1 - 0.8).unwrap()).collect();
+    let ids: Vec<u32> = (0..m * ids_per_row).map(|i| ((i * 37 + 11) % vocab) as u32).collect();
+
+    // Dense (`n = vocab`, no ids) and gather (`n = ids_per_row`, ids) share `a`/`weights`/soft-cap.
+    let make = |n: usize, ids: Option<Box<[u32]>>| Input {
+        a: a.clone().into_boxed_slice(),
+        b: weights.clone().into_boxed_slice(),
+        m,
+        k,
+        n,
+        ids,
+        soft_cap,
+    };
+    let dense_input = make(vocab, None);
+    let gather_input = make(ids_per_row, Some(ids.clone().into_boxed_slice()));
+
+    check_gather!(m, vocab, ids, ids_per_row, eps, |B| {
+        (get_output::<T, B>(&dense_input), get_output::<T, B>(&gather_input))
+    });
+}
+
+#[uzu_test]
+fn gemv_gather() {
+    // Full precision: one call per dtype (generic over T, so bf16/f32 can't be a runtime loop).
+    for soft_cap in [None, Some(15.0)] {
+        fp_gather_case::<bf16>(soft_cap, 0.1);
+        fp_gather_case::<f32>(soft_cap, 0.01);
+    }
+    // Quantized (bf16, per bits/method) — inline, since it isn't type-generic.
+    for (bits, method) in [
+        (4, QuantizationMethod::ScaleBias),
+        (4, QuantizationMethod::ScaleZeroPoint),
+        (4, QuantizationMethod::ScaleSymmetric),
+        (8, QuantizationMethod::ScaleZeroPoint),
+    ] {
+        let (m, k, vocab, ids_per_row, group_size) = (4usize, 128usize, 64usize, 8usize, 32u32);
+        let input = QuantInput::<bf16>::new(m, k, vocab, group_size, bits, method, 0x5EED);
+        let ids: Vec<u32> = (0..m * ids_per_row).map(|i| ((i * 37 + 11) % vocab) as u32).collect();
+        // K_SPLIT == 1 keeps k in one reduction, so gather and dense share the exact accumulation.
+        check_gather!(m, vocab, ids, ids_per_row, 0.05, |B| {
+            let context = <B as Backend>::Context::new().expect("context");
+            let buffers = QuantBuffers::<B, bf16>::allocate(&context, &input);
+            let ids_alloc = alloc_allocation_with_data::<B, u32>(&context, &ids);
+            let variant = || match method {
+                QuantizationMethod::ScaleBias => MatmulB::ScaleBiasDequant {
+                    b: &buffers.w,
+                    scales: &buffers.scales,
+                    biases: buffers.bias.as_ref().expect("bias buffer"),
+                    mode: input.mode,
+                    group_size: input.group_size,
+                },
+                QuantizationMethod::ScaleZeroPoint => MatmulB::ScaleZeroPointDequant {
+                    b: &buffers.w,
+                    scales: &buffers.scales,
+                    zero_points: buffers.zp.as_ref().expect("zp buffer"),
+                    mode: input.mode,
+                    group_size: input.group_size,
+                },
+                QuantizationMethod::ScaleSymmetric => MatmulB::ScaleSymmetricDequant {
+                    b: &buffers.w,
+                    scales: &buffers.scales,
+                    mode: input.mode,
+                    group_size: input.group_size,
+                },
+            };
+            (
+                run_gemv::<B, bf16>(&context, &buffers.x, variant(), None, m, vocab, k, None),
+                run_gemv::<B, bf16>(&context, &buffers.x, variant(), Some(&ids_alloc), m, ids_per_row, k, None),
+            )
+        });
+    }
+}
diff --git a/crates/backend-uzu/unit/backends/common/kernel/matmul/quant_dispatch_test.rs b/crates/backend-uzu/unit/backends/common/kernel/matmul/quant_dispatch_test.rs
index 2b336c65f..d9422a7f1 100644
--- a/crates/backend-uzu/unit/backends/common/kernel/matmul/quant_dispatch_test.rs
+++ b/crates/backend-uzu/unit/backends/common/kernel/matmul/quant_dispatch_test.rs
@@ -199,10 +199,8 @@ fn parity_bf16_gs32_4bit_mlx_with_bias() {
     let mut encoder = Encoder::<Metal>::new(&context).expect("encoder");
     let mut args = quant_arguments(&mut buffers, &input);
     args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
-        accumulate: false,
         bias: Some(&bias_pp_buf),
-        rht_factors: None,
+        ..MatmulDOps::none()
     };
     matmul.gemm.encode_dispatch_path(args, GemmDispatchPath::Simdgroup, &mut encoder).expect("encode quant with bias");
     encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -252,9 +250,8 @@ fn parity_bf16_gemv_qmv_fused_scale_bias() {
     let mut args = quant_arguments(&mut buffers, &input);
     args.d_transform = MatmulDOps {
         ab_scale: scale,
-        accumulate: false,
         bias: Some(&bias_buf),
-        rht_factors: None,
+        ..MatmulDOps::none()
     };
     matmul.encode(args, &mut encoder).expect("encode quant gemv with scale+bias");
     encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -330,10 +327,9 @@ fn parity_bf16_gemv_quant_rht_with_bias() {
     let mut cpu_encoder = Encoder::<Cpu>::new(&cpu_context).expect("cpu encoder");
     let mut cpu_args = quant_arguments(&mut cpu_buffers, &input);
     cpu_args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
-        accumulate: false,
         bias: Some(&cpu_bias),
         rht_factors: Some(&cpu_rht),
+        ..MatmulDOps::none()
     };
     cpu_matmul.encode(cpu_args, &mut cpu_encoder).expect("cpu encode quant+rht+bias");
     cpu_encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -352,10 +348,9 @@ fn parity_bf16_gemv_quant_rht_with_bias() {
     let mut encoder = Encoder::<Metal>::new(&context).expect("encoder");
     let mut args = quant_arguments(&mut buffers, &input);
     args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
-        accumulate: false,
         bias: Some(&metal_bias),
         rht_factors: Some(&metal_rht),
+        ..MatmulDOps::none()
     };
     matmul.encode(args, &mut encoder).expect("encode quant gemv with rht+bias");
     encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -393,10 +388,8 @@ fn parity_bf16_gemv_quant_rht() {
     let mut cpu_encoder = Encoder::<Cpu>::new(&cpu_context).expect("cpu encoder");
     let mut cpu_args = quant_arguments(&mut cpu_buffers, &input);
     cpu_args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
-        accumulate: false,
-        bias: None,
         rht_factors: Some(&cpu_rht),
+        ..MatmulDOps::none()
     };
     cpu_matmul.encode(cpu_args, &mut cpu_encoder).expect("cpu encode quant+rht");
     cpu_encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -415,10 +408,8 @@ fn parity_bf16_gemv_quant_rht() {
     let mut encoder = Encoder::<Metal>::new(&context).expect("encoder");
     let mut args = quant_arguments(&mut buffers, &input);
     args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
-        accumulate: false,
-        bias: None,
         rht_factors: Some(&metal_rht),
+        ..MatmulDOps::none()
     };
     matmul.encode(args, &mut encoder).expect("encode quant gemv with rht");
     encoder.end_encoding().submit().wait_until_completed().unwrap();
@@ -443,10 +434,8 @@ fn quant_gemm_accumulate_returns_unsupported_dop() {
     let mut encoder = Encoder::<Metal>::new(&context).expect("encoder");
     let mut args = quant_arguments(&mut buffers, &input);
     args.d_transform = MatmulDOps {
-        ab_scale: 1.0,
         accumulate: true,
-        bias: None,
-        rht_factors: None,
+        ..MatmulDOps::none()
     };
     let result = matmul.encode(args, &mut encoder);