論文の概要: SystolicAttention: Fusing FlashAttention within a Single Systolic Array
- arxiv url: http://arxiv.org/abs/2507.11331v2
- Date: Wed, 16 Jul 2025 13:36:16 GMT
- ステータス: 翻訳完了
- システム内更新日: 2025-07-17 12:30:11.791448
- Title: SystolicAttention: Fusing FlashAttention within a Single Systolic Array
- Title(参考訳): SystolicAttention:単一のSystolicアレー内でFlashAttentionを使用する
- Authors: Jiawei Lin, Guokai Chen, Yuanlong Li, Thomas Bourgeat,
- Abstract要約: Transformer Modelは、SDPA(Scaled dot-product attention)に大きく依存している。
現在のsystolic-arrayベースのアクセラレータは、FlashAttentionの実行において重大な課題に直面している。
我々は、FlashAttentionアルゴリズム全体を単一のsystolic配列内で完全に動作させることができる拡張されたsystolic配列アーキテクチャであるFSAを提案する。
- 参考スコア(独自算出の注目度): 2.8650887057567864
- License: http://creativecommons.org/licenses/by/4.0/
- Abstract: 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.
- Abstract(参考訳): トランスフォーマーモデルは、通常FlashAttentionアルゴリズムを用いて実装されるSDPA (Scaled dot-product attention) に大きく依存している。
しかし、現在のsystolic-arrayベースのアクセラレータは、FlashAttentionの実行において重大な課題に直面している。
シストリックアレイは連続行列乗算と大規模行列乗算の高利用しか達成できない。
対照的に、FlashAttentionは頻繁にインターリーブされた行列乗算とソフトマックス演算を必要とする。
サイストリックアレイと外部ベクトルユニット間の頻繁なデータスワップは、サイストリックアレイの利用率を低くする。
これは、ソフトマックスが多くの非行列演算を伴っているという事実によりさらに悪化する。
さらに,シストリックアレイ上の行列乗算とベクトル単位上のソフトマックスの同時実行により,レジスタファイルとSRAMポート競合が発生し,さらに性能が低下する。
これらの制限を克服するため,FlashAttention アルゴリズム全体が単一のsystolic 配列内で完全に動作できるように拡張されたsystolic 配列アーキテクチャである FSA を提案する。
FSAのコアとなるSystolicAttentionは、FlashAttention操作を細かな要素レベルでオーバーラップしたsystolic配列にマッピングする新しいスケジューリングアルゴリズムである。
これにより、数値安定性を維持するために元の浮動小数点演算を保存しながら、配列利用を大幅に改善する。
我々は、合成可能なRTLにFSAを実装し、その性能を最先端の商用アクセラレーターに対して評価する。
その結果,AWS NeuronCore-v2 と Google TPUv5e と比較して FSA は FLOPs/s の利用率を 1.77x と 4.83x に向上し,領域オーバーヘッドは 10% に過ぎなかった。
関連論文リスト
- Orthogonal Finetuning Made Scalable [87.49040247077389]
OFT(Orthogonal Finetuning)は、壊滅的な忘れ込みを防止しつつ、パラメータ効率の高い適応を提供するが、実行時とメモリの要求が高いため、実際のデプロイメントが制限される。
ここでは,OFTの計算ボトルネックを重み中心の実装とみなす。
本稿では,行列ベクトル乗法(行列フリー計算)を用いて,計算コストを2次に削減する入力中心の変換法OFTv2を提案する。
これらの修正により、OFTv2はパフォーマンスを損なうことなく、最大10倍高速なトレーニングと3倍のGPUメモリ使用率を達成することができる。
論文 参考訳(メタデータ) (2025-06-24T17:59:49Z) - FLASH-D: FlashAttention with Hidden Softmax Division [3.668018928502405]
FlashAttentionはオンラインのソフトマックス計算に基づいており、ソフトマックス計算と行列演算を統合している。
この研究は、FLASH-Dを数学的に等価だが単純化した定式化として、 (a) ソフトマックス分割を他の非線形関数評価内に隠蔽し、 (b) 本質的に指数関数の数値的に安定な計算を行い、 (c) FlashAttentionカーネルに数値近似を導入することなく計算コストを削減した。
論文 参考訳(メタデータ) (2025-05-20T11:01:33Z) - Periodic Online Testing for Sparse Systolic Tensor Arrays [0.0]
モダン機械学習(ML)アプリケーションは、しばしば構造化されたスパーシティの恩恵を受ける。これは、モデルの複雑さを効率的に低減し、ハードウェア内のスパースデータの処理を単純化するテクニックである。
本稿では,ベクトルの開始前にスパルス・シストリック・テンソルアレイ内の永久断層を検出し,検出するオンラインエラーチェック手法を提案する。
論文 参考訳(メタデータ) (2025-04-25T18:10:45Z) - Tiled Flash Linear Attention: More Efficient Linear RNN and xLSTM Kernels [14.756974816917584]
ゲーティングを備えた線形RNNは、最近、言語モデリングにおけるトランスフォーマーと比較して、競合する性能を示した。
線形RNNのための新しいカーネルアルゴリズムである Tiled Flash Linear Attention (TFLA) を提案する。
高速化ベンチマークでは、TFLAに基づく新しいmLSTMカーネルが、高度に最適化されたFlashアテンション、線形アテンション、およびMambaカーネルより優れていることを示す。
論文 参考訳(メタデータ) (2025-03-18T16:09:47Z) - Sliding Window Attention Training for Efficient Large Language Models [55.56483740523027]
SWATを導入し,スライディングウインドウ・アテンション・トレーニング(Sliding Window Attention Training)を用いて,より効率的な長文処理を実現する。
本稿ではまず,変圧器の非効率性を注目シンク現象とみなす。
我々は、ソフトマックスをシグモイド関数に置き換え、効率的な情報圧縮と保持のためにバランスの取れたALiBiとRotary Position Embeddingを利用する。
論文 参考訳(メタデータ) (2025-02-26T05:31:44Z) - FlashSparse: Minimizing Computation Redundancy for Fast Sparse Matrix Multiplications on Tensor Cores [6.404201720333765]
我々は、スパースワークロードとTCUアーキテクチャのギャップを埋める新しいアプローチであるFlashSparseを提案する。
特に、FlashSparseは、新しいスワップ・アンド・トランスポーション行列乗算戦略により、TCUs上のSpMMとSDDMMのスパース粒度を最小化する。
我々はFlashSparseがスパース行列乗算のための新しい最先端技術(幾何学平均はDTC-SpMMより5.5倍、RoDeより3.22倍)をセットしていることを示す。
論文 参考訳(メタデータ) (2024-12-15T01:12:33Z) - EoRA: Fine-tuning-free Compensation for Compressed LLM with Eigenspace Low-Rank Approximation [84.70637613266835]
EoRAは、圧縮されたLarge Language Modelを低ランク行列で拡張する微調整不要な手法である。
EoRAは、圧縮LDMの精度を回復するために、トレーニングなしの低ランク法よりも一貫して優れている。
論文 参考訳(メタデータ) (2024-10-28T17:59:03Z) - Fast Matrix Multiplications for Lookup Table-Quantized LLMs [58.11584672945781]
FLUTEはLUT量子化LLM用のフレキシブルなルックアップテーブルエンジンである。
バッチサイズ32と量子化グループサイズ128では、FLUTEカーネルは既存のGEMMカーネルよりも2〜4倍高速である。
論文 参考訳(メタデータ) (2024-07-15T17:55:42Z) - HiRE: High Recall Approximate Top-$k$ Estimation for Efficient LLM
Inference [68.59839755875252]
HiREは2つの新しいコンポーネントから構成される: (i) (i) (i) (i) (i) (i) (i) (i) (i) (i) (ii) DA-TOP-$k$: 効率的なマルチデバイス近似トップ-k$演算子) (i) (i) (i) (i) (i) (i) (i) DA-TOP-$k$演算子) 。
我々は、10億のパラメータモデルにおいて、HiREがソフトマックスとフィードフォワード層の両方に適用され、ほぼ一致した事前学習と下流の精度を実現し、1台のTPUv5eデバイスで1.47Times$の推論遅延を高速化することを示した。
論文 参考訳(メタデータ) (2024-02-14T18:04:36Z) - Adaptable Butterfly Accelerator for Attention-based NNs via Hardware and
Algorithm Co-design [66.39546326221176]
多くのAIタスクにおいて、注意に基づくニューラルネットワークが普及している。
注意機構とフィードフォワードネットワーク(FFN)の使用は、過剰な計算とメモリ資源を必要とする。
本稿では,注目機構とFFNの両方を近似するために,バタフライの分散パターンを統一したハードウェアフレンドリーな変種を提案する。
論文 参考訳(メタデータ) (2022-09-20T09:28:26Z) - RASA: Efficient Register-Aware Systolic Array Matrix Engine for CPU [6.436294460697506]
RASA, Register-Aware Systolic Arrayを提案する。
我々は,実行段階を複数のサブステージに分割し,命令を重複させてオーバーヘッドを隠蔽し,同時に実行する手法を開発した。
RASAをベースとした設計では、無視できる面積と電力オーバーヘッドで性能が大幅に向上した。
論文 参考訳(メタデータ) (2021-10-05T00:01:31Z) - FlexSA: Flexible Systolic Array Architecture for Efficient Pruned DNN
Model Training [1.718730454558804]
一般的なトレーニングアクセラレーターを大きなシストリックアレイでプルーニングすることは、非常に性能的に非効率であることがわかった。
本研究では,シストリックアレイを効率的なプルーニングとトレーニングを行うために,フレキシブルなシストリックアレイアーキテクチャであるFlexSAを提案する。
また、FlexSAのリソースを最大限活用するために、トレーニング作業負荷における行列乗算および累積演算のコンパイルを提案する。
論文 参考訳(メタデータ) (2020-04-27T15:51:20Z)
関連論文リストは本サイト内にある論文のタイトル・アブストラクトから自動的に作成しています。
指定された論文の情報です。
本サイトの運営者は本サイト(すべての情報・翻訳含む)の品質を保証せず、本サイト(すべての情報・翻訳含む)を使用して発生したあらゆる結果について一切の責任を負いません。