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

Question

我想添加带有进位的 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;
}

Answer 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。