仔细检查对 CUDA 中内存合并的理解

Question

假设我定义了一些对 GPU 可见的数组：

double* doubleArr = createCUDADouble(fieldLen);
float* floatArr = createCUDAFloat(fieldLen);
char* charArr = createCUDAChar(fieldLen);

现在，我有以下 CUDA 线程：

void thread(){
  int o = getOffset(); // the same for all threads in launch
  double d = doubleArr[threadIdx.x + o];
  float f = floatArr[threadIdx.x + o];
  char c = charArr[threadIdx.x + o];
}

我不太确定我是否正确解释了文档，它对我的设计非常关键：double、float 和 char 的内存访问是否可以很好地合并？ （猜测：是的，它将适合 sizeof(type) * blockSize.x / (transaction size) 笔交易，再加上可能在上下边界处增加一笔交易。）

Answer 1

是的，对于您展示的所有情况，假设 createCUDAxxxxx 转化为某种普通的 cudaMalloc 类型操作，一切都应该很好地合并。

如果我们有通过 cudaMalloc 分配的普通一维设备数组，如果我们的加载模式包含以下形式的数组索引，通常我们应该具有良好的跨线程合并行为：

data_array[some_constant + threadIdx.x];

数组是什么数据类型真的无关紧要 - 它会很好地合并。

但是，从性能的角度来看，全局加载（假设 L1 未命中）将以最小 128 字节的粒度发生。因此，每个线程加载更大的尺寸（例如，int、float、double、float4 等）可能会提供稍微更好的性能。如果负载跨越足够多的扭曲，缓存往往会减轻任何差异。

使用 profiler 在特定代码段上验证这一点也很容易。有很多方法可以做到这一点，具体取决于您选择的分析器，但是例如使用 nvprof，您可以这样做：

nvprof --metric gld_efficiency ./my_exe

并且它将 return 一个平均百分比数字，或多或少准确地反映了在全局负载上发生的最佳合并的百分比。

This 是我经常引用的演示文稿，以获取有关内存优化的其他背景信息。

我想有人会注意到这个模式：

data_array[some_constant + threadIdx.x];

大致对应于上述演示文稿的幻灯片 40-41 中显示的访问类型。 aha!! 效率下降到 50%-80%。这是真的，如果只考虑一个单一的扭曲负载。但是，参考幻灯片 40，我们看到 "first" 加载将需要加载两个缓存行。 在那之后 但是，额外的加载（为简单起见向右移动）每个 warp-load 只需要一个 additional/new 缓存行（假设存在 L1 或 L2 缓存，和合理的地点，即缺乏颠簸）。因此，在相当大的数组（超过 128 字节）上，平均要求每个 warp 一个新的缓存行，这对应于 100% 的效率。

仔细检查对 CUDA 中内存合并的理解

Double-checking understanding of memory coalescing in CUDA

c++

cuda