在 CUDA C++ 中实现 Flash Attention 的历险记

本文详细介绍了作者如何在 CUDA C++ 中实现 Flash Attention，并对其性能进行了优化。作者从一个基础版本开始，逐步改进，使用了共享内存交换、两阶段流水线、以及更有效的 ldmatrix 使用等技术。通过不断地性能分析和优化，最终实现了接近硬件理论极限的性能。文章还深入探讨了在线 softmax 的实现细节以及共享内存冲突的解决方法，对于 CUDA C++ 开发者具有很高的参考价值。