---
id: pr-cutlass-2466
repo: NVIDIA/cutlass
pr: 2466
title: Example 77 add blackwell fmha bwd for MLA shape
author: uchihatmtkinu
date: '2025-07-24'
url: https://github.com/NVIDIA/cutlass/pull/2466
source_category: upstream-code
architectures:
- sm100
tags:
- mla
- attention
- flash-attention
- fp8
techniques:
- warp-specialization
- double-buffering
- pipeline-stages
hardware_features:
- tcgen05
- tmem
- tma
- fp8
kernel_types:
- mla
- attention
- flash-attention
languages:
- cuda-cpp
- cute-dsl
captured_at: 2026-04-17
status: merged
merge_sha: 0e026982ce2ed10b27ec569c6e42035cb9118f62
inclusion_reason: manually curated, Blackwell kernel relevance
changed_paths:
- examples/77_blackwell_fmha/77_blackwell_fmha.cu
- examples/77_blackwell_fmha/77_blackwell_fmha_bwd.cu
- examples/77_blackwell_fmha/77_blackwell_mla.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/device/fmha_device_bwd.hpp
- examples/77_blackwell_fmha/device/sm100_mla.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_kernel_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_mla_kernel_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_mla_reduction.hpp
- examples/77_blackwell_fmha/kernel/sm100_fmha_mla_tma_warpspecialized.hpp
- examples/77_blackwell_fmha/kernel/sm100_mla_tile_scheduler.hpp
artifact_dir: evidence/pull-bundles/cutlass/gh-2466
head_ref_oid: ac13cf44c023c517b0916dcb2494712d5da3df44
---
# Example 77 add blackwell fmha bwd for MLA shape

- Repository: `NVIDIA/cutlass`
- PR: [2466](https://github.com/NVIDIA/cutlass/pull/2466)
- 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_fmha_bwd.cu`
- `examples/77_blackwell_fmha/77_blackwell_mla.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/device/fmha_device_bwd.hpp`
- `examples/77_blackwell_fmha/device/sm100_mla.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_kernel_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_bwd_mla_kernel_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_mla_reduction.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_fmha_mla_tma_warpspecialized.hpp`
- `examples/77_blackwell_fmha/kernel/sm100_mla_tile_scheduler.hpp`

## Evidence Bundle

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