---
id: pr-flash-attention-2360
repo: Dao-AILab/flash-attention
pr: 2360
title: '[Fwd,Sm90] Add paged KV attention support (tma and cp.async)'
author: henrylhtsang
date: '2026-03-18'
url: https://github.com/Dao-AILab/flash-attention/pull/2360
source_category: upstream-code
architectures:
- sm100
- sm90
tags:
- arch_pipeline
- attention_kv
- paged
- sm100
- sm90
- tma
techniques:
- async-copy
- cache-policy
- pipeline-stages
hardware_features:
- blackwell
- hopper
- tma
kernel_types:
- attention
- runtime
languages:
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: 0372975669168fbc044507663e4600e2edee0d13
inclusion_reason: Preserve page/KV/layout and plan-run contracts; profile memory traffic separately from math utilization.
changed_paths:
- flash_attn/cute/block_sparse_utils.py
- flash_attn/cute/flash_fwd_sm90.py
- flash_attn/cute/interface.py
- flash_attn/cute/paged_kv.py
- tests/cute/test_flash_attn.py
artifact_dir: evidence/pull-bundles/flash-attention/gh-2360
head_ref_oid: d2f6bc2cff45fc0536f53bcc65666c368319ac7e
---
# [Fwd,Sm90] Add paged KV attention support (tma and cp.async)

- Repository: `Dao-AILab/flash-attention`
- PR: [2360](https://github.com/Dao-AILab/flash-attention/pull/2360)
- Decision: `Preserve page/KV/layout and plan-run contracts; profile memory traffic separately from math utilization.`
- Reason: Preserve page/KV/layout and plan-run contracts; profile memory traffic separately from math utilization.

## Changed Paths

- `flash_attn/cute/block_sparse_utils.py`
- `flash_attn/cute/flash_fwd_sm90.py`
- `flash_attn/cute/interface.py`
- `flash_attn/cute/paged_kv.py`
- `tests/cute/test_flash_attn.py`

## Evidence Bundle

`evidence/pull-bundles/flash-attention/gh-2360`
