CUDA学习笔记 —— （三）GPU计算性能与线程关系

概述

我们通过一个例子，相同的计算，我们分别在CPU，GPU单线程，GPU单block多线程，GPU多block多线程来运行对比。看GPU是如何大幅提升运算能力的。
参考文章https://developer.nvidia.com/blog/even-easier-introduction-cuda/

CPU运行

一个最简单的例子：将两个数组中的元素相加放到第二个数组中。
我们使用CPU来运行

#include <iostream>
#include <math.h>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 1M elements

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  delete [] x;
  delete [] y;

  return 0;
}

GPU运行单线程

下面作为对比，我们看如何将计算放入到GPU中来完成
下面注释中有具体说明每个函数使用情况。

#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  //这个Unified Memory内存空间使用cudaMallocManaged创建，该内存可以在CPU和GPU之间共享
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  //使用1block，1个thread来执行kernel
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  //等GPU执行完毕，写会内存，因为上面函数是异步过程，所以如果不等GPU写完则会出现数据竞争情况。
  cudaDeviceSynchronize();

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  cudaFree(x);
  cudaFree(y);
  
  return 0;
}

GPU 1block 多线程运行

我们修改运行的线程数量

add<<<1, 256>>>(N, x, y);

然后修改kernel函数
因为线程数量为256个，总共有n个数需要计算
这里将N个数平均分到256个线程去运行

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x; //0
  int stride = blockDim.x; //线程数
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

多block多线程

因为GPU有多个并行的处理器，叫做SM。而每个SM可以并发运行blocks。
所以我们开启多个block来最大化GPU运行能力。

int blockSize = 256; //开启256个线程
int numBlocks = (N + blockSize - 1) / blockSize; //block数为N/256，但是要注意四舍五入
add<<<numBlocks, blockSize>>>(N, x, y);

我们可以调用变量来获取线程id，一维对应关系如下

还需要改造一下kernel函数，让每个线程对应于一个数组中的元素

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x; //线程号
  int stride = blockDim.x * gridDim.x; //总线程数
  //这里用一个for循环来处理，我们叫做grid-stride loop.
  //他是为了即使保证线程数量如果小于元素N数，可以均匀分配
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

结果比对

下面我们对比一下结果。
gpu 1block 1 thread : 15s
gpu 1block 256 thread :65.721ms
gpu Nblock 256 thread : 26.175ms

可以发现，并且执行时间极大的缩短了。