ヒストグラムを構築する
共有メモリのプライベート化とアトミック操作を使います。
「ヒストグラムを構築する」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
このレッスンの一部はまだ翻訳されておらず、英語で表示されています。
What a Histogram Counts
A histogram counts how many inputs fall into each bin. Many threads will want to increment the same bin, so this is an atomics problem. 📊
The Naive Approach
Each thread reads one element, finds its bin, and increments that bin. Without protection, popular bins lose counts to races.
The Global Atomic Version
The simplest correct fix is one atomicAdd per element straight into global memory. It works, but hot bins serialize threads.
atomicAdd(&hist[bin], 1);The Contention Problem
When data clusters into a few bins, thousands of threads pile onto the same address. That contention can make global atomics painfully slow.
Privatization to the Rescue
Privatization gives each block its own private histogram in fast shared memory. Threads collide only inside their block, not across the whole grid.
Declare the Shared Histogram
Each block declares a shared array sized to the number of bins. It lives on-chip, so atomics there are far cheaper than global ones.
__shared__ int local[NBINS];Step 1: Clear the Bins
Threads cooperatively zero the shared histogram, then call __syncthreads so no one counts before clearing is done.
local[tid] = 0;
__syncthreads();Step 2: Count Locally
Now each thread atomically bumps its bin in shared memory. Same atomicAdd, but on the fast on-chip copy instead of global.
atomicAdd(&local[bin], 1);Step 3: Merge to Global
After a barrier, threads add each shared bin into the global histogram with one atomicAdd per bin. Far fewer global atomics than before.
atomicAdd(&hist[i], local[i]);Why This Is Faster
Shared-memory atomics are quick, and the costly global atomics now fire once per bin per block instead of once per element.
Watch the Bin Count
The private histogram must fit in shared memory. With too many bins, split them into passes or fall back to global atomics.
Quick Check
One question on the histogram strategy.
Recap: Building a Histogram
You built a histogram with global atomics, then sped it up using shared-memory privatization: clear, count locally, merge. ✅
よくある質問
「ヒストグラムを構築する」レッスンは無料ですか？
はい。「ヒストグラムを構築する」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「ヒストグラムを構築する」で何を学びますか？
共有メモリのプライベート化とアトミック操作を使います。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。
「ヒストグラムを構築する」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。