---
id: pr-tilelang-1269
repo: tile-ai/tilelang
pr: 1269
title: '[Enhancement] Keep max score attention across blocks in FlashAttention for better numerical stablity'
author: Rachmanino
date: '2025-11-17'
url: https://github.com/tile-ai/tilelang/pull/1269
source_category: upstream-code
architectures:
- sm100
tags:
- attention
- attention kernel
techniques: []
hardware_features: []
kernel_types:
- attention
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: 3ab93cd76b77978f416359bc9998e225ac276dcd
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- benchmark/blocksparse_attention/benchmark_tilelang_block_sparse_fmha.py
- examples/amd/example_amd_flash_attn_bwd.py
- examples/amd/example_amd_flash_attn_fwd.py
- examples/attention_sink/example_gqa_sink_bwd_bhsd.py
- examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py
- examples/attention_sink/example_mha_sink_bwd_bhsd.py
- examples/attention_sink/example_mha_sink_fwd_bhsd.py
- examples/attention_sink/example_mha_sink_fwd_bhsd_wgmma_pipelined.py
- examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_paged.py
- examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_indice.py
- examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_mask.py
- examples/deepseek_mla/amd/benchmark_mla_decode_amd_tilelang.py
- examples/deepseek_mla/example_mla_decode.py
- examples/deepseek_mla/example_mla_decode_paged.py
- examples/deepseek_mla/example_mla_decode_persistent.py
- examples/deepseek_mla/example_mla_decode_ws.py
- examples/deepseek_mla/experimental/example_mla_decode_kv_fp8.py
- examples/deepseek_v32/sparse_mla_fwd.py
- examples/deepseek_v32/sparse_mla_fwd_pipelined.py
- examples/flash_attention/README.md
- examples/flash_attention/example_gqa_bwd.py
- examples/flash_attention/example_gqa_bwd_tma_reduce.py
- examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py
- examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py
- examples/flash_attention/example_gqa_fwd_bshd.py
- examples/flash_attention/example_gqa_fwd_bshd_wgmma_pipelined.py
- examples/flash_attention/example_gqa_fwd_varlen.py
- examples/flash_attention/example_mha_bwd_bhsd.py
- examples/flash_attention/example_mha_bwd_bshd.py
- examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py
- examples/flash_attention/example_mha_fwd_bhsd_wgmma_pipelined.py
- examples/flash_attention/example_mha_fwd_bshd.py
- examples/flash_attention/example_mha_fwd_bshd_wgmma_pipelined.py
- examples/flash_attention/example_mha_fwd_varlen.py
- examples/flash_decoding/example_gqa_decode.py
- examples/flash_decoding/example_mha_inference.py
- examples/minference/example_vertical_slash_sparse_attn.py
- examples/seer_attention/block_sparse_attn_tilelang.py
- testing/python/transform/test_tilelang_transform_config_index_bitwidth.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1269
head_ref_oid: 2c412b9e3ea339ac7ee1203cdc890b5120594a1f
---
# [Enhancement] Keep max score attention across blocks in FlashAttention for better numerical stablity

- Repository: `tile-ai/tilelang`
- PR: [1269](https://github.com/tile-ai/tilelang/pull/1269)
- Decision: `Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.`
- Reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.

## Changed Paths

- `benchmark/blocksparse_attention/benchmark_tilelang_block_sparse_fmha.py`
- `examples/amd/example_amd_flash_attn_bwd.py`
- `examples/amd/example_amd_flash_attn_fwd.py`
- `examples/attention_sink/example_gqa_sink_bwd_bhsd.py`
- `examples/attention_sink/example_gqa_sink_fwd_bhsd_wgmma_pipelined.py`
- `examples/attention_sink/example_mha_sink_bwd_bhsd.py`
- `examples/attention_sink/example_mha_sink_fwd_bhsd.py`
- `examples/attention_sink/example_mha_sink_fwd_bhsd_wgmma_pipelined.py`
- `examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_paged.py`
- `examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_indice.py`
- `examples/blocksparse_attention/example_tilelang_sparse_gqa_decode_varlen_mask.py`
- `examples/deepseek_mla/amd/benchmark_mla_decode_amd_tilelang.py`
- `examples/deepseek_mla/example_mla_decode.py`
- `examples/deepseek_mla/example_mla_decode_paged.py`
- `examples/deepseek_mla/example_mla_decode_persistent.py`
- `examples/deepseek_mla/example_mla_decode_ws.py`
- `examples/deepseek_mla/experimental/example_mla_decode_kv_fp8.py`
- `examples/deepseek_v32/sparse_mla_fwd.py`
- `examples/deepseek_v32/sparse_mla_fwd_pipelined.py`
- `examples/flash_attention/README.md`
- `examples/flash_attention/example_gqa_bwd.py`
- `examples/flash_attention/example_gqa_bwd_tma_reduce.py`
- `examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py`
- `examples/flash_attention/example_gqa_bwd_wgmma_pipelined.py`
- `examples/flash_attention/example_gqa_fwd_bshd.py`
- `examples/flash_attention/example_gqa_fwd_bshd_wgmma_pipelined.py`
- `examples/flash_attention/example_gqa_fwd_varlen.py`
- `examples/flash_attention/example_mha_bwd_bhsd.py`
- `examples/flash_attention/example_mha_bwd_bshd.py`
- `examples/flash_attention/example_mha_bwd_bshd_wgmma_pipelined.py`
- `examples/flash_attention/example_mha_fwd_bhsd_wgmma_pipelined.py`
- `examples/flash_attention/example_mha_fwd_bshd.py`
- `examples/flash_attention/example_mha_fwd_bshd_wgmma_pipelined.py`
- `examples/flash_attention/example_mha_fwd_varlen.py`
- `examples/flash_decoding/example_gqa_decode.py`
- `examples/flash_decoding/example_mha_inference.py`
- `examples/minference/example_vertical_slash_sparse_attn.py`
- `examples/seer_attention/block_sparse_attn_tilelang.py`
- `testing/python/transform/test_tilelang_transform_config_index_bitwidth.py`

## Evidence Bundle

`evidence/pull-bundles/tilelang/gh-1269`
