0Pricing
CUDA Academy · 课时

使用 float4 进行向量化加载

更宽的事务带来更高带宽

使用 float4 进行向量化加载 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

每条指令移动更多数据

普通加载一次获取一个值。向量化加载通过一条更宽的指令获取多个相邻值,从而每次请求完成更多工作。

认识 float4

CUDA 提供了内置向量类型。float4 将四个浮点数打包到一个 16 字节的数据组中,您可以一次加载或存储它们。

float4 v = make_float4(1, 2, 3, 4);
float first = v.x; // also .y .z .w

一次加载四个浮点数

从内存读取 float4 会通过一次事务获取四个浮点数。这会将线程必须发出的内存指令数量减少为原来的四分之一。

float4 a = reinterpret_cast<float4*>(in)[i];

重新解释指针

要在普通浮点数组上使用向量化加载,需要转换指针。reinterpret_cast 允许您将相同的字节读取为 float4 元素。

float4* in4 = reinterpret_cast<float4*>(in);
float4 chunk = in4[idx];

必须满足对齐要求

float4 必须位于 16 字节边界上。如果数据没有正确对齐,加载操作就是非法的,内核可能崩溃或读取无效数据。

cudaMalloc 会为您对齐

好消息是:cudaMalloc 分配的缓冲区至少按 256 字节对齐。因此,新分配的设备数组从起始位置开始就可以安全地作为 float4 读取。

请求更少,带宽更多

更宽的加载意味着在处理相同数据时需要发出的请求更少。当内核受内存而非计算限制时,这可能提高实际达到的带宽。

按向量而不是浮点数索引

使用 float4 后,每个线程现在处理四个元素。全局索引会逐个经过 float4 槽位,因此线程总数减少为原来的四分之一。

int i = blockIdx.x * blockDim.x + threadIdx.x;
float4 d = in4[i]; // covers 4 floats

其他向量宽度

float4 不是唯一的选择。像 float2 和 int4 这样的类型可以执行 8 字节或 16 字节的加载,因此您可以选择适合数据的宽度。

注意剩余元素

如果数组长度不是四的倍数,请在向量化部分之后使用普通的标量循环处理额外的 余数 元素。

它会消耗寄存器

一个 float4 包含四个值,因此每个线程会使用更多寄存器。一如既往,请确认带宽增益超过占用率下降带来的损失。

快速检查

您将一个 float 数组重新解释为 float4,但内核发生了崩溃。为什么?

回顾:更宽就更快

您已经了解到,float4 可以一次加载四个浮点数,从而减少指令数量并提升带宽。请保持数据对齐,并处理剩余元素。📦

常见问题解答

「使用 float4 进行向量化加载」课时是免费的吗?

是的 — 「使用 float4 进行向量化加载」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。

「使用 float4 进行向量化加载」这节课中我会学到什么?

更宽的事务带来更高带宽 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗?

无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 3 节课,共 4 节。

「使用 float4 进行向量化加载」课时需要多长时间?

大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。

我能在这节 CUDA Academy 课中编写并运行代码吗?

能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。

此课程中的所有课时

  1. 指令级并行
  2. 使用 #pragma unroll 展开循环
  3. 使用 float4 进行向量化加载
  4. 寄存器压力与溢出
← 返回 CUDA Academy