バンク競合を防ぐ
パディングで共有メモリへのアクセスを高速化できる理由を学びます。
「バンク競合を防ぐ」は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 for each thread in a warp. Spreading accesses across them lets all 32 threads read at full speed.
How Addresses Map
Consecutive 4-byte words land in consecutive banks, wrapping around after 32. So word 0 is bank 0, word 32 is bank 0 again.
What a Bank Conflict Is
A bank conflict happens when two threads in a warp hit the same bank but different words. The hardware must serialize those accesses.
Serialization Costs Speed
An N-way conflict turns one fast access into N slow ones. A 32-way conflict can make shared memory as slow as if it were single-file. 🐢
The Happy Case: Stride One
When each thread reads tile[threadIdx.x], every thread hits a distinct bank. That is conflict-free and runs at full bandwidth.
float v = tile[threadIdx.x];The Trap: Stride of 32
Reading with a stride of 32 sends every thread to the same bank. That is the worst case, a full 32-way conflict.
float v = tile[threadIdx.x * 32];Even Strides Hurt Too
Any even stride that shares a factor with 32 causes partial conflicts. A stride of 2, for example, gives a 2-way conflict across the warp.
The Broadcast Exception
Good news: if all threads read the same address, the hardware broadcasts it in one shot. Same word is fine, only same bank with different words conflicts.
Padding to the Rescue
For 2D tiles, add one extra column with +1 padding. This shifts each row so column accesses no longer all land in one bank.
__shared__ float tile[32][33];Why +1 Works
The extra column makes the row length coprime with 32. Now stepping down a column visits a different bank each time, killing the conflict.
Profile, Do Not Guess
Bank conflicts are invisible in source code. Use Nsight Compute to measure shared-memory conflicts before spending effort fixing them.
Quick Check
Let us check what makes shared access fast.
Recap
You learned that shared memory has 32 banks, that same-bank different-word access serializes, and that +1 padding fixes column conflicts. Next: dynamic shared memory. 🎯
よくある質問
「バンク競合を防ぐ」レッスンは無料ですか？
はい。「バンク競合を防ぐ」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと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フィードバックを取得できます。ローカル設定は不要です。