使用 Cuda 进行 128 位向量加法，性能问题

Posted 2023-03-23

【中文标题】使用 Cuda 进行 128 位向量加法，性能问题

【英文标题】：128-bit vector addition with Cuda, performance issue

【发布时间】：2020-08-25 03:42:47

【问题描述】：

我想添加带有进位的 128 位向量。我的 128 位版本（下面代码中的addKernel128）比基本的 32 位版本（下面的addKernel32）慢两倍。 我有内存合并问题吗？如何获得更好的性能？

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define UADDO(c, a, b) asm volatile("add.cc.u32 %0, %1, %2;" : "=r"(c) : "r"(a) , "r"(b));
#define UADDC(c, a, b) asm volatile("addc.cc.u32 %0, %1, %2;" : "=r"(c) : "r"(a) , "r"(b));

__global__ void addKernel32(unsigned int *c, const unsigned int *a, const unsigned int *b, const int size)

  int tid = blockIdx.x * blockDim.x + threadIdx.x;

  while (tid < size)
  
    c[tid] = a[tid] + b[tid];
    tid += blockDim.x * gridDim.x;
  


__global__ void addKernel128(unsigned *c, const unsigned *a, const unsigned *b, const int size)

  int tid = blockIdx.x * blockDim.x + threadIdx.x;

  while (tid < size / 4)
  
    uint4 a4 = ((const uint4 *)a)[tid],
          b4 = ((const uint4 *)b)[tid],
          c4;

    UADDO(c4.x, a4.x, b4.x)
    UADDC(c4.y, a4.y, b4.y) // add with carry
    UADDC(c4.z, a4.z, b4.z) // add with carry
    UADDC(c4.w, a4.w, b4.w) // add with carry (no overflow checking for clarity)

    ((uint4 *)c)[tid] = c4;

    tid += blockDim.x * gridDim.x;
  


int main()

  const int size = 10000000; // 10 million

  unsigned int *d_a, *d_b, *d_c;

  cudaMalloc((void**)&d_a, size * sizeof(int));
  cudaMalloc((void**)&d_b, size * sizeof(int));
  cudaMalloc((void**)&d_c, size * sizeof(int));

  cudaMemset(d_a, 1, size * sizeof(int)); // dummy init just for the example
  cudaMemset(d_b, 2, size * sizeof(int)); // dummy init just for the example
  cudaMemset(d_c, 0, size * sizeof(int));

  int nbThreads = 512;
  int nbBlocks = 1024; // for example

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start);

  addKernel128<<<nbBlocks, nbThreads>>>(d_c, d_a, d_b, size);

  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  float m = 0;
  cudaEventElapsedTime(&m, start, stop);

  cudaFree(d_c);
  cudaFree(d_b);
  cudaFree(d_a);
  cudaDeviceReset();
  printf("Elapsed = %g\n", m);
  return 0;

【问题讨论】：

128 位添加需要更多的工作，不是吗？你不希望它比 32 位版本慢吗？

在全球范围内，工作量是相同的：确实，每次添加都需要多 4 倍的工作，但必要添加的数量减少了 4 倍。

你没有提到你在哪个平台上（windows 或 linux，哪个 GPU，CUDA 版本，是 WDDM 等），这对于时序分析非常重要。对于计时，在计时测量之前进行“热身”也很常见。当我运行代码的修改版本here（cuda 5.0、rhel 5.5、Quadro5000 GPU）时，我得到的结果是经过的时间大致相同（实际上 128 版本要快一些）。我猜你是在 Windows 上，如果使用 WDDM GPU，准确计时可能会非常困难。

如果你在windows下，我猜，也请确定你是在构建项目的调试版本还是发布版本。

怀疑，我在带有 WDDM 驱动程序的 Windows 7 下。 VS2010、Cuda 5.5、驱动程序版本 320.57 和 Quadro 600 GPU。在调试版本中观察到性能不佳。发布版本要好得多，但仍然慢了 20%。也许我们可以得出结论，在这种环境下测量时间是不可预测的。感谢您测试我的代码。

【参考方案1】：

由于各种原因，在 WDDM GPU 上对 CUDA 代码进行计时可能非常困难。其中大部分都围绕着这样一个事实，即 Windows 将 GPU 作为显示设备进行管理，这可能会在时序中引入各种伪影。一个例子是 windows 驱动程序和 WDDM 将为 GPU 批量工作，并可能在 CUDA GPU 工作的中间交错显示工作。

如果可能，请在 linux 或 windows GPU 上为您的 cuda 代码计时 在 TCC 模式下。 为了提高性能，总是在没有-G 开关的情况下构建。在 Visual Studio 中，这通常对应于构建版本，而不是项目的调试版本。 为了获得良好的性能比较，通常建议在实际测量计时结果之前进行一些“热身运行”。这些将消除“启动”和其他一次性测量问题，您更有可能获得明智的结果。您可能还希望多次运行代码并对结果取平均值。 通常还建议使用与您的 GPU 对应的 arch 标志进行编译，例如 -arch=sm_20 用于 cc2.0 GPU。

【讨论】：

并且通过使用 sm_20 标志进行编译，128 位版本比 32 位版本快一点。谢谢！

好点，更新答案作为附加建议