How to Implement MQA (Multi-Query Attention) Logits Computation with DeepGEMM

DeepGEMM provides highly-optimized CUDA kernels for MQA logits computation using Tensor-Memory Access (TMA), Unified Matrix-Multiply-Accumulate (UMMA) instructions, and multi-stage pipelining on NVIDIA SM100 architecture.

The deepseek-ai/DeepGEMM library implements high-throughput Multi-Query Attention (MQA) logits computation for large language models like DeepSeek-V3. This article explains how to implement MQA logits computation with DeepGEMM using both standard dense and paged KV cache layouts.

Core Architecture Components

DeepGEMM implements MQA logits computation through several specialized CUDA components that work together to maximize SM100 utilization:

• Tensor-Memory Access (TMA) – Streams query (q), key-value cache (kv), scaling factors, and per-head weights from global memory to shared memory in a single transaction. This prefetching occurs in sm100_fp8_mqa_logits.cuh and sm100_fp4_mqa_logits.cuh at lines 54-61.

• Shared-Memory Buffers – Holds Q, KV, scaling factors, and per-head weights aligned to 512 bytes for the SM100 "swizzle-64B" pattern. The buffer is declared as extern __shared__ __align__(512) uint8_t smem_buffer[] at lines 71-76 in the same files.

• Pipeline Stages – Three pipelines (Q, KV, TMEM) overlap loads, matrix-multiply-accumulate (UMMA), and reduction operations. Stage creation and barrier management use cutlass::arch::ClusterTransactionBarrier at lines 18-30 (SM100) and barrier initialization at lines 12-27 (FP8).

• Unified Matrix-Multiply-Accumulate (UMMA) – Uses SM100 UMMA instructions to multiply Q (FP8/E4M3) with KV (FP8/E4M3) while applying per-head scaling. The UMMA descriptor is built in make_umma_desc at lines 61-70 (FP8) and lines 49-58 (FP4) via mma::sm100::make_umma_desc.

• Math Warp Reduction – After UMMA, the math warp loads partial results from TMEM, applies ReLU, multiplies by per-head weights, and reduces across heads to produce scalar logits. This logic resides at lines 94-124 (FP8) and lines 120-165 (FP4).

• Compressed Logits – When kIsCompressedLogits is true, the output stores only the valid range [seq_k_start, seq_k_end) for each query token, saving memory. This code path is guarded by if constexpr (kIsCompressedLogits) at lines 136-140 (FP8) and lines 130-136 (FP4).

Implementing Standard MQA Logits Computation

The Python API exposes the C++ kernels through the deep_gemm module. The primary entry point for non-paged (dense) MQA computation is deep_gemm.fp8_fp4_mqa_logits.

import torch
import deep_gemm as dg

# -------------------------------------------------

# 1️⃣  Fake data (replace with real model tensors)

# -------------------------------------------------

seq_len      = 2048          # tokens in the query sequence

seq_len_kv   = 4096          # tokens in the KV cache

num_heads    = 32
head_dim     = 64            # FP8/E4M3 uses 64-dim per head

block_q      = 128           # BLOCK_Q from the kernel config

block_kv     = 128           # BLOCK_KV from the kernel config

# Q: [seq_len, num_heads, head_dim]  (E4M3)

q = torch.randn(seq_len, num_heads, head_dim, dtype=torch.float8_e4m3fn, device='cuda')

# KV cache: [seq_len_kv, head_dim] (E4M3)

kv = torch.randn(seq_len_kv, head_dim, dtype=torch.float8_e4m3fn, device='cuda')

# Per-head weights: [seq_len, num_heads] (float32)

weights = torch.randn(seq_len, num_heads, dtype=torch.float32, device='cuda')

# Optional compressed-logits bookkeeping

cu_seq_len_k_start = torch.arange(seq_len, device='cuda', dtype=torch.int32)
cu_seq_len_k_end   = torch.full((seq_len,), seq_len_kv, dtype=torch.int32, device='cuda')

# Output tensor: logits shape = [seq_len, seq_len_kv] (float16)

logits = torch.empty(seq_len, seq_len_kv, dtype=torch.float16, device='cuda')

# -------------------------------------------------

# 2️⃣  Launch the kernel

# -------------------------------------------------

dg.fp8_fp4_mqa_logits(
    seq_len=seq_len,
    seq_len_kv=seq_len_kv,
    max_seqlen_k=seq_len_kv,         # max KV length (used internally)

    stride_logits=seq_len_kv,        # contiguous stride in the output

    cu_seq_len_k_start=cu_seq_len_k_start,
    cu_seq_len_k_end=cu_seq_len_k_end,
    logits=logits,
    tensor_map_q=dg.utils.make_tma_descriptor(q),          # helper in utils

    tensor_map_kv=dg.utils.make_tma_descriptor(kv),
    tensor_map_kv_scales=dg.utils.make_tma_descriptor(
        torch.ones(seq_len_kv, dtype=torch.float32, device='cuda')   # dummy scale

    ),
    tensor_map_weights=dg.utils.make_tma_descriptor(weights),
    # optional: compressed logits flag

    is_compressed_logits=False,
    clean_logits=True,               # zero-fill unused positions

)

# `logits` now contains the MQA scores (scaled ReLU + per-head weighting)

print(logits.shape)   # torch.Size([2048, 4096])

Key implementation details:

• make_tma_descriptor is a utility that builds a cute::TmaDescriptor from a PyTorch tensor. The wrapper resides in deep_gemm/utils/math.py at line 84.
• The kernel arguments map directly to the C++ signature defined in sm100_fp8_mqa_logits.cuh.
• Setting clean_logits=True forces the kernel to write -inf to padding regions when using compressed logits mode.

