为什么 cuda 指针内存访问比全局设备内存访问慢？

问题描述

#include <vector_functions.h>
#include <vector_types.h>

#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <string>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

__device__ int foo[16];
__device__ int bar[16];

__global__ void go(const int* ptr) {
  printf("device: tid = %d,foo = %p\n",blockIdx.x,foo);
  printf("device: tid = %d,ptr = %p\n",ptr);

  int val = threadIdx.x;
  for (int i = 0; i < (1 << 20); i++) {
    bar[blockIdx.x] = val;
    val = (val * 19 + ptr[threadIdx.x]) % (int)(1e9 + 7); // change ptr to foo for experiment
  }
}

int main() {
  int* ptr = nullptr;
  cudaGetSymbolAddress((void**)&ptr,foo);

  cudaEvent_t start,stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  cudaEventRecord(start);
  go<<<16,16>>>(ptr);
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  cudaDeviceSynchronize();

  float ms;
  cudaEventElapsedTime(&ms,start,stop);
  printf("%.6fms\n",ms);

  return 0;
}

在我的 GeForce GTX 1080 上： 使用 ptr 需要 180 毫秒，但使用 foo 只需要 36 毫秒，尽管 ptr 和 foo 指向完全相同的地址。我认为它们应该以相同的速度执行，因为它们都是由 L2 缓存的全局内存。

我使用的是 Linux，我的编译命令是：

nvcc -gencode=arch=compute_61,code=compute_61 -Xptxas -O3 test.cu -o test

谁能解释一下为什么？

解决方法

这两种情况不同的原因是，当显式使用 foo 时，编译器（在这种情况下为 ptxas）知道 foo 不会 {{3 }} bar，等可以做具体的优化。当改用内核参数 ptr 时，编译器不知道此别名是否发生，并假设它可能发生。这对设备代码生成有重大影响。

作为证明点，使用以下内核原型重新编译您的测试用例：

__global__ void go(const int*  __restrict__ ptr) {

你会看到时差消失了。这是 alias 编译器，ptr 不能为任何其他已知位置（例如 bar）设置别名，因此这允许在两种情况下生成类似的代码。 （在现实世界中，当您准备与编译器签订此类合同时，您才会/应该只使用此类装饰。）

详情：

请务必记住，设备代码编译器是一个优化编译器。此外，从单线程的角度来看，设备代码编译器主要对正确性感兴趣。多线程访问同一个位置不是考虑到这个答案，也确实不是设备代码编译器考虑的。当多个线程访问同一位置时，确保正确性是程序员的责任。

有了那个序言，这里的主要区别似乎是优化之一。知道 foo（或 ptr）没有别名 bar 并且仅考虑单个执行线程，很明显您的内核循环代码可以重写为：

int val = threadIdx.x;
int ptrval = ptr[threadIdx.x];  // becomes a LDG instruction
for (int i = 0; i < ((1 << 20)-1); i++) {
 val = (val * 19 + ptrval) % (int)(1e9 + 7); 
} 
bar[blockIdx.x] = val;          // becomes a STG instruction

此优化的一个主要影响是我们从多次写入 bar 变为仅写入一次。通过这种优化，ptr 的读取也可以“优化到寄存器中”（因为我们现在知道它是循环不变的）。最终效果是消除了循环中的所有全局加载和存储。另一方面，如果 ptr 可能会或可能不会别名 bar，那么我们必须考虑这种可能性，并且上述优化将不成立。

这似乎是编译器正在做的事情。在我们使用 foo（或 __restrict__）的情况下，编译器（在 sass 代码中）在开头安排了一个全局加载，在结尾安排了一个全局存储，以及部分展开的充满整数运算的循环。

然而，当我们将代码保持原样/发布时，编译器也部分展开了循环，但在部分展开的循环中散布了 LDG 和 STG 指令。>

您可以使用 informing 自行观察，例如：

cuobjdump -sass test

（对于每种情况）

设备代码 printf 语句不会实质性地改变这里的任何观察结果，因此为了分析的简单起见，我将删除它们。