我正在尝试在 CUDA C++ 代码上运行 vector 步长加法函数，但对于大小为 5,000,000 的大型 float 组，它的运行速度也比我的 CPU 版本慢。以下是我正在谈论的相关 CUDA 和 cpu 代码:

#define THREADS_PER_BLOCK 1024
typedef float real;
__global__ void vectorStepAddKernel2(real*x, real*y, real*z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < size)
    {
        x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep];
    }
}

cudaError_t vectorStepAdd2(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{

    cudaError_t cudaStatus;
    int threadsPerBlock = THREADS_PER_BLOCK;
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock;
    vectorStepAddKernel2<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size, xstep, ystep, zstep);

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching vectorStepAddKernel!\n", cudaStatus);
        exit(1);
    }

    return cudaStatus;
}

//CPU function:

void vectorStepAdd3(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{
    for(int i=0;i<size;i++)
    {
        x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep];
    }
}

当 3 个数组的大小均为 5,000,000 且大小为 50,000(即，以这种逐步方式将 50,000 个元素相加)时，调用 vectorStepAdd2 的计算速度比 vectorStepAdd3 慢。

关于我可以做什么来加速 GPU 代码有什么想法吗？ 我的设备是 Tesla M2090 GPU

谢谢

最佳答案

回答您的问题“我可以做些什么来加速 GPU 代码？”

首先，让我先说明所提议的操作 X = alpha * Y + beta * Z 对每字节数据传输所需的计算强度并不高。结果，我无法在这个特定代码上击败 CPU 时间。然而，涵盖 2 个想法以加速此代码可能是有益的:

1. 使用 page-locked用于数据传输操作的存储器。这使 GPU 版本的数据传输时间减少了大约 2 倍，这在 GPU 版本的整体执行时间中占主导地位。

2. 在 cudaMemcpy2D 中使用跨步复制技术正如@njuffa 所提议的那样here .结果是双重的:我们可以将数据传输量减少到仅计算所需的数量，并且然后我们可以重新编写内核以按照建议连续地对数据进行操作评论(再次由 njuffa 发表)。这使数据传输时间进一步缩短了 3 倍，内核计算时间缩短了约 10 倍。

此代码提供了这些操作的示例:

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


#define THREADS_PER_BLOCK 1024
#define DSIZE 5000000
#define WSIZE 50000
#define XSTEP 47
#define YSTEP 43
#define ZSTEP 41
#define TOL 0.00001f


#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

typedef float real;

__global__ void vectorStepAddKernel2(real *x, real *y, real *z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < size)
    {
        x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep];
    }
}

__global__ void vectorStepAddKernel2i(real *x, real *y, real *z, real alpha, real beta, int size)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < size)
    {
        x[i] = alpha* y[i] + beta*z[i];
    }
}

void vectorStepAdd2(real *x, real *y, real *z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{

    int threadsPerBlock = THREADS_PER_BLOCK;
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock;
    vectorStepAddKernel2<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size, xstep, ystep, zstep);
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel2 fail");
}


void vectorStepAdd2i(real *x, real *y, real *z, real alpha, real beta, int size)
{

    int threadsPerBlock = THREADS_PER_BLOCK;
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock;
    vectorStepAddKernel2i<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size);
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel3 fail");
}

//CPU function:

void vectorStepAdd3(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep)
{
    for(int i=0;i<size;i++)
    {
        x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep];
    }
}

