XiaosaYin Blog

Thinking will not overcome fear but action will.

Flash Attention 2 Chapter6

FP 指令融合与 Auto-Tuning

简介 在 Chapter5 中，我们已经完成了 Cutlass 风格的三项核心优化（eager loading、fragment interleaving、double buffering），在 RTX 3090 上达到接近 reference 的性能。 这一章继续做最后两步： Kernel 6：FP instruction fusion（优化 softmax 的 FP32 指令开...

Posted by BY on April 18, 2026

Flash Attention 2 Chapter5

Cutlass GEMM 优化

简介 在 Chapter4 中，我们通过 swizzling 消除了主要的 SMEM bank conflicts，性能已经接近 reference。 这一章进入 Cutlass 风格的 GEMM 优化阶段，重点是进一步提升 overlap 与可扩展性。 本章实现 3 组渐进优化： Kernel 3：GMEM -> SMEM eager loading（double buf...

Posted by BY on April 18, 2026

Flash Attention Appendix

实验、配置与指令补充说明

说明 本文是 Flash Attention from Scratch 系列的附录整理版，主要聚焦： benchmarking / profiling 的实验设置 arithmetic intensity 的计算脚本 nvcc 与寄存器 spill 相关编译参数 Ampere 设备关键规格与吞吐数据 mma、cp.async、wmma 相关指令级补充 原文参考：...

Posted by BY on April 18, 2026

Flash Attention Appendix B

Block Size 配置的性能权衡

简介 本附录讨论一个在实战中非常关键但常被“经验化”的问题： block size 配置到底如何影响 instruction pattern 与最终性能？ 分析对象使用 kernel 16（最终优化版本），目的是尽量减少代码生成噪声，让结论更聚焦在本质 tradeoff 上。 关注四个维度： Q 是 persist 在 RF，还是每轮 reload B_r（query ti...

Posted by BY on April 18, 2026

Flash Attention Appendix A

Ampere 微架构与吞吐约束

简介 本附录聚焦 Ampere 微架构层面的执行约束，解释为什么同样的 kernel 在 A100 与 RTX 3090 上会表现出完全不同的瓶颈形态。 核心关注点： SM / sub-core 级执行单元如何竞争资源 INT 与 FP32 pipeline 的发射与交错规则 mma 与 FP32 吞吐比（16x vs 2x）如何放大指令开销 原文参考：Appendix...

Posted by BY on April 18, 2026

Flash Attention 2 Chapter4

Bank Conflicts 与 Swizzling

简介 在 Chapter3（Kernel 1）中，我们已经拿到了一个可运行的 baseline kernel，但性能仅约为参考实现的一半。 这一章的核心目标是定位瓶颈并解决它：bank conflicts。 通过 Nsight Compute 分析可以看到，Kernel 1 在 SMEM 上存在严重冲突，导致大量访存串行化。 本章将引入并实现 swizzling，最终把性能从 33.28...

Posted by BY on April 17, 2026

Flash Attention 2 Chapter3

Kernel 1 基础实现

简介 在 Part 2 里，我们已经讨论了 Flash Attention 的基础 CUDA building blocks：Tensor Core mma 与高效内存传输（cp.async、ldmatrix）。这一章进入“组装阶段”：把这些低层组件拼成第一个完整、可运行的 Flash Attention kernel（Kernel 1）。 我们会按 3 个步骤推进： 先确定 C...

Posted by BY on April 17, 2026

Flash Attention 2 Chapter2

基本子块

简介 在这一章, 我们将会去探索 CUDA 中组成 Falsh Attention 核的基础操作.Flash Attention 的性能取决于两点：其一，通过 Tensor Core 充分提升计算吞吐；其二，通过高效的数据搬运降低内存瓶颈。 为了说明这些因素为何关键，先来看一个典型 attention 切片的算术复杂度：设有 64 个 query 向量，关注 4096 个 key-val...

Posted by BY on April 15, 2026

My First Post

Hello World, Hello Blog

Hey 这是我的第一篇博客。 简单试一下, 之后我就要开始我的 flash attention 的奇妙旅程了!

Posted by BY on April 10, 2026

FEATURED TAGS
Cuda, Flash Attention
ABOUT ME

Goals determine what you going to be!

✉️ xiaosayin@gmail.com

FRIENDS