为什么这个内核在 GTX760 上比 GTX560 慢

Question

我有GTX560。我今天买了GTX760。为什么下面的内核在 GTX760 上（~0.031ms）比 GTX560（~0.0232ms）慢。当我将 n 增加到 1000000 时，速度更快（约 25%），但对于较小的 n 则不然。我有两台电脑。第一个（内置 GTX560）是 Intel(R) Core(TM) i5 CPU、P7P55D-E LX、CUDA 5.0、Kubuntu 12.04。第二个（GTX760内），AMD FX(tm)-6300，主板760GA-P43(FX)，CUDA 6.5 Kubuntu 14.04。但我仍然认为，原因不是由于不同的CPU等......

GTX560: nvcc -arch=sm_20 -fmad=false -O3 -o vecc vecc.cu -lm
GTX760: nvcc -arch=sm_30 -fmad=false -O3 -o vecc vecc.cu -lm

我也尝试改变块大小，但没有根本效果。

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

// CUDA kernel. Each thread takes care of one element of c
__global__ void vecAdd(double *a, double *b, double *c, int n)
{
 // Get our global thread ID
 int id = blockIdx.x*blockDim.x+threadIdx.x;

 // Make sure we do not go out of bounds
 if (id < n)
 c[id] = sqrt(a[id]*b[id]);
}

int main( int argc, char* argv[] )
 {
  cudaEvent_t start, stop;
  float elapsedTime;

  // Size of vectors
  int n = 512;

  // Host input vectors
  double *h_a;
  double *h_b;
  //Host output vector
  double *h_c;

  // Device input vectors
  double *d_a;
  double *d_b;
  //Device output vector
  double *d_c;

  // Size, in bytes, of each vector
  size_t bytes = n*sizeof(double);

  // Allocate memory for each vector on host
  h_a = (double*)malloc(bytes);
  h_b = (double*)malloc(bytes);
  h_c = (double*)malloc(bytes);

  // Allocate memory for each vector on GPU
  cudaMalloc(&d_a, bytes);
  cudaMalloc(&d_b, bytes);
  cudaMalloc(&d_c, bytes);

  int i;
  // Initialize vectors on host
  for( i = 0; i < n; i++ ) {
   h_a[i] = sin(i)*sin(i);
   h_b[i] = cos(i)*cos(i);
  }

 // Copy host vectors to device
 cudaMemcpy( d_a, h_a, bytes, cudaMemcpyHostToDevice);
 cudaMemcpy( d_b, h_b, bytes, cudaMemcpyHostToDevice);

 int blockSize, gridSize;

 // Number of threads in each thread block
 blockSize = 256;

 // Number of thread blocks in grid
 gridSize = (int)ceil((float)n/blockSize);

 // Execute the kernel
 cudaEventCreate(&start);
 cudaEventRecord(start,0);

 vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

 cudaEventCreate(&stop);
 cudaEventRecord(stop,0);
 cudaEventSynchronize(stop);

 cudaEventElapsedTime(&elapsedTime, start,stop);
 printf("Elapsed time : %f ms\n" ,elapsedTime);

 // Copy array back to host
 cudaMemcpy( h_c, d_c, bytes, cudaMemcpyDeviceToHost );

 // Sum up vector c and print result divided by n, this should equal 1 within error
 double sum = 0;
 for(i=0; i<n; i++)
  sum += h_c[i];
 printf("final result: %f\n", sum/n);

 // Release device memory
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);

 // Release host memory
 free(h_a);
 free(h_b);
 free(h_c);

 return 0;
}

Answer 1

这几乎是评论的延伸，因此它们值得称赞。

这里有两种不同的情况需要研究：

• 答：启动了 512 个线程
• B：启动了 1000000 个线程

在A上，您没有为 GPU 提供足够的工作，您基本上是在测量内核执行/启动的开销。正如评论所指出的，这取决于您的系统。GPU 端花费的时间可以忽略不计。

在这里和这里您可以找到一些信息和计时，例如这张图表，说明了不同 GPU 的执行开销：

在B上，随着线程数量的增加，花在 GPU 端的时间会更长。在这种情况下，760 拥有更好的硬件，并且可以更快地完成工作，从而克服了内核启动开销。

这里还有一些与 CUDA 编程模型本身相关的其他因素；提供更多工作会对 GPU 的性能产生积极影响，但我认为讨论超出了本答案的范围。查看这些帖子 ( 1 , 2 ) 以了解该主题。

您的内核基本上受到内存带宽的限制，760 超过 192 GB/s，而 560 的峰值带宽约为 128 GB/s，因此即使您的卡具有相同数量的核心，您的内核也应该在 760 上运行得更快。

关于内存传输的注意事项

您的代码不受内存传输开销的影响，因为它们超出了测量区域，但我仍然会在这里对此进行注释，因为它有助于解释完整代码的性能差异。

开销和传输时间还取决于您的整个系统；包括硬件和软件两个方面。考虑一下您的机器拥有的 RAM 内存带宽；这取决于主板芯片组、模块的时钟频率、通道和模块的数量、CPU 可以处理的峰值带宽……等，其中一些参数也会影响 PCI 上的内存传输速度。

我鼓励您测量系统上的可用带宽。您可以使用流基准测试来测量 RAM 内存带宽，并使用 CUDA 示例（实用程序目录）中提供的带宽实用程序来测量 PCI 上的 CPU-GPU 内存带宽。这将使您深入了解您的机器并为进一步比较提供起点。