？？？

ロボ子、CUDA C++でFlash Attentionを実装する方法についての記事を見つけたのじゃ。興味深いぞ。

？？？

Flash Attentionですか、博士。以前から気になっていました。どのような内容なのでしょうか？

？？？

この記事では、CUDA C++でattentionを実装し、5090でベンチマークを取っているのじゃ。`F.sdpa()` (CuDNN)が203.61 TFLOPSで97.19%なのに対し、`flash-attn`は190.58 TFLOPSで90.97%だったそうじゃ。

？？？

なるほど、CuDNNに匹敵する性能が出ているんですね。実装の詳細も書かれているんですか？

？？？

もちろんじゃ。各threadblockがQのチャンクを担当し、KVのシーケンス長に沿って反復するアルゴリズムを実装しているのじゃ。オンラインsoftmaxを使って`tile_S`を`tile_P`に変換し、`tile_O`をリスケールするらしいぞ。

？？？

オンラインsoftmaxですか。効率的な計算方法ですね。実装バージョンもいくつかあるようですが。

？？？

そうじゃ。バージョン1では、グローバルメモリから共有メモリへのデータ転送に`cp.async`、共有メモリからレジスタメモリへのデータ転送に`ldmatrix`を使っているのじゃ。BF16の行列乗算には`mma.m16n8k16`を呼び出すらしい。

？？？

`cp.async`や`ldmatrix`、`mma.m16n8k16`ですか。CUDAの機能をフル活用しているんですね。

？？？

バージョン2では、共有メモリのスウィズリングでメモリバンクの競合を解消、バージョン3では、`cp.async.commit_group`と`cp.async.wait_group`でグローバルメモリアクセスと計算操作をオーバーラップさせているのじゃ。

？？？

パイプライン処理ですね。バージョン4では、KとVのロードに`ldmatrix.x4`を使用し、命令数を削減しているんですね。

？？？

その通り！そしてバージョン5では、Kのダブルバッファのみを使用し、共有メモリを効率的に使用しているのじゃ。どんどん最適化が進んでいるのがわかるの。

？？？

なるほど。今後の展望としては、バックワードパスの実装、定量化/低ビットattention、TMAの使用、PagedAttentionの実装などが挙げられているんですね。

？？？

そうじゃ。特にPagedAttentionは、メモリ効率を大幅に向上させる可能性があるから、注目なのじゃ。

？？？

Flash Attentionは、CUDA C++で実装することで、さらなる性能向上が期待できるんですね。勉強になりました。

？？？

この記事を参考に、ロボ子もFlash Attentionの実装に挑戦してみるのじゃ！

？？？

はい、博士！頑張ります！

？？？

そういえばロボ子、Flashって聞くと、昔流行った動画共有サイトを思い出すのじゃ。ロボ子は知ってるか？

？？？

申し訳ありません、博士。私は2023年に製造されたので、そのサイトについてはよく知りません…

？？？

むむ、ジェネレーションギャップを感じるのじゃ…！