顺序寻址
共享内存中的无冲突步长
顺序寻址 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 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. ✨
常见问题解答
「顺序寻址」课时是免费的吗？
是的 — 「顺序寻址」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「顺序寻址」这节课中我会学到什么？
共享内存中的无冲突步长 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 3 节课，共 4 节。
「顺序寻址」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。