标题： SystolicAttention：在单个脉冲阵列中融合FlashAttention 显示英文标题

标题： SystolicAttention: Fusing FlashAttention within a Single Systolic Array

摘要： Transformer模型高度依赖缩放点积注意力（SDPA），通常使用FlashAttention算法实现。 然而，当前基于行波阵列的加速器在执行FlashAttention时面临重大挑战。 行波阵列只能在连续且大规模的矩阵乘法中实现高利用率。 相比之下，FlashAttention需要频繁交错的矩阵乘法和softmax操作。 行波阵列与外部向量单元之间的频繁数据交换导致行波阵列利用率低下。 此外，softmax涉及大量非矩阵操作，这使得其不太适合行波阵列，进一步加剧了这一问题。 而且，在行波阵列上同时执行矩阵乘法和在向量单元上执行softmax会导致寄存器文件和SRAM端口冲突，进一步降低性能。 为了克服这些限制，我们提出了FSA，一种增强的行波阵列架构，使整个FlashAttention算法能够在单个行波阵列内运行，无需外部向量单元。 FSA的核心是SystolicAttention，这是一种新的调度算法，能够以细粒度、逐元素重叠的方式将FlashAttention操作映射到行波阵列上。 这显著提高了阵列利用率，同时保留了原始的浮点运算顺序，以保持数值稳定性。 我们在可综合的RTL中实现了FSA，并将其性能与最先进的商用加速器进行了对比。 我们的结果表明，与AWS NeuronCore-v2和Google TPUv5e相比，FSA的注意力FLOPs/s利用率分别提高了1.77倍和4.83倍，仅增加了约10%的面积开销。 显示英文摘要

摘要： Transformer models rely heavily on scaled dot-product attention (SDPA), typically implemented using the FlashAttention algorithm. However, current systolic-array-based accelerators face significant challenges when executing FlashAttention. Systolic arrays can only achieve high utilization for consecutive and large matrix multiplications. In contrast, FlashAttention requires frequently interleaved matrix multiplications and softmax operations. The frequent data swaps between the systolic array and external vector units result in low systolic array utilization. This is further exacerbated by the fact that softmax involves numerous non-matrix operations, which are not well-suited for systolic arrays. Moreover, the concurrent execution of matrix multiplication on systolic arrays and softmax on vector units leads to register file and SRAM port contention, further degrading performance. To overcome these limitations, we propose FSA, an enhanced systolic array architecture that enables the entire FlashAttention algorithm to run entirely within a single systolic array, eliminating the need for external vector units. At the core of FSA is SystolicAttention, a novel scheduling algorithm that maps FlashAttention operations onto systolic arrays with fine-grained, element-wise overlap. This significantly improves array utilization while preserving the original floating-point operation order to maintain numerical stability. We implement FSA in synthesizable RTL and evaluate its performance against state-of-the-art commercial accelerators. Our results show that FSA achieves 1.77x and 4.83x higher attention FLOPs/s utilization compared to AWS NeuronCore-v2 and Google TPUv5e, respectively, with only about 10% area overhead.