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

I have a simple CUDA kernel which I thought was accessing global memory efficiently. The Nvidia profiler however reports that I am performing inefficient global memory accesses.

举个例子，你的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 数据重组，并进行适当的代码调整。这在其含义和操作方法方面都有很好的记录（例如 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;
}

尽管您可能认为此代码比您的代码 "less efficient"，因为您的代码 "appears" 仅加载和存储 .x、.y、.z，而我的 "appears" 也加载和存储 .w，实际上，由于 GPU 加载和存储 to/from 全局内存的方式，本质上没有区别。尽管您的代码似乎没有触及 .w，但在访问相邻元素的过程中，GPU 将从全局内存中加载 .w 元素，并且（最终）存储 .w元素回到全局内存。

What I don't understand is why the profiler reports that the ideal access would involve 8 L2 transactions per access for line 247

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

CUDA 探查器报告低效的全局内存访问

CUDA profiler reports inefficient global memory access

memory

profiler

caching

cuda