CUDA 全局内存中带有 memcpy 的奇怪"银行冲突"类型行为

Question

我已经将一个性能问题提炼为如下代码。此代码采用128000个64字节结构（"规则"）的数组，并将它们分散在另一个数组中。例如，如果SCATTERSIZE为10，则代码将从"小"数组中复制（"散射"）128000个这些结构，这些结构在索引0、1、2…处连续存储。。。，127999，并将它们放置在索引0、10、20、30…处。。。，"大"数组中的1279990。

以下是我无法理解的：在计算能力为1.3的设备（特斯拉C1060）上，只要散射尺寸是16的倍数，性能就会急剧下降。在计算能力为2.0的设备（特斯拉C2075）上，每当散射尺寸为24的倍数时，性能就会受到很大影响。

我不认为这是一个共享内存库的事情，因为我没有使用共享内存。我不认为这与聚结有关。使用命令行探查器并检查"gputime"条目，我发现1.3设备的运行时增加了300%，2.0设备的运行时间增加了40%，因为SCATTERSIZE不好。我被难住了。这是代码：

#include <stdio.h>
#include <cuda.h>
#include <stdint.h>
typedef struct{
  float a[4][4];
} Rule;
#ifndef SCATTERSIZE
#define SCATTERSIZE 96
#endif
__global__ void gokernel(Rule* b, Rule* s){
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  memcpy(&b[idx * SCATTERSIZE], &s[idx], sizeof(Rule));
}

int main(void){
  int blocksPerGrid = 1000;
  int threadsPerBlock = 128;
  int numThreads = blocksPerGrid * threadsPerBlock;
  printf("blocksPerGrid = %d, SCATTERSIZE = %dn", blocksPerGrid, SCATTERSIZE);
  Rule* small;      
  Rule* big;        
  cudaError_t err = cudaMalloc(&big, numThreads * 128 * sizeof(Rule));
  printf("Malloc big: %sn",cudaGetErrorString(err));
  err = cudaMalloc(&small, numThreads * sizeof(Rule));
  printf("Malloc small: %sn",cudaGetErrorString(err));
  gokernel <<< blocksPerGrid, threadsPerBlock >>> (big, small);
  err = cudaThreadSynchronize();
  printf("Kernel launch: %sn", cudaGetErrorString(err));
}

Answer 1

因为__device__ memcpy的实现是隐藏的（它是内置的编译器），所以很难说确切的原因是什么。一种预感（多亏了njuffa）是，这就是所谓的分区露营，来自许多线程的地址映射到一个或几个物理DRAM分区，而不是分布在它们之间。

在SM 1_2/1_3 GPU上，分区驻留可能会非常糟糕，这取决于内存访问步长，但从SM_2_0设备开始，这一点已经得到了改善，这可以解释为什么效果不那么明显。

您通常可以通过在数组中添加一些填充来避免偏移，来解决这种效果，但这可能不值得，具体取决于您的计算。