int main() {

  real *h_x, *h_y, *h_z, *c_x, *h_x1;
  real *d_x, *d_y, *d_z, *d_x1, *d_y1, *d_z1;

  int dsize = DSIZE;
  int wsize = WSIZE;
  int xstep = XSTEP;
  int ystep = YSTEP;
  int zstep = ZSTEP;
  real alpha = 0.5f;
  real beta = 0.5f;
  float et;

/*
  h_x = (real *)malloc(dsize*sizeof(real));
  if (h_x == 0){printf("malloc1 fail\n"); return 1;}
  h_y = (real *)malloc(dsize*sizeof(real));
  if (h_y == 0){printf("malloc2 fail\n"); return 1;}
  h_z = (real *)malloc(dsize*sizeof(real));
  if (h_z == 0){printf("malloc3 fail\n"); return 1;}
  c_x = (real *)malloc(dsize*sizeof(real));
  if (c_x == 0){printf("malloc4 fail\n"); return 1;}
  h_x1 = (real *)malloc(dsize*sizeof(real));
  if (h_x1 == 0){printf("malloc1 fail\n"); return 1;}
*/

  cudaHostAlloc((void **)&h_x, dsize*sizeof(real), cudaHostAllocDefault);
  cudaCheckErrors("cuda Host Alloc 1 fail");
  cudaHostAlloc((void **)&h_y, dsize*sizeof(real), cudaHostAllocDefault);
  cudaCheckErrors("cuda Host Alloc 2 fail");
  cudaHostAlloc((void **)&h_z, dsize*sizeof(real), cudaHostAllocDefault);
  cudaCheckErrors("cuda Host Alloc 3 fail");
  cudaHostAlloc((void **)&c_x, dsize*sizeof(real), cudaHostAllocDefault);
  cudaCheckErrors("cuda Host Alloc 4 fail");
  cudaHostAlloc((void **)&h_x1, dsize*sizeof(real), cudaHostAllocDefault);
  cudaCheckErrors("cuda Host Alloc 5 fail");


  cudaMalloc((void **)&d_x, dsize*sizeof(real));
  cudaCheckErrors("cuda malloc1 fail");
  cudaMalloc((void **)&d_y, dsize*sizeof(real));
  cudaCheckErrors("cuda malloc2 fail");
  cudaMalloc((void **)&d_z, dsize*sizeof(real));
  cudaCheckErrors("cuda malloc3 fail");
  cudaMalloc((void **)&d_x1, wsize*sizeof(real));
  cudaCheckErrors("cuda malloc4 fail");
  cudaMalloc((void **)&d_y1, wsize*sizeof(real));
  cudaCheckErrors("cuda malloc5 fail");
  cudaMalloc((void **)&d_z1, wsize*sizeof(real));
  cudaCheckErrors("cuda malloc6 fail");

  for (int i=0; i< dsize; i++){
    h_x[i] = 0.0f;
    h_x1[i] = 0.0f;
    c_x[i] = 0.0f;
    h_y[i] = (real)(rand()/(real)RAND_MAX);
    h_z[i] = (real)(rand()/(real)RAND_MAX);
    }


  cudaEvent_t t_start, t_stop, k_start, k_stop;
  cudaEventCreate(&t_start);
  cudaEventCreate(&t_stop);
  cudaEventCreate(&k_start);
  cudaEventCreate(&k_stop);
  cudaCheckErrors("event fail");

  // first test original GPU version

  cudaEventRecord(t_start);
  cudaMemcpy(d_x, h_x, dsize * sizeof(real), cudaMemcpyHostToDevice);
  cudaCheckErrors("cuda memcpy 1 fail");
  cudaMemcpy(d_y, h_y, dsize * sizeof(real), cudaMemcpyHostToDevice);
  cudaCheckErrors("cuda memcpy 2 fail");
  cudaMemcpy(d_z, h_z, dsize * sizeof(real), cudaMemcpyHostToDevice);
  cudaCheckErrors("cuda memcpy 3 fail");


  cudaEventRecord(k_start);
  vectorStepAdd2(d_x, d_y, d_z, alpha, beta, wsize, xstep, ystep, zstep);
  cudaEventRecord(k_stop);

  cudaMemcpy(h_x, d_x, dsize * sizeof(real), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cuda memcpy 4 fail");
  cudaEventRecord(t_stop);
  cudaEventSynchronize(t_stop);
  cudaEventElapsedTime(&et, t_start, t_stop);
  printf("GPU original version total elapsed time is: %f ms.\n", et);
  cudaEventElapsedTime(&et, k_start, k_stop);
  printf("GPU original kernel elapsed time is: %f ms.\n", et);

  //now test CPU version

  cudaEventRecord(t_start);
  vectorStepAdd3(c_x, h_y, h_z, alpha, beta, wsize, xstep, ystep, zstep);
  cudaEventRecord(t_stop);
  cudaEventSynchronize(t_stop);
  cudaEventElapsedTime(&et, t_start, t_stop);
  printf("CPU version total elapsed time is: %f ms.\n", et);
  for (int i = 0; i< dsize; i++)
    if (fabsf((float)(h_x[i]-c_x[i])) > TOL) {
      printf("cpu/gpu results mismatch at i = %d, cpu = %f, gpu = %f\n", i, c_x[i], h_x[i]);
      return 1;
      }


  // now test improved GPU version

  cudaEventRecord(t_start);
//  cudaMemcpy2D(d_x1, sizeof(real),  h_x, xstep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice);
//  cudaCheckErrors("cuda memcpy 5 fail");
  cudaMemcpy2D(d_y1, sizeof(real),  h_y, ystep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice);
  cudaCheckErrors("cuda memcpy 6 fail");
  cudaMemcpy2D(d_z1, sizeof(real),  h_z, zstep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice);
  cudaCheckErrors("cuda memcpy 7 fail");

  cudaEventRecord(k_start);
  vectorStepAdd2i(d_x1, d_y1, d_z1, alpha, beta, wsize);
  cudaEventRecord(k_stop);

  cudaMemcpy2D(h_x1, xstep*sizeof(real), d_x1, sizeof(real), sizeof(real), wsize, cudaMemcpyDeviceToHost);
  cudaCheckErrors("cuda memcpy 8 fail");
  cudaEventRecord(t_stop);
  cudaEventSynchronize(t_stop);
  cudaEventElapsedTime(&et, t_start, t_stop);
  printf("GPU improved version total elapsed time is: %f ms.\n", et);
  cudaEventElapsedTime(&et, k_start, k_stop);
  printf("GPU improved kernel elapsed time is: %f ms.\n", et);

  for (int i = 0; i< dsize; i++)
    if (fabsf((float)(h_x[i]-h_x1[i])) > TOL) {
      printf("gpu/gpu improved results mismatch at i = %d, gpu = %f, gpu imp = %f\n", i, h_x[i], h_x1[i]);
      return 1;
      }

  printf("Results:i   CPU     GPU     GPUi \n");
  for (int i = 0; i< 20*xstep; i+=xstep)
    printf("    %d         %f      %f     %f    %f    %f\n",i, c_x[i], h_x[i], h_x1[i]);


  return 0;
}

如前所述，我仍然无法超越 CPU 时间，我将这归因于我自己缺乏编码技能，或者这个操作根本没有足够的计算复杂性，无法在 GPU 上发挥作用。不过这里有一些示例结果:

GPU original version total elapsed time is: 13.352256 ms.
GPU original kernel elapsed time is: 0.195808 ms.
CPU version total elapsed time is: 2.599584 ms.
GPU improved version total elapsed time is: 4.228288 ms.
GPU improved kernel elapsed time is: 0.027392 ms.
Results:i   CPU     GPU     GPUi
    0         0.617285      0.617285     0.617285
    47         0.554522      0.554522     0.554522
    94         0.104245      0.104245     0.104245
....

我们可以看到，与原始内核相比，改进后的内核总体上减少了大约 3 倍，这几乎都是由于数据复制时间的减少。数据复制时间的减少是由于使用改进的 2D memcpy，我们只需要复制我们实际使用的数据。 (没有页面锁定内存，这些数据传输时间大约是原来的两倍)。我们还可以看到内核计算时间比原始内核的 CPU 计算快约 10 倍，比改进内核的 CPU 计算快约 100 倍。然而，考虑到数据传输时间，我们无法克服 CPU 速度。

最后一个评论是 cudaMemcpy2D 操作的“成本”仍然很高。对于减少 100 倍的 vector 大小，我们只看到复制时间减少了 3 倍。因此，跨步访问仍然是使用 GPU 的一种相对昂贵的方式。如果我们只是简单地传输包含 50,000 个连续元素的 vector ，我们预计复制时间几乎线性减少 100 倍(与包含 5000000 个元素的原始复制 vector 相比)。这意味着复制时间将少于 1 毫秒，并且我们的 GPU 版本将比 CPU 更快，至少在这个简单的单线程 CPU 代码中是这样。

关于c++ - cuda 上的 vector 步长加法较慢，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/15194798/