我的第一份CUDA代码

作者丨xcyuyuyu

来源丨https://zhuanlan.zhihu.com/p/507678214

编辑丨极市平台

注：内容仅做学术分享之用，若涉及侵权等行为，请联系删除，万分感谢！

1. 前言

这是一份简单的CUDA编程入门，主要参考英伟达的官方文档进行学习，本人也是刚开始学习，如有表述错误，还请指出。官方文档链接如下：

https://developer.nvidia.com/blog/even-easier-introduction-cuda/

本文先从一份简单的C++代码开始，然后逐步介绍如何将C++代码转换为CUDA代码，以及对转换前后程序的运行时间进行对比，本文代码放在我的github中，有需要可以自取。

https://github.com/xcyuyuyu/My-First-CUDA-Code

本文所使用的CPU为i7-4790，GPU为GTX 1080，那就开始吧。

2. 一份简单的C++代码

首先是一份简单的C++代码，主要的运行函数为add函数，该函数实现功能为30M次的for循环，每次循环进行一次加法。

// add.cpp
#include <iostream>
#include <math.h>
#include <sys/time.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<<25; // 30M 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;
  }

  struct timeval t1,t2;
  double timeuse;
  gettimeofday(&t1,NULL);
  // Run kernel on 30M elements on the CPU
  add(N, x, y);
  gettimeofday(&t2,NULL);
  timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000.0;

  std::cout << "add(int, float*, float*) time: " << timeuse << "ms" << std::endl;
  // 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;
}

编译以及运行代码：

g++ add.cpp -o add
./add

不出意外的话，你应该得到下面的结果：

第一行表示add函数的运行时间，第二行表示每个for循环里的计算是否与预期结果一致。

这个简单的C++代码在CPU端运行，运行时间为85ms，接下来介绍如何将主要运算的add函数迁移至GPU端。

3. 把C++代码改成CUDA代码

将C++代码改为CUDA代码，目的是将add函数的计算过程迁移至GPU端，利用GPU的并行性加速运算，需要修改的地方主要有3处：

1.首先需要做的是将add函数变为GPU可运行函数，在CUDA中称为kernel，为此，仅需将变量声明符添加到函数中，告诉 CUDA C++ 编译器这是一个在 GPU 上运行并且可以从 CPU 代码中调用的函数。

__global__ 
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

那么修改后的add函数的调用也比较简单，仅需要在add函数名后面加上三角括号语法<<<i,j>>>指定CUDA内核启动即可，<<<i,j>>>称为执行配置（execution configuration），用于配置程序运行时的线程，后续会讲到，目前先将其设置为<<<i,j>>>：

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

2. 那么为了在GPU进行计算，需要在GPU上分配可访问的内存。CUDA中通过Unified Memory（统一内存）机制来提供可同时供GPU和CPU访问的内存，使用cudaMallocManaged()函数进行分配：

cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

同时，在程序最后使用cudaFree()进行内存释放：

cudaFree(x);
cudaFree(y);

其实就相当于C++中的new跟delete。

3. add函数在GPU端运行之后，CPU需要等待cuda上的代码运行完毕，才能对数据进行读取，因为CUDA内核启动时并未对CPU的线程进行固定，需要使用cudaDeviceSynchronize()函数进行同步。

4. 整体的程序如下所示：

// add.cu
#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
// __global__ 变量声明符，作用是将add函数变成可以在GPU上运行的函数
// __global__ 函数被称为kernel，
// 在 GPU 上运行的代码通常称为设备代码（device code），而在 CPU 上运行的代码是主机代码（host code）。
__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<<25;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  // 内存分配，在GPU或者CPU上统一分配内存
  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
  // execution configuration, 执行配置
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  // CPU需要等待cuda上的代码运行完毕，才能对数据进行读取
  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;
}

使用nvcc对程序进行编译并运行：

nvcc add.cu -o add_cuda 
./add_cuda

或者使用nvprof进行速度测试：

nvprof ./add_cuda

不出意外的话，你会得到以下输出：

