合并读取与跨步读取
了解线程到地址的映射为何重要
合并读取与跨步读取 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 2 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
关键在于映射
合并访问取决于每个线程访问哪个地址。线程到地址的映射决定了一次事务能否服务整个线程束。
合并访问模式
当相邻线程读取相邻地址时,线程束覆盖的是一段连续区域。这种整齐的布局称为合并访问。
int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = data[i];它为何更高效
线程 0 访问索引 0,线程 1 访问索引 1,依此类推。全部 32 个地址位于同一个对齐的内存行中,因此 GPU 只需要一次事务。
引入步长
步长是连续线程访问的地址之间的固定间隔。这个间隔一旦超过一个元素,合并访问就会开始失效。
float v = data[i * stride];跨步读取会分散
步长为 2 时,线程 0 读取 0,线程 1 读取 2,线程 2 读取 4。此时线程束覆盖的内存范围扩大了一倍,因此需要更多事务。
代价会不断增加
步长翻倍,触及的内存行数量大致也会翻倍。较大的步长可能让线程束需要执行许多事务,却只使用每个内存行中的一小部分。
常见陷阱
让每个线程负责行主序矩阵中的一整列,会产生巨大的步长。代码看起来很整齐,却会悄悄地把每个线程束的访问分散开。
float v = matrix[threadIdx.x * width + row];转置映射关系
通常只需交换随线程变化最快的索引即可修复。让连续线程遍历连续内存,读取就会再次合并。
float v = matrix[row * width + threadIdx.x];偏移也会造成影响
即使访问模式本身是合并的,如果从内存行中间开始,也会受到影响。未对齐的偏移量会让线程束跨越边界,把一次读取拆成两次。
从线程束角度思考
判断一种访问模式时,不要只想象一个线程。请同时想象全部 32 个线程,并思考它们的地址合计覆盖了多少个内存行。🔍
经验法则
让变化最快的索引跟随 threadIdx.x。只要养成这个习惯,大多数全局内存读取都能自动实现合并。
快速检查
请选择合并效果最好的访问模式。
回顾
您已经了解到,合并读取会让相邻访问保持在一起,而步长会把它们分散到不同内存行。请让 threadIdx.x 驱动变化最快的索引。🎉
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「合并读取与跨步读取」课时是免费的吗?
是的 — 「合并读取与跨步读取」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「合并读取与跨步读取」这节课中我会学到什么?
了解线程到地址的映射为何重要 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 2 节课,共 4 节。
「合并读取与跨步读取」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- 什么是内存事务
- 合并读取与跨步读取
- 数组结构与结构体数组
- 测量有效带宽