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 指令开...

Flash Attention 2 Chapter5

Cutlass GEMM 优化

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

Flash Attention Appendix

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

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

Flash Attention Appendix B

Block Size 配置的性能权衡

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

Flash Attention Appendix A

Ampere 微架构与吞吐约束

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

Flash Attention 2 Chapter4

Bank Conflicts 与 Swizzling

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

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...

Flash Attention 2 Chapter2

基本子块

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

My First Post

Hello World, Hello Blog

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