框出来的就是add函数在GPU端的运行时间，为4s。没错，就是比CPU端85ms还要慢，那还学个锤子。

4. 使用CUDA代码并行运算

好的回过头看看，问题出现在这个执行配置 <<<i,j>>> 上。不急，先看一下一个简单的GPU结构示意图，按照层次从大到小可将GPU按照 grid -> block -> thread划分，其中最小单元是thread，并行的本质就是将程序的计算模块拆分成多个小模块扔给每个thread并行计算。

再看一下前面执行配置 `<<<i,j>>>` 的含义，`<<<i,j>>>` 应该写成 `<<<numBlocks, blockSize>>>` ，即表示函数运行时使用的block数量以及每个block的大小，前面我们将其设置为`<<<1,1>>>` ，说明程序是单线程运行的，那当然慢了~~。下面我们以单个block为例，将其改为`<<<1,256>>>`，add函数也需要适当修改：

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x; // threadIdx.x表示当前在第几个thread上运行
  int stride = blockDim.x; // blockDim.x表示每个block的大小
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

修改的部分也比较好理解，不赘述了，接下来运行看看结果：

你看，开始加速了吧，4s加速到了77ms。

那么，`<<<numBlocks, blockSize>>>` 的两个参数应该怎么设置好呢。首先，CUDA GPU 使用大小为 32 的倍数的线程块运行内核，因此 `blockSize` 的大小应该设置为32的倍数，例如128、256、512等。确定 `blockSize` 之后，可以根据for循环的总个数`N`确定 `numBlock` 的大小（注意四舍五入的误差）：

int numBlock = (N + blockSize - 1) / blockSize;

当然因为变成了多个`block`，所以此时add函数需要再改一下：

__global__ 
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i+=stride)
    y[i] = x[i] + y[i];
}

这里index跟stride的计算可以参考上面GPU结构图以及下面的图（图取自An Even Easier Introduction to CUDA | NVIDIA Technical Blog），自行推算，较好理解。

搞定之后再编译运行一下：

看看，又加速了不是，通过提升并行度而加速，相比于CPU端（85ms）加速了接近一倍左右。

5. 结论

以上仅是一份简单的CUDA入门代码，看起来还算比较简单，不过继续深入肯定有更多的坑，期待后面有时间继续学习。

本文代码：

GitHub - xcyuyuyu/My-First-CUDA-Code: The introduction to cuda, a simple and easy cuda project

https://github.com/xcyuyuyu/My-First-CUDA-Code

参考文献

[1] An Even Easier Introduction to CUDA | NVIDIA Technical Blog（https://developer.nvidia.com/blog/even-easier-introduction-cuda/）

推荐阅读

• 西电IEEE Fellow团队出品！最新《Transformer视觉表征学习全面综述》

• 来自谷歌、Meta工程师联合出品的模型部署秘籍（附pdf下载）

• 如何做好科研？这份《科研阅读、写作与报告》PPT，手把手教你做科研

• 最新 2022「深度学习视觉注意力 」研究概述，包括50种注意力机制和方法！

• 【重磅】斯坦福李飞飞《注意力与Transformer》总结，84页ppt开放下载！

• 分层级联Transformer！苏黎世联邦提出TransCNN: 显著降低了计算/空间复杂度！

• 清华姚班教师劝退文：读博，你真的想好了吗？

• 2021李宏毅老师最新40节机器学习课程！附课件+视频资料

欢迎大家加入DLer-计算机视觉技术交流群！

大家好，这是DLer-计算机视觉技术交流群，群里会第一时间发布计算机视觉相关技术的论文前沿论文以及工程项目实践等等，主要方向有：图像分类、Transformer、目标检测、目标跟踪、点云与语义分割、GAN、超分辨率、人脸检测与识别、动作行为与时空运动、模型压缩和量化剪枝、迁移学习、人体姿态估计等内容。

进群请备注：研究方向+学校/公司+昵称（如图像分类+上交+小明）

👆 长按识别，邀请您进群