CUDA 8.0，GTX 1080，为什么向量加法比矩阵乘法慢 5 倍？

Question

我正在使用最新的 CUDA 8.0 和 GTX 1080，以及运行样本来测试速度。（我知道它们并不能反映最佳速度，但我只是想横向比较一下。）

在0_Simple/matrixMul中，速度由代码测量，给出：

Performance= 1029.91 GFlop/s, Time= 0.127 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block

然后我运行 0_Simple/vectorAdd，并从上面的示例中复制速度测试代码。即：

// Measure speed
    cudaEvent_t start;
    cudaEventCreate(&start);
    cudaEvent_t stop;
    cudaEventCreate(&stop);

    cudaEventRecord(start, NULL);
    int nIter = 300;
    for (int i = 0; i < nIter; i++) {
        vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
    }
    cudaEventRecord(stop, NULL);

    cudaEventSynchronize(stop);
    float msecTotal = 0.0f;
    cudaEventElapsedTime(&msecTotal, start, stop);
    float msecPerAdd = msecTotal / nIter;
    double flopsPerAdd = numElements;
    double gigaFlops = (flopsPerAdd * 1.0e-9f) / (msecPerAdd / 1000.0f);
    printf("Performance= %.2f GFLOPS, Time= %.3f ms, Size= %.0f Ops\n", gigaFlops, msecPerAdd, flopsPerAdd);

我还把numElements从50000扩大到了67108864。速度结果是：

Performance= 19.85 GFLOPS, Time= 3.380 ms, Size= 67108864 Ops

几乎慢了 5 倍。

我知道示例代码可能不是最优的，所以谁能告诉我为什么 vectorAdd 代码这么慢，以及如何优化它？

我正在使用 CUDA 8.0 和 GTX 1080

Answer 1

与矩阵乘法不同，向量加法是内存带宽限制运算。衡量其性能的正确方法是衡量全局内存访问的带宽。对于向量加法，它包括2个输入向量和1个输出向量，可以计算如下。

3 * numElements * sizeof(d_A[0]) / kernel_running_time

您可以将其与简单的D2D副本的带宽进行比较，看看您是否达到了峰值。

Answer 2

编辑：为 GPU 和 CPU

添加了代码

我运行下面做实验

***Cuda 内核

__global__ void add(float *a, float *c , size_t N)
{
  int tid = blockIdx.x*blockDim.x + threadIdx.x;    
  if(tid < N) a[tid] += c[tid];
}

***CPU 原始版本（基线）

void naiveAdd(float *h_a, float *h_c, size_t N)
{
   for (size_t i=0; i<N; i++)
      h_a[i] += h_c[i];
}

用于在具有不同 CUDA GCC 版本的同一台机器（740M GPU，第 4 代 i7）中添加两个向量的性能比较。

在 CUDA 8.0 和 GCC 5 上，GPU 版本速度较慢

在 CUDA 7.5 和 GCC 4.8 上，GPU 版本更快

有了这些证据，我可以得出结论，暂时最好避免使用 CUDA 8.0。

CUDA 8.0，GTX 1080，为什么向量加法比矩阵乘法慢 5 倍？

CUDA 8.0, GTX 1080, why is vector addition slower than matrix multiplication by 5x?

cuda

gpu