CUDA Academy · レッスン

リダクションに__shfl_down_syncを使う

バリアなしでワープ内のリダクションを行います。

レッスン 2/413 ステップ

「リダクションに__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フィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

  1. ワープ、レーン、マスク
  2. リダクションに__shfl_down_syncを使う
  3. BallotとVote関数
  4. Cooperative Groups
← CUDA Academyに戻る