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

【问题标题】：Why is cuda pointer memory access slower than global device memory access?为什么 cuda 指针内存访问比全局设备内存访问慢？
【发布时间】：2021-06-28 03:04:52
【问题描述】：

#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", blockIdx.x, 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

谁能解释一下原因？

【问题讨论】：

标签： c++ pointers caching cuda restrict-qualifier

【解决方案1】：

这两种情况不同的原因是，当显式使用foo时，编译器（本例中为ptxas）知道foo不知道aliasbar，所以可以做具体的优化。当使用内核参数ptr 代替时，编译器不知道这种别名是否正在发生，并假设它可能发生。这对设备代码生成有重大影响。

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

__global__ void go(const int*  __restrict__ ptr) {

你会看到时差消失了。这是informing 编译器，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，也可以不别名bar，那么我们必须考虑这种可能性，上述优化将不成立。

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

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

您可以使用cuda binary utilities 自己观察这一点，例如：

cuobjdump -sass test

（针对每种情况）

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

【讨论】：