CUDA 分析器报告全局内存访问效率低下

Question

我有一个简单的 CUDA 内核，我认为它可以有效地访问全局内存。然而，Nvidia 分析器报告说我正在执行低效的全局内存访问。我的内核代码是：

__global__ void update_particles_kernel
(
    float4 *pos, 
    float4 *vel, 
    float4 *acc, 
    float dt, 
    int numParticles
)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
int offset = 0;

while(index + offset < numParticles)
{
    vel[index + offset].x += dt*acc[index + offset].x;   // line 247
    vel[index + offset].y += dt*acc[index + offset].y;
    vel[index + offset].z += dt*acc[index + offset].z;

    pos[index + offset].x += dt*vel[index + offset].x;   // line 251
    pos[index + offset].y += dt*vel[index + offset].y;
    pos[index + offset].z += dt*vel[index + offset].z;

    offset += blockDim.x * gridDim.x;
}

特别是分析器报告以下内容：

来自 CUDA best practices guide 它说：

“对于计算能力为 2.x 的设备，可以很容易地总结要求：warp 线程的并发访问将合并为与服务所需的缓存行数相等的事务数warp 的所有线程。默认情况下，所有访问都通过 L1 缓存，作为 128 字节的行。对于分散的访问模式，为了减少过度读取，有时只缓存在 L2 中是有用的，它缓存较短的 32-字节段（参见 CUDA C 编程指南）。

对于计算能力 3.x 的设备，对全局内存的访问仅缓存在 L2 中； L1 保留用于本地内存访问。一些计算能力为 3.5、3.7 或 5.2 的设备也允许在 L1 中选择加入缓存全局变量。"

现在在我的内核中，基于此信息，我预计需要 16 次访问才能为 32 线程扭曲提供服务，因为 float4 是 16 字节，并且在我的卡（770m 计算能力 3.0）上，从 L2 缓存读取是在 32 中执行的字节块（16 字节 * 32 线程 / 32 字节缓存线 = 16 次访问）。确实，正如您所看到的分析器报告我正在进行 16 次访问。我不明白为什么分析器报告理想的访问将涉及第 247 行的每次访问 8 个 L2 事务，而其余行的每次访问仅涉及 4 个 L2 事务。有人可以解释我在这里缺少什么吗？

Answer 1

我有一个简单的 CUDA 内核，我认为它可以有效地访问全局内存。然而，Nvidia 分析器报告说我正在执行低效的全局内存访问。

举个例子，你的float4vel数组是这样存储在内存中的：

0.x 0.y 0.z 0.w 1.x 1.y 1.z 1.w 2.x 2.y 2.z 2.w 3.x 3.y 3.z 3.w ...
  ^               ^               ^               ^             ...
  thread0         thread1         thread2         thread3

所以当你这样做时：

vel[index + offset].x += ...;   // line 247

您正在我上面标记的位置 (.x) 访问（存储）。每个^ 标记之间的间隙表示访问模式效率低下，分析器指出了这一点。 （在下一行代码中，您将存储到.y 位置并不重要。）

至少有 2 种解决方案，其中一种是经典的 AoS -> SoA 重组数据，并进行适当的代码调整。就其含义和操作方法而言，这是有据可查的（例如 here 在 cuda 标签和其他地方），所以我会让你查一下。

另一个典型的解决方案是在需要时为每个线程加载一个float4 数量，并在需要时为每个线程存储一个float4 数量。可以对您的代码进行微不足道的修改来执行此操作，这应该会改进分析结果：

//preceding code need not change
while(index + offset < numParticles)
{
    float4 my_vel = vel[index + offset];
    float4 my_acc = acc[index + offset];
    my_vel.x += dt*my_acc.x;   
    my_vel.y += dt*my_acc.y;
    my_vel.z += dt*my_acc.z;
    vel[index + offset] = my_vel;

    float4 my_pos = pos[index + offset];
    my_pos.x += dt*my_vel.x; 
    my_pos.y += dt*my_vel.y;
    my_pos.z += dt*my_vel.z;
    pos[index + offset] = my_pos;

    offset += blockDim.x * gridDim.x;
}

即使您可能认为此代码比您的代码“效率低”，因为您的代码“似乎”只是加载和存储 .x、.y、.z，而我的“似乎”是还加载和存储.w，事实上，由于GPU加载和存储到/从全局内存的方式，本质上没有区别。尽管您的代码似乎没有触及.w，但在访问相邻元素的过程中，GPU 将从全局内存中加载.w 元素，并且（最终）将.w 元素存储回全局内存。

我不明白为什么探查器报告理想的访问将涉及第 247 行的每次访问 8 个 L2 事务

对于原始代码中的第 247 行，对于 acc.x 的加载操作，每个线程访问一个 float 数量，对于 vel.x 的加载操作，每个线程访问一个 float 数量。每个线程的 float 数量本身应该需要 128 字节的 warp，即 4 个 32 字节的 L2 缓存线。两个加载一起需要 8 个 L2 缓存线加载。这是理想的情况，假设数量很好地打包在一起（SoA）。但这不是您所拥有的（您拥有 AoS）。