ワープダイバージェンスをなくす
インデックスを組み替えてワープを効率よく動かします。
「ワープダイバージェンスをなくす」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン2/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
warpは足並みをそろえて実行する
warpは、同じ命令を一緒に実行する32個のスレッドです。処理の分岐が一致していれば、ハードウェアは最大速度で実行できます。
ダイバージェンスのコスト
warp内のスレッドが異なる分岐へ進むと、それはdivergenceです。ハードウェアは各経路を順番に実行するため、一部のレーンがアイドルになり、サイクルが無駄になります。
素朴なリダクションでは分岐が起きる
単純な実装ではtid % (2*s)を使ってアクティブなスレッドを選びます。各warp内でアクティブなスレッドとアイドルなスレッドが交互に並ぶため、warpで大きな分岐が発生します。
if (tid % (2 * s) == 0)
data[tid] += data[tid + s];アイドルなレーンにもコストがかかる
半分のスレッドが何もしなくても、それらはwarpを占有し続けます。アクティブな経路とアイドルな経路の両方を処理するまで、warpは完了できません。
スレッドIDで再インデックスする
解決策は、分散したスレッドではなく最小のスレッドIDにアクティブな処理を割り当てることです。tidとstrideからインデックスを計算します。
int index = 2 * s * tid;
if (index < blockDim.x)
data[index] += data[index + s];それで改善する理由
これで処理中のスレッドが連続します。tid 0、1、2、…がすべて処理し、それ以外はすべて待機します。warp全体がアクティブか、全体がアイドルになります。
完全にアイドルなwarpは無料
すべてのレーンがアイドルなwarpは、処理なしで終了します。レーンごとの逐次実行がないため、ダイバージェンスのコストはほぼなくなります。
剰余演算の罠
隠れた原因はmodulo条件でした。これによって各warp内にアクティブなスレッドが分散し、まさにダイバージェンスが発生していました。
同じ処理、よりよい割り当て
計算内容も加算回数も変えていません。各加算をどのスレッドが担当するかを割り当て直しただけで、warpの実行効率が向上します。
規模が大きいほど効果が積み重なる
数千のブロックと多数のステップにわたってダイバージェンスを取り除くと、実際のspeedupにつながります。素朴なカーネルより数倍高速になることもあります。
まだ1つ問題が残っている
この実装では、shared memory上でインターリーブされた位置にある隣接要素を読み取るため、バンクコンフリクトが起こる可能性があります。次のレッスンではこれも解決します。
確認問題
素朴なリダクションでwarpダイバージェンスを引き起こす原因を考えてみてください。
まとめ
最も小さいスレッド IDに処理を割り当てることで発散をなくし、warp 全体がアクティブかアイドルのどちらかになるようにしました。計算は同じまま、リダクションは高速になります。次はバンクコンフリクトです。🚀
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「ワープダイバージェンスをなくす」レッスンは無料ですか？
はい。「ワープダイバージェンスをなくす」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「ワープダイバージェンスをなくす」で何を学びますか？
インデックスを組み替えてワープを効率よく動かします。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン2/4です。
「ワープダイバージェンスをなくす」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- リダクションツリーの考え方
- ワープダイバージェンスをなくす
- 逐次アドレッシング
- 複数ブロックで最終リダクションする