リダクションに__shfl_down_syncを使う
バリアなしでワープ内のリダクションを行います。
「リダクションに__shfl_down_syncを使う」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン2/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
Shuffle でレジスタを移動する
shuffleを使うと、あるレーンが別のレーンのレジスタ値を直接読み取れます。shared memory や途中のバリアなしで、スレッド間をデータが移動します。
shfl_down_sync の紹介
リダクションの中心となるのがshfl_down_syncです。各レーンは、ワープ内で固定個数分だけ上にあるレーンから値を取得します。
float v = __shfl_down_sync(mask, val, offset);引数を読む
呼び出しには、アクティブなマスク、共有する値、オフセットを指定します。レーン i は、レーン i + offset が保持する値を受け取ります。
受け取った値を加算する
ワープの合計は、shuffle で取得した値を加算するだけです。レーン i は隣のレーンの値を取得し、自分の累積値に加えます。
val += __shfl_down_sync(mask, val, offset);オフセットを半分にする
木構造のリダクションと同様に、オフセットは 16 から始まり、各ステップで半分になります。16、8、4、2、1 の5ステップで、ワープ全体の合計が求まります。
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);5ステップでよい理由
ワープには32個のレーンがあり、2 の5乗は32です。そのため、5回半分にすることで、32個すべての値を1つにまとめるのにちょうど足ります。
答えはレーン 0 に入る
ループの後、完全なワープ合計はlane 0に格納されています。他のレーンには不完全な値が入っているため、結果を書き出すのはlane 0だけにしてください。
バリアは不要です
交換はロックステップでレジスタを介して行われるため、syncthreadsは完全に省略できます。syncサフィックスが、マスクされたレーンをすでに調整しています。
共有メモリより高速です
ワープシャッフルによるリダクションは、共有メモリ版よりも高速です。命令数が少なく、バンク競合もbarrierによる停止もありません。最後の32要素にはこれが高速な経路です。
適切なマスクを渡します
一部のレーンがすでに終了している場合、全レーン用のマスクは不適切です。activemaskで実行中の集合を取得し、各シャッフルが存在するレーンだけを対象にするようにします。
2段階のリダクション
大規模なリダクションでは、各ワープをシャッフルでリダクションしてから、ワープごとの結果をshared memoryで結合することがよくあります。シャッフルは、この低コストな内側の段階を見事に処理します。
簡単な確認
ループが終了したとき、どのレーンが答えを保持しているか考えてみてください。
まとめ
shfl_down_syncを使ってワープを合計しました。オフセットを5回半分にすると、バリアなしでlane 0に合計が格納されます。次はballotとvoteです。✨
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「リダクションに__shfl_down_syncを使う」レッスンは無料ですか?
はい。「リダクションに__shfl_down_syncを使う」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「リダクションに__shfl_down_syncを使う」で何を学びますか?
バリアなしでワープ内のリダクションを行います。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか?
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン2/4です。
「リダクションに__shfl_down_syncを使う」レッスンにはどのくらい時間がかかりますか?
ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか?
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- ワープ、レーン、マスク
- リダクションに__shfl_down_syncを使う
- BallotとVote関数
- Cooperative Groups