2

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

这是我无法弄清楚的：在计算能力为 1.3（特斯拉 C1060）的设备上，只要 SCATTERSIZE 是 16 的倍数，性能就会显着下降。而在计算能力为 2.0（特斯拉 C2075）的设备上，只要 SCATTERSIZE 性能就会受到很大影响是 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 = %d\n", blocksPerGrid, SCATTERSIZE);

  Rule* small;      
  Rule* big;        

  cudaError_t err = cudaMalloc(&big, numThreads * 128 * sizeof(Rule));
  printf("Malloc big: %s\n",cudaGetErrorString(err));

  err = cudaMalloc(&small, numThreads * sizeof(Rule));
  printf("Malloc small: %s\n",cudaGetErrorString(err));

  gokernel <<< blocksPerGrid, threadsPerBlock >>> (big, small);
  err = cudaThreadSynchronize();
  printf("Kernel launch: %s\n", cudaGetErrorString(err));
}
4

1 回答 1

1

因为 的实现__device__ memcpy是隐藏的（它是内置的编译器），所以很难说到底是什么原因。一种预感（感谢 njuffa 对此）是所谓的分区驻留，其中来自许多线程的地址映射到一个或几个物理 DRAM 分区，而不是分布在它们之间。

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

您通常可以通过在数组中添加一些填充以避免违反偏移量来解决此效果，但根据您的计算可能不值得。

于 2012-08-29T06:05:24.803 回答