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

> Learn how to implement MQA logits computation with DeepGEMM. Discover optimized CUDA kernels leveraging TMA, UMMA, and pipelining on NVIDIA SM100. Accelerate your AI models.

- Repository: [DeepSeek/DeepGEMM](https://github.com/deepseek-ai/DeepGEMM)
- Tags: how-to-guide
- Published: 2026-04-19

---

**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`.

```python
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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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.

```python
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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/utils/math.py) – Utility functions for building TMA descriptors and helper math operations used by the Python wrappers.

- [`tests/test_attention.py`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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.