複数ブロックで最終リダクションする
ブロックごとの部分和を統合します。
「複数ブロックで最終リダクションする」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン4/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
ブロック同士は通信できない
ブロック内のリダクションは簡単ですが、ブロックは独立して実行され、カーネルの途中で互いに同期できません。そのため、1回の起動ですべてを合計することはできません。
各ブロックが部分結果を生成する
そこで各ブロックが自分のチャンクを1つの値、つまり部分和にリダクションし、blockIdx をインデックスとして小さな出力配列に書き込みます。
if (tid == 0)
out[blockIdx.x] = data[0];値の数を減らせる
1000ブロックを使うと、100万個の入力が1000個の部分結果になります。難しい処理は終わり、あとは小さな配列を結合するだけです。
方法1：もう一度起動する
最も簡単な仕上げ方は、部分結果に対して同じカーネルを2回目に起動することです。値が1つだけになるまで繰り返します。
1つになるまで再帰的に処理する
各パスで、配列のサイズがブロックサイズ分だけ縮小します。数回の再帰的な起動で、数百万個の値を最終的な1つの合計に減らせます。
方法2：アトミック演算
別の方法として、各ブロックのスレッド0が部分結果を1つの global total にatomicAddで直接加算できます。これにより、2回目のカーネルを避けられます。
if (tid == 0)
atomicAdd(total, data[0]);アトミック演算のトレードオフ
アトミック演算は単純で起動も1回だけですが、多数のブロックが同じアドレスで競合すると、処理が直列化されます。部分結果が少ない場合は通常問題ありません。
方法3：グリッドストライドループ
grid-stride loopを使うと、各スレッドが最初に多数の要素をレジスタへ合計できます。そのため、最終ステップまでに必要なブロック数を大幅に減らせます。
for (int i = gid; i < n; i += gridDim.x * blockDim.x)
sum += in[i];ブロックを減らしてオーバーヘッドを抑える
最初にスレッドあたりの処理量を増やすと、部分結果と起動回数を減らせます。これは要素ごとに1スレッドを生成するより高速なことがよくあります。
最初に合計値をゼロにする
アトミック演算を使う場合は、起動前に global total をゼロにすることを忘れないでください。そうしないと、そのメモリに残っていた不定な値から合計が始まってしまいます。
問題のサイズで選ぶ
小さな入力では単純なアトミック演算が適しています。大きな入力では2パス方式またはグリッドストライド方式が有利です。実際のデータで測定して選んでください。
クイックチェック
1回のカーネル起動だけでは配列全体を直接合計できない理由を考えてみてください。
まとめ
各ブロックが部分和を生成し、それらを2回目の起動、アトミック演算、またはグリッドストライド方式で結合します。これで、どのようなサイズの配列でもリダクションできるようになりました。🏁
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は初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン4/4です。
「複数ブロックで最終リダクションする」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- リダクションツリーの考え方
- ワープダイバージェンスをなくす
- 逐次アドレッシング
- 複数ブロックで最終リダクションする