Implementing Paged MQA Logits Computation for Decoding

For inference with paged KV caches, DeepGEMM provides fp8_fp4_paged_mqa_logits and a metadata helper to manage block-table translations.

import torch
import deep_gemm as dg

# ----- 1️⃣ Build the KV cache in paged layout (simplified) -----

batch_size = 1
max_len    = 4096
block_kv   = 128                     # must match the kernel config

num_blocks = max_len // block_kv

# Random KV values and per-token scales (float32)

kv_cache   = torch.randn(batch_size, max_len, 64, dtype=torch.float8_e4m3fn, device='cuda')
kv_scales  = torch.ones(batch_size, max_len, dtype=torch.float32, device='cuda')

# ----- 2️⃣ Metadata for the *next* decoding step -----

# `context_lens` = lengths of each sequence in the batch (here all same)

context_lens = torch.tensor([1024], dtype=torch.int32, device='cuda')
metadata = dg.get_paged_mqa_logits_metadata(
    context_lens_nextn=context_lens,   # length after adding the new token(s)

    block_kv=block_kv,
    num_sms=dg.get_num_sms(),          # runtime-detected SM count

)

# ----- 3️⃣ Call the paged kernel -----

logits = torch.empty(1024, max_len, dtype=torch.float16, device='cuda')
dg.fp8_fp4_paged_mqa_logits(
    seq_len=1024,
    seq_len_kv=max_len,
    max_seqlen_k=max_len,
    stride_logits=max_len,
    cu_seq_len_k_start=context_lens,   # same as context lens for the simple case

    cu_seq_len_k_end=context_lens,
    logits=logits,
    tensor_map_q=dg.utils.make_tma_descriptor(q),
    tensor_map_kv=dg.utils.make_tma_descriptor(kv_cache),
    tensor_map_kv_scales=dg.utils.make_tma_descriptor(kv_scales),
    tensor_map_weights=dg.utils.make_tma_descriptor(weights),
    metadata=metadata,
    clean_logits=True,
)

print(logits.shape)   # torch.Size([1024, 4096])

Key differences from the standard path:

• get_paged_mqa_logits_metadata (exposed in deep_gemm/__init__.py at line 64) builds a struct that translates block-table indices into physical KV page addresses.
• The paged kernel reuses the same compute pipeline as the non-paged version but first looks up the correct KV slice via the block table metadata.

Key Source Files

The implementation spans several files that define the complete stack from low-level CUDA to Python wrappers:

• deep_gemm/include/deep_gemm/impls/sm100_fp8_mqa_logits.cuh – Main FP8 + FP4 MQA kernel (SM100) implementing prefetch, TMA, UMMA, and reduction logic.

• deep_gemm/include/deep_gemm/impls/sm100_fp4_mqa_logits.cuh – FP4-only variant for KV caches stored in FP4 format.

• deep_gemm/__init__.py – Python-level API exposing fp8_fp4_mqa_logits, fp8_fp4_paged_mqa_logits, and metadata helpers.

• deep_gemm/utils/math.py – Utility functions for building TMA descriptors and helper math operations used by the Python wrappers.

• tests/test_attention.py – Reference implementation (ref_fp8_mqa_logits) and exhaustive tests that verify numerical correctness and benchmark performance.

Summary

• DeepGEMM accelerates MQA logits computation through custom SM100 CUDA kernels that combine TMA prefetching, UMMA matrix operations, and pipelined execution.
• Two primary entry points exist: fp8_fp4_mqa_logits for dense KV caches and fp8_fp4_paged_mqa_logits for paged decoding scenarios.
• TMA descriptors constructed via dg.utils.make_tma_descriptor are required to map PyTorch tensors to the kernel's shared memory layout.
• Compressed logits mode reduces memory usage by writing only valid token ranges [seq_k_start, seq_k_end) rather than full attention matrices.
• The implementation resides primarily in sm100_fp8_mqa_logits.cuh and sm100_fp4_mqa_logits.cuh, with Python bindings in deep_gemm/__init__.py.

Frequently Asked Questions

What hardware is required to run DeepGEMM MQA kernels?

DeepGEMM's MQA kernels are optimized for NVIDIA's SM100 architecture (Blackwell generation). The code utilizes SM100-specific instructions including UMMA (Unified Matrix-Multiply-Accumulate) and TMA (Tensor-Memory Access), which require compatible hardware to execute.

How does DeepGEMM handle quantization scales in MQA computation?

The kernels accept per-token scaling factors through the tensor_map_kv_scales parameter. During the UMMA computation (lines 61-70 in sm100_fp8_mqa_logits.cuh), these scales are applied to the FP8/E4M3 Q and KV matrices to maintain numerical precision throughout the attention computation.

What is the difference between compressed and standard logits output?

Standard mode writes the full [seq_len, seq_len_kv] logits matrix. Compressed mode, enabled by setting is_compressed_logits=True, stores only the valid range [seq_k_start, seq_k_end) for each query token (lines 136-140 in sm100_fp8_mqa_logits.cuh). This reduces memory bandwidth and storage when the full attention matrix is not required.

Can DeepGEMM MQA kernels be used with Grouped-Query Attention (GQA)?

While the current implementation specifically targets Multi-Query Attention (single KV head per query group), the same kernel architecture can theoretically support GQA by adjusting the num_heads parameter and ensuring the KV cache layout matches the grouped-head structure. The fp8_fp4_mqa_logits function accepts arbitrary head counts through the tensor shapes, but you must ensure the KV data is properly broadcast or repeated for each query group.

{
  "mcpServers": {
    "instagit": {
      "command": "npx",
      "args": ["-y", "instagit@latest"]
    }
  }
}