Formats and kernels: FP8, INT8, INT4, Marlin, Machete

Take one weight matrix from Llama-3-70B: gate_proj, shape $N \times K = 28672 \times 8192$, 234.9 M parameters. At batch $M = 1$ the GEMM is $2MNK = 0.47$ GFLOP — 0.47 microseconds of H100 tensor-core time at the 989.4 TFLOP/s bf16 peak §0.4 fixed. The weight read is 0.47 GB, which at 3.35 TB/s costs 140 µs. The layer is 300× memory bound; only the weight bytes matter.

INT4 with group size 128 and fp16 scales stores $b_{\text{eff}} = 4 + 16/128 = 4.125$ bits per weight (§0.5 owns that formula), so 121 MB, so 36 µs. That is the prize: 3.9×. Now count what the naive implementation actually moves.

The middle row is not a strawman — it is what you get from torch.matmul(x, dequant(w)), and it is what every "we support INT4" implementation does before somebody writes a real kernel. The quantised weight only pays if it stays packed all the way from HBM into the registers that feed the tensor core. That constraint — never materialise the dequantised weight in memory — is the entire design brief for Marlin, Machete, and every kernel in this chapter.

A tensor core accepts a fixed menu of operand types. If your format is on the menu, the GEMM is an ordinary GEMM and the only interesting question is where the scale factors get applied. If your format is not on the menu — and INT4 weights against bf16 activations are not — the kernel must convert, and the only question that matters is where. Materializing expanded weights in HBM can erase bandwidth savings. Shared-memory and register conversion trade storage, occupancy and scheduling differently. Register conversion can overlap integer/ALU work with tensor-core instructions, but its cost is not automatically zero: measure stalls, register pressure, repeated conversion and achieved throughput.

Two register buffers indexed [k % 2] make the overlap explicit: the kernel issues fetch_to_registers(k+1) before matmul(k), so the loads and the LOP3 conversions for tile $k{+}1$ are in flight while the tensor cores chew on tile $k$. This is classic software pipelining; what is unusual is that a whole numeric format conversion has been hidden inside the slack.

