CUDA Academy · 课时

合并读取与跨步读取

了解线程到地址的映射为何重要

第 2 / 4 课13 个步骤

合并读取与跨步读取 是 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 反馈 — 无需本地设置。

此课程中的所有课时

  1. 什么是内存事务
  2. 合并读取与跨步读取
  3. 数组结构与结构体数组
  4. 测量有效带宽
← 返回 CUDA Academy