逐次アドレッシング
共有メモリで競合のないストライドを実現します。
「逐次アドレッシング」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
このレッスンの一部はまだ翻訳されておらず、英語で表示されています。
Shared Memory Has Banks
Shared memory is split into 32 banks, one per warp lane. When 32 threads hit 32 different banks, all reads happen in a single fast cycle.
Bank Conflicts Slow You Down
If two threads in a warp touch the same bank, that is a bank conflict. The hardware serializes those accesses, costing extra cycles.
Interleaved Addressing
The previous reduction used interleaved addressing: stride starts small and doubles, so partners are close together in shared memory.
int index = 2 * s * tid;
data[index] += data[index + s];Why Interleaving Conflicts
With small, doubling strides, several lanes in a warp map onto the same bank. Those accesses can no longer happen in one cycle.
Flip the Stride Order
Sequential addressing starts the stride large and halves it each step, the reverse of interleaving. This single change removes the conflicts.
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s)
data[tid] += data[tid + s];
__syncthreads();
}Big Stride, Clean Banks
A large stride spreads partner addresses far apart, so each lane lands on its own bank. The warp reads conflict-free in one cycle.
The tid < s Guard
Only the lower half of threads work each step, written as tid < s. That keeps active threads contiguous, so warps stay non-divergent too.
Two Wins at Once
Sequential addressing fixes bank conflicts and avoids warp divergence in the same kernel. One layout change, two performance problems solved.
Still Sync Each Step
You still need a __syncthreads after each step. Threads must see the previous level's writes before they read for the next level.
Result Lands at Index 0
As the stride halves toward zero, all partial sums fold into data[0]. Thread 0 then writes that block's result back to global memory.
A Classic Optimization
This pattern comes straight from NVIDIA's famous reduction guide. Sequential addressing is a textbook step toward a conflict-free kernel.
Quick Check
Think about why a large, halving stride beats a small, doubling one.
Recap
You swapped interleaved for sequential addressing: stride starts large and halves, killing bank conflicts and divergence at once. Up next: multi-block sums. ✨
よくある質問
「逐次アドレッシング」レッスンは無料ですか？
はい。「逐次アドレッシング」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと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フィードバックを取得できます。ローカル設定は不要です。