The inline-PTX strings in Marlin's mma wrapper identify the instructions this kernel emits. They are evidence for this implementation, not an exhaustive inventory of every format the GPU ISA supports. The wrapper is short:

    } else if constexpr (std::is_same<scalar_t, nv_bfloat16>::value) {
      float* c = reinterpret_cast<float*>(&frag_c);
      asm volatile(
          "mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 "
          "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n"
          : "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3])
          : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b[0]), "r"(b[1]),
            "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]));
    } else if constexpr (std::is_same<scalar_t, __nv_fp8_e4m3>::value) {
      float* c = reinterpret_cast<float*>(&frag_c);
      asm volatile(
          "mma.sync.aligned.m16n8k16.row.col.f32.e4m3.e4m3.f32 "
          "{%0,%1,%2,%3}, {%4,%5}, {%6}, {%7,%8,%9,%10};\n"
          : "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3])
          : "r"(a[idx * 2]), "r"(a[idx * 2 + 1]), "r"(b[idx]), "f"(c[0]),
            "f"(c[1]), "f"(c[2]), "f"(c[3]));
    } else if constexpr (std::is_same<scalar_t, int8_t>::value) {
      int32_t* c = reinterpret_cast<int32_t*>(&frag_c);
      asm volatile(
          "mma.sync.aligned.m16n8k16.row.col.s32.s8.s8.s32.satfinite "
          "{%0,%1,%2,%3}, {%4,%5}, {%6}, {%7,%8,%9,%10};\n"
          : "=r"(c[0]), "=r"(c[1]), "=r"(c[2]), "=r"(c[3])
          : "r"(a[idx * 2]), "r"(a[idx * 2 + 1]), "r"(b[idx]), "r"(c[0]),
            "r"(c[1]), "r"(c[2]), "r"(c[3]));
    }
  } else if (k_size == 32) {

The wrapper covers fp16, bf16, e4m3 FP8 and s8 operands, mostly through the $m16n8k$ family with $k \in \{8,16,32\}$. SM75 uses m16n8k8 for fp16 and an m8n8k16 sequence for the INT8 k32 case. Floating-point paths accumulate in fp32 except for the explicit fp16-accumulation option; INT8 paths accumulate in s32 (marlin_mma.h:L23-L133). Both operands come from registers. This wrapper has no INT4 entry. A hypothetical s4×s4 instruction would not directly implement W4A16 either: the activation operand remains fp16 or bf16, so that instruction would require an additional activation quantization step and a different numerical contract.

The table separates formats consumed directly by the listed tensor-core instructions from stored weight formats that these kernels first convert. Native operand support does not imply epilogue-only scaling: a scale independent of the reduction index can be factored outside the sum, but a scale varying along K must participate inside that reduction. Among these listed implementations, Blackwell's block-scaled tensor op consumes the sub-byte NVFP4 operands natively, with block-scale tensors laid out by Sm100BlkScaledConfig::tile_atom_to_shape_SFA (nvfp4_scaled_mm_kernels.cu:L155-L168), so the hardware incorporates the per-16-element block scales into the multiplication. This is distinct from applying one global output multiplier in an epilogue.

Converting int4 to fp16 the obvious way costs a shift, a mask, an integer-to-float convert and a multiply per value. The trick vLLM inherits from FasterTransformer costs roughly one instruction per pair of values, by constructing the fp16 bit pattern directly:

template <>
__device__ inline void dequant<half2, vllm::kU4B8.id(), true>(int q,
                                                              half2* frag_b) {
  const int MASK = 0x000f000f;
  const int EX = 0x64006400;
  // Guarantee that the `(a & b) | c` operations are LOP3s.
  int lo = lop3<(0xf0 & 0xcc) | 0xaa>(q, MASK, EX);
  q >>= 4;
  int hi = lop3<(0xf0 & 0xcc) | 0xaa>(q, MASK, EX);

  frag_b[0] = *reinterpret_cast<half2*>(&lo);
  frag_b[1] = *reinterpret_cast<half2*>(&hi);
}

0x6400 is fp16 for 1024. Splicing a 4-bit integer into the low mantissa bits of that pattern yields $1024 + v$ exactly, for $v \in [0,15]$, with no rounding. One LOP3.LUT instruction — a single three-input lookup on the ALU — produces two such fp16 values packed in one register. The $-1024$ correction and the $-8$ zero point are then folded into one __hsub2 against the constant 0x64086408 (dequant.h:L121-L143), and the group scale into one __hmul2. Cost: about 4 ALU instructions per 4 weights, issued on pipes the tensor cores are not using.

Group scaling is why this must happen inside the k-loop. With per-tensor or per-channel scales the scale factors out of the sum, $\sum_k (s_a a_{ik})(s_b b_{kj}) = s_a s_b \sum_k a_{ik} b_{kj}$, and can be applied once at the end. With group size 128 the scale changes every 128 values of k and cannot be hoisted out of the reduction. Marlin therefore scales the register fragment before every mma:

// Multiply dequantized values by the corresponding quantization scale; used
// only for grouped quantization.
template <vllm::ScalarTypeId type_id>
__device__ inline void scale(typename MarlinScalarType<type_id>::FragB& frag_b,
                             typename MarlinScalarType<type_id>::FragS& frag_s,
                             int i) {
  using scalar_t = typename MarlinScalarType<type_id>::scalar_t;
  using scalar_t2 = typename MarlinScalarType<type_id>::scalar_t2;
  scalar_t2 s = MarlinScalarType<type_id>::num2num2(
      reinterpret_cast<scalar_t*>(&frag_s)[i]);
  frag_b[0] = __hmul2(frag_b[0], s);
  frag_b[1] = __hmul2(frag_b[1], s);
}

Where the batch-size crossover is

A weight-only kernel wins while the weight stream dominates and loses once the tensor cores saturate. Equate the two: streaming $NK$ weights at $b_{\text{eff}}$ bits costs $NK b_{\text{eff}}/(8\,\text{BW})$; the math costs $2MNK/P$ where $P$ is the peak of the instruction the kernel actually issues. Setting them equal:

with $b_{\text{eff}}$ the effective bits per stored weight (§0.5), $P$ the tensor-core peak in FLOP/s, $\mathrm{BW}$ HBM bandwidth. On H100 SXM ($P_{\text{bf16}} = 989.4$ TFLOP/s, $P_{\text{fp8}} = 1978.9$ TFLOP/s dense, $\mathrm{BW} = 3.35$ TB/s):

These are the same crossover §4.1.4 derives as $B^{*} = I^{*}b_w/2$, refined by the group metadata: §4.1 charges W4A16 a flat 0.5 B per weight and gets 74, this section charges the 4.125 effective bits and gets 76. The scale bytes you also have to stream keep you memory-bound a fraction longer. W8A16 (148) and W8A8 FP8 (295) are identical in both, since neither carries per-group metadata.

These are derived numbers, and they are the reason every dispatcher in this chapter buckets on $M$ at 16 / 64 / 128 / 256. Marlin's own switch between its small-batch and large-batch tile tables fires at $M > 16$; CUTLASS's FP8 dispatcher has boundaries at 16, 64, 128; Machete's generated heuristic has 16, 32, 64, 128, 256. The hardware ratio puts the interesting region at a few tens of rows, and that is where the tuning effort goes.

Marlin (Elias Frantar, IST-DASLab; the vLLM copy carries "Adapted from https://github.com/IST-DASLab/marlin" at marlin_template.h:L19) starts from an observation about the mma.m16n8k16 B fragment. For a $16(k) \times 8(n)$ B tile, lane $t$ of the warp owns exactly four elements: column $n = t/4$, rows $k \in \{2r,\, 2r{+}1,\, 2r{+}8,\, 2r{+}9\}$ where $r = t \bmod 4$. That is a strided, interleaved scatter — nothing like row-major order.

Every W4A16 kernel has to land its dequantised values in those lanes. Most do it with shared-memory gathers or warp shuffles after conversion. Marlin does it once, offline, in a repack pass, so that at runtime one contiguous 32-bit load plus two LOP3s produce the eight fp16 values a lane needs, in the order it needs them, with zero cross-lane movement. The repack kernel is where the mapping is written down:

    auto warp_id = threadIdx.x / 32;
    auto th_id = threadIdx.x % 32;

    if (warp_id >= 4) {
      return;
    }

    int tc_col = th_id / 4;
    int tc_row = (th_id % 4) * (is_a_8bit ? 4 : 2);

    constexpr int tc_offsets[4] = {0, 1, 8, 9};

    int cur_n = (warp_id / (is_a_8bit ? 2 : 1)) * 16 + tc_col;

tc_col, tc_row and tc_offsets = {0,1,8,9} are the mma fragment map, transcribed. The thread gathers four values for column cur_n into vals[0..3] and four for column cur_n + 8 into vals[4..7], then writes them in permuted nibble order:

    if constexpr (!is_a_8bit && num_bits == 4) {
      int pack_idx[8] = {0, 2, 4, 6, 1, 3, 5, 7};

      uint32_t res = 0;
#pragma unroll
      for (int i = 0; i < 8; i++) {
        res |= vals[pack_idx[i]] << (i * 4);
      }

      out_ptr[out_offset + th_id * 4 + warp_id] = res;

{0,2,4,6,1,3,5,7} looks arbitrary until you line it up against the LOP3. The mask 0x000f000f selects nibbles 0 and 4 into frag_b[0]; after q >>= 4 it selects nibbles 1 and 5 into frag_b[1]. So the dequant emits, in order, nibbles $\langle 0, 4, 1, 5 \rangle$. The permutation places $\text{vals}[0..3]$ at exactly those nibble positions, and $\text{vals}[4..7]$ at nibbles $\langle 2,6,3,7 \rangle$ — which is what a second dequant of q >> 8 reads. One word in, two fragments out, no shuffles.

The pipeline and its budget

The rest of Marlin is a four-stage cp.async pipeline sized to fit shared memory. The constants are deliberately small:

static constexpr int default_threads = 256;

static constexpr int pipe_stages =
    4;  // 4 pipeline stages fit into shared memory

static constexpr int min_thread_n = 64;
static constexpr int min_thread_k = 64;
static constexpr int max_thread_n = 256;

static constexpr int tile_size = 16;
static constexpr int max_par = 16;

Shared-memory demand is computed exactly, not guessed, in marlin.cu:L193-L226: sh_b_size = stages * (tb_k * tb_n / pack_factor) * 4. For the default 128×128 tile at 4 bits that is $4 \times (128 \times 128 / 8) \times 4 = 32{,}768$ bytes of packed B plus $4 \times (16 \times 128) \times 2 = 16{,}384$ bytes of bf16 A — 48 KB against the 228 KB/SM Hopper ceiling §0.3 quotes. That headroom buys four stages instead of two. is_valid_config (marlin.cu:L229-L260) rejects any tile whose computed cache_size exceeds the device limit — which is why a too-large tile shows up as a kernel-selection failure rather than a launch error. The overlap itself is in the main loop:

  #pragma unroll
    for (int pipe = 0; pipe < stages;) {
  #pragma unroll
      for (int k = 0; k < b_sh_wr_iters; k++) {
        fetch_to_registers(k + 1, pipe % stages);
        fetch_scales_to_registers(k + 1, pipe);
        fetch_zp_to_registers(k + 1, pipe);
        if (k == b_sh_wr_iters - 2) {
          fetch_to_shared((pipe + stages - 1) % stages, pipe,
                          slice_iters >= stages);
          pipe++;
          wait_for_stage();
          init_same_group(pipe % stages);
        }
        # ...
          matmul(k, pipe - (k >= b_sh_wr_iters - 2 ? 1 : 0));

Three fetch depths run concurrently: cp.async global→shared two stages ahead, shared→register one k-step ahead, and the matmul consuming the buffer filled last iteration. wait_for_stage() waits on cp_async_wait<stages - 2> with the comment "We only have stages - 2 active fetches since we are double buffering" (marlin_template.h:L924-L931).

Tile tables, and what happens when your shape does not fit

The tile choice is a two-table lookup keyed on nothing more than "is M above one 16-row block":

thread_config_t small_batch_thread_configs[] = {
    // Ordered by priority

    // thread_k, thread_n, num_threads
    {128, 128, 256},
    {64, 128, 128},
    {128, 64, 128}};

thread_config_t large_batch_thread_configs[] = {
    // Ordered by priority

    // thread_k, thread_n, num_threads
    {64, 256, 256},
    {64, 128, 128},
    {128, 64, 128}};

Small batch prefers deep K (128) and moderate N; large batch prefers wide N (256), so each loaded A fragment is reused across more output columns. determine_exec_config (marlin.cu:L276-L325) takes the first entry that both divides the problem and fits in shared memory. Divisibility is the constraint that bites. The Python guard is stricter than the CUDA one — GPTQ_MARLIN_MIN_THREAD_K = 128 against the kernel's min_thread_k = 64 — because a usable tile family needs $n \bmod 64 = 0$ and $k \bmod 128 = 0$, or the transposed pair:

def marlin_padded_nk(size_n: int, size_k: int, group_size: int = -1) -> tuple[int, int]:
    """Minimal (padded_n, padded_k) satisfying a Marlin thread-tile family.

    Marlin GEMM and repack require (n % 64, k % 128) or (n % 128, k % 64);
    shapes satisfying neither are zero-padded up to the cheaper family. K
    stays divisible by group_size so padded scales keep an integral group
    count. Padded weight regions contribute nothing to the GEMM output:
    quantized value 0 decodes to 0.0 (FP4/FP8) or is cancelled by the
    zero-padded scales/zero-points (INT).
    """
    group = group_size if group_size > 0 else 1
    candidates = (
        (round_up(size_n, 64), round_up(size_k, math.lcm(128, group))),
        (round_up(size_n, 128), round_up(size_k, math.lcm(64, group))),
    )
    padded_nk = min(candidates, key=lambda nk: (nk[0] * nk[1], nk[0] + nk[1]))
    if padded_nk != (size_n, size_k):
        logger.warning_once(
            "Marlin requires thread-tile padding for some weight shapes in "
            "this model. Activations and/or outputs of the padded layers are "
            "padded/sliced on every forward; performance may be degraded."
        )
    return padded_nk

As of a556f3f, an awkward shape gets zero-padded rather than rejected. Two cases still cannot be padded away: act-order (g_idx) couples K to the whole-model group layout, so the strict check stays (mixed_precision/marlin.py:L62-L69); and a group straddling a tensor-parallel boundary is fatal, because the rank does not hold a whole group. The strict check produces the error you will see:

    if output_size_per_partition % GPTQ_MARLIN_MIN_THREAD_N != 0:
        raise ValueError(
            f"Weight output_size_per_partition = "
            f"{output_size_per_partition} is not divisible by "
            f" min_thread_n = {GPTQ_MARLIN_MIN_THREAD_N}. "
            "Consider reducing tensor_parallel_size or running "
            "with --quantization gptq."
        )

    # Validate input_size_per_partition
    if input_size_per_partition % GPTQ_MARLIN_MIN_THREAD_K != 0:
        raise ValueError(
            f"Weight input_size_per_partition = "
            f"{input_size_per_partition} is not divisible "
            f"by min_thread_k = {GPTQ_MARLIN_MIN_THREAD_K}. "
            "Consider reducing tensor_parallel_size or running "
            "with --quantization gptq."
        )

Read the suggested fix carefully: "Consider reducing tensor_parallel_size or running with --quantization gptq." The second half is stale at this SHA. As §4.3.8 shows, "gptq" and "gptq_marlin" both map to AutoGPTQConfig, so re-running with --quantization gptq takes exactly the same path and raises exactly the same error. The real fallback is the rest of the priority list — ExLlama or Triton W4A16 — reached only if the shape check passes and the faster kernels decline for another reason. The supported quantisation grid is narrow by design: MARLIN_SUPPORTED_GROUP_SIZES = [-1, 32, 64, 128] at marlin_utils.py:L36, and a group-96 checkpoint has no Marlin kernel at any shape.

Marlin is an Ampere-era kernel: cp.async, warp-level mma, hand-written PTX. Hopper added two things it cannot use — the TMA (a hardware descriptor-driven bulk copy engine) and wgmma (warpgroup-wide MMA, 128 threads issuing one instruction). Machete is the CUTLASS-based rewrite that does. Its README states the relationship plainly:

# Machete (Mixed Precision Cutlass-Based GEMM)

Machete is a spiritual successor to the Marlin kernel but optimized for Hopper architectures and based on Cutlass. Being based on Cutlass, new type pairs and epilogues are easier to add compared to Marlin.

The core trick is the same — prepack the weights to match the tensor-core layout — but the target layout is a wgmma atom rather than an mma fragment, and the block is bigger:

  // TODO (LucasWilkinson): compare the performance for other sizes
  // Prepacked block shape, smallest layout atom for loading into registers
  //   (can contain multiple wgmma instructions worth of data in one block)
  // We ideally want this to be configured such that a thread can perform 128bit
  // loads, i.e. we amount of data associated with each thread within a
  // prepacked block is a multiple of 128bits, when using a cooperative sechdule
  // we have 256 threads working a single block at a time, this means each
  // thread works on `sizeof_bits_v<ElementB> * (128*64) / 256` bits of data,
  // for a 4bit type this would be 128bits

Do the arithmetic in the comment: $4 \times (128 \times 64) / 256 = 128$ bits per thread. The 128×64 block size is chosen precisely so that a cooperative two-warpgroup schedule gives every thread a full 128-bit vector load from shared memory. Marlin's equivalent granularity is one 32-bit word per lane; Machete gets 4× the load width because it has 4× the threads cooperating on one block.

Three further Hopper-specific changes, all visible in machete_mainloop.cuh:

Machete's constraints are tighter than Marlin's, in exchange for the Hopper features:

check_machete_supports_shape is two divisibility checks and is the whole story (vllm/model_executor/layers/quantization/utils/machete_utils.py:L37-L51): if in_features % 64 or out_features % 128 is non-zero, Machete declines and the selector falls through to Marlin. Note that Machete is checked before Marlin in vLLM's CUDA priority list, so on an H100 a divisible shape gets Machete and an awkward one silently gets Marlin.

FP8 needs no dequantisation in the mainloop, because e4m3 is on the tensor-core menu and because the scaling is per-tensor or per-row/column — which, as shown in §4.3.4, factors out of the k-sum entirely. So the whole scale story lives in the CUTLASS epilogue:

template <typename ElementAcc, typename ElementD, typename TileShape>
struct ScaledEpilogue
    : private ScaledEpilogueBase<ElementAcc, ElementD, TileShape> {
 private:
  using SUPER = ScaledEpilogueBase<ElementAcc, ElementD, TileShape>;
  using Accum = typename SUPER::Accum;
  using ScaleA = typename SUPER::template ColOrScalarLoad<float>;
  using ScaleB = typename SUPER::template RowOrScalarLoad<float>;

  using Compute0 = cutlass::epilogue::fusion::Sm90Compute<
      cutlass::multiplies, float, float,
      cutlass::FloatRoundStyle::round_to_nearest>;

  using EVTCompute0 =
      cutlass::epilogue::fusion::Sm90EVT<Compute0, ScaleB, Accum>;

  using Compute1 = cutlass::epilogue::fusion::Sm90Compute<
      cutlass::multiplies, ElementD, float,
      cutlass::FloatRoundStyle::round_to_nearest>;

 public:
  using EVTCompute =
      cutlass::epilogue::fusion::Sm90EVT<Compute1, ScaleA, EVTCompute0>;

Read the type tree bottom-up: EVTCompute0 = ScaleB * Accum in fp32, then EVTCompute = ScaleA * EVTCompute0 cast to ElementD. Two fused multiplies on the fp32 accumulator after the k-loop has finished, with the scale tensors streamed as an epilogue broadcast. ColOrScalarLoad and RowOrScalarLoad are the same code path for per-tensor (a scalar) and per-token / per-channel (a vector) — which is why vLLM converts a fused QKV module's per-tensor scales to per-channel at load time (scaled_mm/cutlass.py:L58-L69) rather than specialising the kernel. The comment above the struct states the requirement: "A and B must have symmetric quantization (zero point == 0)" (scaled_mm_epilogues_c3x.hpp:L141). Asymmetric INT8 needs a separate azp epilogue and a separate kernel file.

Block-wise FP8 (DeepSeek-V3's 128×128 weight blocks) breaks the factoring argument the same way group-wise INT4 does, and gets its own kernels: scaled_mm_blockwise_sm90_fp8.cu, …sm100…, …sm120…. That is three separate compiled dispatch paths for one format, which is a good measure of how much a non-factorable scale costs.

Start with a framing correction that the source forces. As of a556f3f, Marlin is not a quantisation method. The names survive as aliases, but they all resolve to one config class per algorithm:

    method_to_config: dict[str, type[QuantizationConfig]] = {
        "awq": AutoAWQConfig,
        "awq_marlin": AutoAWQConfig,
        "auto_awq": AutoAWQConfig,
        "fp8": Fp8Config,
        "fbgemm_fp8": FBGEMMFp8Config,
        "fp_quant": FPQuantConfig,
        "modelopt": ModelOptFp8Config,
        "modelopt_fp4": ModelOptNvFp4Config,
        "modelopt_mxfp8": ModelOptMxFp8Config,
        "modelopt_mixed": ModelOptMixedPrecisionConfig,
        "auto_gptq": AutoGPTQConfig,
        "gptq": AutoGPTQConfig,
        "gptq_marlin": AutoGPTQConfig,

--quantization gptq and --quantization gptq_marlin load the same AutoGPTQConfig and take the same code path; the flag no longer selects a kernel. Whether you get Marlin is decided per layer, at weight-creation time, by capability. The decision point is four lines long:

        mp_linear_kernel_config = MPLinearLayerConfig(
            full_weight_shape=(input_size, output_size),
            partition_weight_shape=(
                input_size_per_partition,
                output_size_per_partition,
            ),
            weight_type=self.quant_config.quant_type,
            act_type=params_dtype if input_dtype is None else input_dtype,
            group_size=self.quant_config.group_size,
            zero_points=False,
            has_g_idx=self.quant_config.desc_act,
        )

        kernel_type = choose_mp_linear_kernel(mp_linear_kernel_config)

        if kernel_type.__name__ not in self._kernel_backends_being_used:
            logger.info("Using %s for AutoGPTQLinearMethod", kernel_type.__name__)
            self._kernel_backends_being_used.add(kernel_type.__name__)

AutoAWQConfig does the identical thing at auto_awq.py:L430-L448, differing only in zero_points=self.quant_config.zero_point. That logger.info("Using %s …") line, emitted once per distinct backend, reports the selected implementation for this linear method. At this pin, --linear-backend filters the candidate family, while VLLM_DISABLED_KERNELS removes named candidates. The backend filter falls back to the unfiltered list with a warning when that family has no kernel for the layer type (vllm/model_executor/kernels/linear/__init__.py:L353-L383); an explicit preference therefore does not bypass capability and shape checks or guarantee every layer uses that family. One check does run earlier: AutoGPTQLinearMethod.__init__ calls verify_marlin_supported (auto_gptq.py:L320-L324), so a checkpoint whose type or group size is outside Marlin's grid fails at config time, before any kernel gets a vote.

Below that, vLLM's linear kernels declare their own eligibility through a two-method protocol: is_supported(compute_capability) for hardware and can_implement(config) for the quantisation config, both returning (bool, reason) (vllm/model_executor/kernels/linear/base.py:L168-L210). Selection is then a first-match walk down a hand-ordered list:

_POSSIBLE_KERNELS: dict[PlatformEnum, list[type[MPLinearKernel]]] = {
    PlatformEnum.CUDA: [
        CutlassW4A8LinearKernel,
        MacheteLinearKernel,
        AllSparkLinearKernel,
        MarlinLinearKernel,
        ConchLinearKernel,
        ExllamaLinearKernel,
        TritonW4A16LinearKernel,
        HummingLinearKernel,
    ],

If nothing matches, the raised ValueError concatenates every kernel's refusal reason (kernels/linear/__init__.py:L826-L837) — the single most useful diagnostic in the whole quantisation stack, because it tells you in one message why each of eight kernels declined. VLLM_DISABLED_KERNELS knocks one out by class name to see what the next choice does (kernels/linear/__init__.py:L809-L813).

That selection happens once, at load time. The M-bucketing happens per call, inside the chosen kernel. CUTLASS FP8 on Hopper dispatches on both M and N:

  if (m <= 16) {
    // m in [1, 16]
    if (n <= 1280) {
      return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM16_N1280>(
          out, a, b, b_scales, a_scales, std::forward<EpilogueArgs>(args)...);
    }
    return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM16_N8192>(
        out, a, b, b_scales, a_scales, std::forward<EpilogueArgs>(args)...);
  } else if (m <= 64) {
    # ...
  } else if (m <= 128) {
    // m in (64, 128]
    return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM128>(
        out, a, b, a_scales, b_scales, std::forward<EpilogueArgs>(args)...);
  } else if (m >= 8192 && k >= 6144) {
    return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM8192_K6144>(
        out, a, b, a_scales, b_scales, std::forward<EpilogueArgs>(args)...);
  } else {
    // m in (128, inf)
    return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmDefault>(
        out, a, b, a_scales, b_scales, std::forward<EpilogueArgs>(args)...);
  }

Look at the argument order: for $M \le 64$ the scales are passed b_scales, a_scales, swapped. Those configs set the swap_ab template flag — "enable swap AB for M < 64" (:L207) — and compute $C^{\mathsf T} = B^{\mathsf T}A^{\mathsf T}$ so the skinny dimension becomes N, where a tile of 16 is legal. Tile shapes go from _128,_128,_128 at large M down to _64,_16,_256 at M ≤ 16: as M shrinks the tile narrows in N and deepens in K, maximising the number of concurrent CTAs streaming weights.

Machete's heuristic is generated rather than written, from a Python dict compiled into a Jinja dispatch chain:

    # Stored as "condition": ((tile_shape_mn), (cluster_shape_mnk))
    default_tile_heuristic_config = {
        #### M = 257+
        "M > 256 && K <= 16384 && N <= 4096": ((128, 128), (2, 1, 1)),
        "M > 256": ((128, 256), (2, 1, 1)),
        #### M = 129-256
        "M > 128 && K <= 4096 && N <= 4096": ((128, 64), (2, 1, 1)),
        "M > 128 && K <= 8192 && N <= 8192": ((128, 128), (2, 1, 1)),
        "M > 128": ((128, 256), (2, 1, 1)),
        #### M = 65-128
        "M > 64 && K <= 4069 && N <= 4069": ((128, 32), (2, 1, 1)),
        "M > 64 && K <= 4069 && N <= 8192": ((128, 64), (2, 1, 1)),
        "M > 64 && K >= 8192 && N >= 12288": ((256, 128), (2, 1, 1)),
        "M > 64": ((128, 128), (2, 1, 1)),
        #### M = 33-64
        "M > 32 && K <= 6144 && N <= 6144": ((128, 16), (1, 1, 1)),
        "M > 32 && K >= 16384 && N >= 12288": ((256, 64), (2, 1, 1)),
        "M > 32": ((128, 64), (2, 1, 1)),
        #### M = 17-32
        "M > 16 && K <= 12288 && N <= 8192": ((128, 32), (2, 1, 1)),
        "M > 16": ((256, 32), (2, 1, 1)),
        #### M = 1-16
        "N >= 26624": ((256, 16), (1, 1, 1)),
        None: ((128, 16), (1, 1, 1)),
    }

    # For now we use the same heuristic for all types
    # Heuristic is currently tuned for H100s

The comment "Heuristic is currently tuned for H100s" and the typo 4069 (twice, for 4096) are honest evidence about what these tables are: measured on one machine, generated into C++, never re-tuned. Read them as a record of an afternoon's benchmarking, not as physics.

The best-documented example in either repo is the ROCm MXFP8 selector, whose docstring explains itself:

def _select_cfg(M, N, K):
    """(BLOCK_M, BLOCK_N, BLOCK_K, num_warps, num_stages) — graph-tuned on gfx950.

    The M-bucketed, shape-adaptive tile selection here is the speedup over the
    upstream 2-bucket launcher. Tiles are pipelined (num_stages>=2, larger BLOCK_K)
    and occupancy- and shape-aware: keyed on the LOCAL (M, N, K), so it adapts to the
    TP-sharded shapes (e.g. MiniMax-M3 TP=4 vs TP=8, where local N and K differ) —
    large-K prefill uses BLOCK_K=256; short-K (K=768) widens N. BLOCK_K must divide K
    (the K-loop is unmasked), so every BLOCK_K below is guarded to be K-divisible
    (served K: 384/768/1024/2048/6144).
    """
    if M <= 64:
        # decode (M in {1,32,64}): tiny-M GEMV is weight-BW + GPU-OCCUPANCY bound. The
        # lever is NARROW BLOCK_N=16 (maximize N-tiles so more CUs stream the weight in
        # parallel) + LARGE BLOCK_K (fewer K-iters, bigger coalesced weight loads).
        # Tuned by CUDA-graph replay latency. Optimal at both TP=4 and TP=8.
        if K % 1024 == 0:  # K=2048, 6144 -> graph-best 16x16x1024 (all M)
            return 16, 16, 1024, 2, 2
        if K % 512 == 0:
            return 16, 16, 512, 2, 3
        if K % 256 == 0:  # K=768 (shared_down) -> graph-best 16x32x256
            return 16, 32, 256, 4, 3
        return 16, 32, 128, 4, 3

"Tiny-M GEMV is weight-BW + GPU-occupancy bound. The lever is narrow BLOCK_N=16 (maximize N-tiles so more CUs stream the weight in parallel) + large BLOCK_K." That is the same conclusion CUTLASS's _64,_16,_256 small-M tile reached independently, on different hardware, in a different language. When M is small the goal is not data reuse — there is none — it is getting every compute unit pulling weight bytes at once.

SGLang's Marlin support is vLLM's, vendored with the attribution in the header (python/sglang/srt/layers/quantization/marlin_utils.py:L3) and the same constants — GPTQ_MARLIN_TILE = 16, MIN_THREAD_N = 64, MIN_THREAD_K = 128, MARLIN_SUPPORTED_GROUP_SIZES = [-1, 32, 64, 128] at :L54-L59. The CUDA is in-tree but JIT-compiled rather than shipped as an AOT extension:

@cache_once
def _jit_gptq_marlin_module(dtype: torch.dtype) -> Module:
    args = make_cpp_args(dtype)
    return load_jit(
        "gptq_marlin",
        *args,
        cuda_files=["gemm/marlin/gptq_marlin.cuh"],
        cuda_wrappers=[("gptq_marlin_gemm", f"gptq_marlin_gemm<{args}>")],
    )

That gemm/marlin/gptq_marlin.cuh resolves to python/sglang/kernels/jit/csrc/gemm/marlin/gptq_marlin.cuh, 1,005 lines against vLLM's 35 KB marlin.cu plus 82 KB template — an earlier snapshot, with the same small_batch_thread_configs/large_batch_thread_configs tables and the same pipe_stages = 4.

There is no Machete in SGLang at this SHA. A repo-wide grep for "machete" over python/ returns exactly one hit — a comment inherited from vLLM at python/sglang/srt/layers/quantization/utils.py:L508 ("For some kernels (namely Machete) the zero-points are applied after the") — and no kernel, config class or registry entry. On Hopper, a GPTQ W4A16 checkpoint that vLLM would run through Machete runs through Marlin in SGLang. The divergence runs in both directions:

The QuantizationMethods literal at vllm/model_executor/layers/quantization/__init__.py:L12-L46 lists twenty-odd names and neither "bitsandbytes" nor "gguf" is among them. A bnb-quantised or GGUF checkpoint that loads under SGLang will not load under vLLM at this SHA. That is a real capability difference, not a kernel-quality one — and it cuts the other way from Machete, where vLLM has the faster Hopper path.

Where SGLang invests instead is FP8, and it invests in an external package rather than in-tree CUTLASS. The auto backend order is explicit:

def _dispatch_auto_backend() -> Callable:
    """Auto-select the best backend based on hardware capabilities."""
    # Priority order for auto selection:
    # 1. DeepGEMM (if enabled and available)
    # 2. FlashInfer TRTLLM (if Blackwell GPU and FlashInfer available)
    # 3. CUTLASS (if SM120 GPU and CUDA 12.8+)
    # 4. AITER (if AMD GPU with AITER enabled)
    # 5. Triton (fallback)

    if deep_gemm_wrapper.ENABLE_JIT_DEEPGEMM:
        return deepgemm_w8a8_block_fp8_linear_with_fallback
    elif is_blackwell_supported() and is_flashinfer_available():
        return flashinfer_gemm_w8a8_block_fp8_linear_with_fallback
    elif is_sm120_supported():
        return cutlass_w8a8_block_fp8_linear_with_fallback
    elif _use_aiter:
        return aiter_w8a8_block_fp8_linear
    else:
        return triton_w8a8_block_fp8_linear

DeepGEMM first, everywhere it is available. The wrapper is thin — deep_gemm.fp8_gemm_nt(lhs, rhs, out) at python/sglang/srt/layers/deep_gemm_wrapper/entrypoint.py:L189-L193 — with the interesting work in the gating: SM ≥ 90 required, SM120 explicitly excluded ("DeepGEMM requires TMEM/tcgen05 (SM100+datacenter), not available on SM120", configurer.py:L21-L23), and DEEPGEMM_SCALE_UE8M0 switched on for Blackwell — the UE8M0 power-of-two block scale that §0.5 describes.

SGLang's M-bucketing lives in Python rather than C++, and is unusually well documented. The FlashInfer FP8 blockscale path is restricted to a 31-row window with the reason spelled out:

    # Keep this backend to 1 <= M < 32, mirroring vLLM's
    # FlashInferFp8DeepGEMMDynamicBlockScaledKernel. fp8_blockscale_gemm_sm90 is
    # one entry point over two kernels and only the M < 32 swapAB half is worth
    # taking:
    #   M >= 32 picks the non-swapAB kernel, which is slower than DeepGEMM (worst
    #     just above the threshold) and, on some checkpoints, less accurate.
    #   M == 0 hard-fails inside the kernel ("Check failed: (input_ptr !=
    #     nullptr)"). Empty batches are a normal steady-state input, not an edge
    #     case: DP attention hands an idle rank a zero-token forward so the
    #     collectives stay in sync (ScheduleBatch.prepare_for_idle).
    # Same shape of guard as the gfx95 CK M bound below.
    m_supported = 1 <= input.view(-1, input.shape[-1]).shape[0] < 32

    if not m_supported and deep_gemm_wrapper.ENABLE_JIT_DEEPGEMM:

Same swap-AB idea as CUTLASS, same M threshold neighbourhood, a different implementation, and a zero-token edge case that only appears once you run data-parallel attention. Both engines converge on "choose by M"; they differ on where that choice lives — vLLM compiles it into the dispatch chain, SGLang keeps it in Python where it can carry a paragraph of justification.

Llama-3-70B, GPTQ 4-bit group 128, no act-order, TP=2, on H100. The down_proj of one decoder layer: full shape $N = 8192$, $K = 28672$; per rank $K_{\text{part}} = 14336$.

1. AutoGPTQLinearMethod.create_weights (auto_gptq.py:L340-L358) builds an MPLinearLayerConfig from the partitioned shapes and calls choose_mp_linear_kernel (kernels/linear/__init__.py:L775-L837), which walks _POSSIBLE_KERNELS[CUDA]. CutlassW4A8LinearKernel declines (activations are bf16, not fp8).
2. MacheteLinearKernel.can_implement (mixed_precision/machete.py:L30-L64): cc is 90 ✓, weight type uint4b8 ✓, group 128 ✓, then check_machete_supports_shape(14336, 8192) — $14336 \bmod 64 = 0$ ✓, $8192 \bmod 128 = 0$ ✓. Machete wins.
3. process_weights_after_loading (machete.py:L70-L104) calls permute_param_layout_(x, input_dim=0, output_dim=1, packed_dim=0) and then ops.machete_prepack_B(...), which rewrites the packed int4 tensor into 128×64 prepacked blocks whose per-thread slices are 128-bit contiguous.
4. Decode step, M = 1. machete_mm enters the generated mm_dispatch_<type_sig> (template at generate.py:L45-L57), reads M = args.A.size(0), falls through every condition to the None default, and selects tile (128, 16) with cluster (1,1,1).
5. Inside the kernel: TMA loads the prepacked B block into shared memory; each thread pulls a 128-bit slice into registers; the upconvert runs (interleaved layout, since 4-bit and not converting to fp8/int8 — machete_prepacked_layout.cuh:L56-L59); wgmma with register-source A applies group scales during upconversion/inner reduction when they vary along K; the GEMM accumulates in fp32 and the epilogue stores bf16. Such group scales cannot be applied just once after the full K sum.
6. Prefill step, M = 2048. Same weights, same kernel object, different compiled schedule: "M > 256" matches first, giving tile (128, 256) cluster (2,1,1).

Now change TP to 3 for the sake of argument: $K_{\text{part}} = 28672/3$ is not an integer, but take a model where it is and the result is not divisible by 64 — Machete's shape check fails, the walk continues, and MarlinLinearKernel.can_implement accepts it because Marlin pads. You get a working server with a warning_once about tile padding and a slightly slower layer. Nothing in the logs says "you lost Machete"; the only signal is the one-off Using MarlinLinearKernel for AutoGPTQLinearMethod line at auto_gptq.py:L356-L358.

Both kernel families ship a microbenchmark. Marlin's sweeps batch sizes across a model's real layer shapes:

# Marlin: default batch sweep is 1,16,32,64,128,256,512,1024,2048,4096,8192
python benchmarks/kernels/benchmark_marlin.py --batch-sizes 1 16 64 256 1024

# Machete vs Marlin vs bf16 on real layer shapes (defaults: Llama-3-8b, Llama-2-70b)
python benchmarks/kernels/benchmark_machete.py model_bench --batch-sizes 1 16 32 64 128 256

# Force a kernel out of the running and see what the selector picks next
VLLM_DISABLED_KERNELS=MacheteLinearKernel vllm serve <gptq-model> --tensor-parallel-size 2

Three things worth measuring, none of which need a paper to justify:

1. Plot Marlin's speedup over bf16 against M and separately locate its own memory/compute ridge. The ideal $M^{*}\approx76$ is the W4A16 own-ridge, not its crossover with bf16. Actual knees and pairwise speedups depend on achieved bandwidth, compute, conversion and shapes; do not assume a fixed direction of shift.
2. Run the same model with VLLM_DISABLED_KERNELS=MacheteLinearKernel and without, at M=1 and M=256. Machete's advantage should be largest at moderate M where TMA and wgmma pay, and smallest at M=1 where both kernels are pure weight streaming.
3. Take a model whose hidden size is not a multiple of 128, load it under TP=2, and read the warning_once about tile padding. Compare per-token latency against the unpadded TP=1 case to price the padding.

• IST-DASLab/marlin — the original kernel and its README, cited in vLLM's header at marlin_template.h:L19. The README carries the authors' own batch-size sweep and the "batch sizes up to 16–32" framing that the M-bucketing inherits.
• vLLM PR #7174 — Machete: the Hopper-optimised mixed-precision GEMM. The discussion is the best available account of why prepacking to a wgmma layout beats Marlin's fragment layout on Hopper.
• csrc/libtorch_stable/quantization/machete/Readme.md — short, and the only place the out = (w_q.to(scale_type) * w_s - w_z.to(scale_type)) @ a contract and the FMA-ordering caveat on zero points are written down.
• csrc/libtorch_stable/quantization/marlin/dequant.h:L1-L64 — a genuinely good file-header essay on fusing zero-point subtraction and scaling into the bit-trick conversion, including which fusions are numerically safe and which are not.
• NVIDIA/cutlass epilogue visitor trees — the Sm90EVT composition used by ScaledEpilogue. vLLM's csrc/quantization/w8a8/cutlass/Epilogues.md documents the vLLM-specific ones (note the path: the markdown stayed under csrc/quantization/ when the sources moved to csrc/libtorch_stable/quantization/).
• deepseek-ai/DeepGEMM — SGLang's first-choice FP8 backend. Not readable from either checkout; see the Unverified callout in §4.3.9.