---
id: pr-cutlass-2472
repo: NVIDIA/cutlass
pr: 2472
title: 'Add Blackwell MLA forward (shape: d=192, dv=128) implementation'
author: dianzhangchen
date: '2025-07-18'
url: https://github.com/NVIDIA/cutlass/pull/2472
source_category: upstream-code
architectures:
- sm100
tags:
- mla
- attention
- prefill
- fp8
techniques:
- warp-specialization
- pipeline-stages
- double-buffering
- tma-multicast
hardware_features:
- tcgen05
- tmem
- tma
- fp8
kernel_types:
- mla
- attention
- prefill
languages:
- cuda-cpp
- cute-dsl
captured_at: 2026-04-17
status: merged
merge_sha: 9baa06dd57804ce8fb5efe9e471b3451341522c6
inclusion_reason: manually curated, Blackwell kernel relevance
changed_paths:
- examples/77_blackwell_fmha/77_blackwell_fmha.cu
- examples/77_blackwell_fmha/77_blackwell_mla_fwd.cu
- examples/77_blackwell_fmha/CMakeLists.txt
- examples/77_blackwell_fmha/collective/fmha_fusion.hpp
- examples/77_blackwell_fmha/collective/sm100_fmha_fwd_epilogue_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/collective/sm100_fmha_fwd_mainloop_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/collective/sm100_fmha_mla_fwd_mainloop_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/collective/sm100_fmha_mla_load_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/common/pipeline_mla.hpp
- examples/77_blackwell_fmha/kernel/fmha_causal_tile_scheduler.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_kernel_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_fwd_kernel_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/reference/fmha_fwd_reference.hpp
artifact_dir: evidence/pull-bundles/cutlass/gh-2472
head_ref_oid: 7f48538bd1a3a13cde9f2282fe4d5528f2db997c
---
# Add Blackwell MLA forward (shape: d=192, dv=128) implementation

- Repository: `NVIDIA/cutlass`
- PR: [2472](https://github.com/NVIDIA/cutlass/pull/2472)
- Decision: `manually curated, Blackwell kernel relevance`
- Reason: manually curated, Blackwell kernel relevance

## Changed Paths

- `examples/77_blackwell_fmha/77_blackwell_fmha.cu`
- `examples/77_blackwell_fmha/77_blackwell_mla_fwd.cu`
- `examples/77_blackwell_fmha/CMakeLists.txt`
- `examples/77_blackwell_fmha/collective/fmha_fusion.hpp`
- `examples/77_blackwell_fmha/collective/sm100_fmha_fwd_epilogue_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/collective/sm100_fmha_fwd_mainloop_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/collective/sm100_fmha_mla_fwd_mainloop_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/collective/sm100_fmha_mla_load_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/common/pipeline_mla.hpp`
- `examples/77_blackwell_fmha/kernel/fmha_causal_tile_scheduler.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_kernel_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_fwd_kernel_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/reference/fmha_fwd_reference.hpp`

## Evidence Bundle

`evidence/pull-bundles/cutlass/gh-2472`
