AIシステム性能エンジニアリング (Japanese Edition)

第 10 章 カーネル内 パイプライン、ワープ特殊化、および 協調的なスレッドブロッククラスタリング

この作品はAIを使って翻訳されている。ご意見、ご感想をお待ちしている：[e-mail]

これまでの章では、メモリアクセス調整、並列性の最大化、計算とデータ転送の重複、占有率の向上、ワープのストール最小化などの基本的な最適化について取り上げてきました。これらは、遅延を隠し、ボトルネックを排除するのに役立ちました。しかし、最新の GPU は、基本的な最適化テクニックをさらに一歩進めることができる、高度なハードウェア機能と実行モデルを提供しています。

この章では、ワープ特殊化パイプライン、グリッドレベルおよびクラスタレベルの同期を備えた協調グループ、動的なワークキューをループする永続的なカーネル、分散共有メモリ (DSMEM または DSM) および Tensor Memory Accelerator (TMA) マルチキャストを使用するスレッドブロッククラスタ (別名、協調スレッド配列クラスタ[CTA]) など、より高度な CUDA テクニックをいくつか紹介します。スレッドブロッククラスタは、大まかに言えば、並行実行が保証されたスレッドブロック群です。これらはDSMEMを使用して、互いの共有メモリに対して読み取り、書き込み、原子性操作を実行できます。

これらのメソッドにより、ホストの介入なしにメモリアクセスと演算処理を並列化できます。また、スレッドブロック間でオンチップデータ共有が可能となり、各SMを完全に活用し続けられます。

これらの現代的なGPU実行モデルを理解することで、次の章に進む準備が整います。次の章では、CUDAストリームを用いたカーネル間パイプラインの探求を通じて、これらの最適化をさらに拡張します。次の章では、本章で説明したカーネル内最適化の基盤の上に、カーネル間パイプラインを構築します。

カーネル内パイプライン化テクニック

カーネル内パイプラインとは、 単一カーネル実行内でメモリ操作と計算を並列化する一連のテクニックを指します。（次章では、異なるストリームで実行される複数カーネル間で作業を並列化するカーネル間パイプラインを解説します。）

中核となる考え方は、カーネルを並列ステージに構造化し、あるデータがロードまたは保存されている間に、以前にロードされたデータを処理できるようにすることです。これらのステージは、異なるタイルまたはデータチャンク上で並行して動作します。これによりスループットが向上し、レイテンシを効率的に隠蔽します。

従来、GPUは遅延を隠蔽するためにワープレベルのマルチスレッディングに依存してきた。あるワープがメモリロードでストールしている間、他のワープは計算を継続する。これが実行モデルにおける単一命令・複数スレッド（SIMT）遅延隠蔽の基盤である。

カーネル内パイプライン処理は、同一ワープまたはカーネル内でメモリ操作と演算を重複させることでこれをさらに推進します。メモリロードと演算を時差させるために、時には単一ワープ内でも、きめ細かい調整を行います。

CUDA Pipeline APIによるカーネル内パイプライン処理では、__syncthreads()を一切使用せずに非同期メモリ転送と演算を並列化します。カーネル内パイプライン処理の一般的な手法は、ダブルバッファリングとワープ特殊化の2つです。

ダブルバッファリング（2段階）パイプライン手法では、全スレッドが均一に協調します。ワープ特化パイプライン手法では、ワープがメモリローダー、演算、メモリストーラーといった異なる役割に特化されます。選択はワークロードと性能要件に依存します。 ...