---
id: pr-flash-attention-1361
repo: Dao-AILab/flash-attention
pr: 1361
title: Fix FA3 Varlen Performance regression
author: kadeng
date: '2024-12-02'
url: https://github.com/Dao-AILab/flash-attention/pull/1361
source_category: upstream-code
architectures:
- sm90
tags:
- FA3
- arch_pipeline
- attention_kv
- benchmark_test
- memory_primitives
- scheduler_autotune
- tma
- wgmma
techniques:
- async-copy
- autotuning
- pipeline-stages
- tile-scheduling
hardware_features:
- hopper
- tma
- wgmma
kernel_types:
- attention
- memory
- runtime
- scheduler
languages:
- cuda-cpp
- cute-dsl
- ptx
captured_at: '2026-05-16'
status: merged
merge_sha: 0823cf7b5d96499c1c79a4f64b1e256a035ba4b4
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- hopper/copy_paged_sm90_tma_cutlass35.hpp
- hopper/copy_paged_sm90_tma_cutlass36.hpp
- hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp
artifact_dir: evidence/pull-bundles/flash-attention/gh-1361
head_ref_oid: 39db2acb22991bee559ea50b635e7d24a4870814
---
# Fix FA3 Varlen Performance regression

- Repository: `Dao-AILab/flash-attention`
- PR: [1361](https://github.com/Dao-AILab/flash-attention/pull/1361)
- Decision: `Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.`
- Reason: Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.

## Changed Paths

- `hopper/copy_paged_sm90_tma_cutlass35.hpp`
- `hopper/copy_paged_sm90_tma_cutlass36.hpp`
- `hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp`

## Evidence Bundle

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