Hello CUDA

2024-5-28

CUDA

和所有编程语言的第一个 "hello world" 一样，我们开始 "hello CUDA"。

编写 CUDA 程序通常经历以下基本步骤：

1. 在 CPU 主机上初始化数据。

2. 在 GPU 设备上分配内存，将数据从 CPU 传输到 GPU。

3. 在 GPU 上执行编写的核函数。

4. 将处理完成的数据从 GPU 传回 CPU。

5. 在 CPU 上进行后处理。释放 GPU 和 CPU 内存。

如代码清单 1 所示，是我们的第一个 CUDA 程序。修饰符 __global__ 表明这个函数是一个核函数。核函数不能返回值，所以返回类型必须要是 void。

核函数的调用语法如下。和 C 语言函数调用语法不同的是，它有三对尖括号，里面指定执行配置。

  • Kernel_name <<< number_of_blocks, thread_per_block >>> (arguments);

比如此处的代码 13 行，第一个 1 表示核函数使用的 block 数量，第二个 1 表示每个 block 中的 thread 数量。所以此时只会打印一次 "Hello CUDA"。

代码清单 1 Hello CUDA
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5.  
  6. __global__ void hello_cuda()
  7. {
  8.     printf("Hello CUDA\n");
  9. }
  10.  
  11. int main()
  12. {
  13.     hello_cuda << <1, 1>> > ();
  14.  
  15.     cudaDeviceSynchronize();
  16.  
  17.     cudaDeviceReset();
  18.     return 0;
  19. }

这边的 printf 肯定不是之前认识的那个 C 库 printf 了。应该涉及打印字符的传输。

可以看到语法是在 C/C++ 基础上的扩展，入门编写比较方便。

我们再进一步理解核函数的执行配置。如图 1 所示，执行配置可以理解成一个三维笛卡尔坐标，里面的小方块就是一个 thread。多个 thread（即图中的相同颜色）组成多组 block。而多个 block 的排布方式叫做 grid。

图1 grid 和 block

在 CUDA 中，dim3 结构体专门用来定义三维变量，可以用于执行参数的指定。

让我们再做几个“习题”理解一下。如果我们想要图 2 的执行“布局”，即每个 block 里要 4 个 thread，并且要 8 个 block。

图2 一维举例

代码清单 2 是对应的实现，因为 dim3 的默认缺省值是 1，所以我们这边其他维度加以省略。运行看看，会打印 32 次 "Hello CUDA"。

代码清单 2 一维举例
  1. int main()
  2. {
  3.     dim3 block(4);
  4.     dim3 grid(8);
  5.  
  6.     hello_cuda << <grid, block>> > ();
  7.  
  8.     cudaDeviceSynchronize();
  9.  
  10.     cudaDeviceReset();
  11.     return 0;
  12. }

如图 3 所示，我们再看一个二维的例子。我们可以像在代码清单 3 中一样，指定对应的执行参数。

图3 二维举例
代码清单 3 一维举例
  1. int main()
  2. {
  3.     int nx = 16, ny = 4;
  4.  
  5.     dim3 block(8, 2);
  6.     dim3 grid(nx / block.x, ny / block.y);
  7.  
  8.     hello_cuda << <grid, block>> > ();
  9.  
  10.     cudaDeviceSynchronize();
  11.  
  12.     cudaDeviceReset();
  13.     return 0;
  14. }

最后了解一下 cudaDeviceSynchronize() 和 cudaDeviceReset() 函数。

cudaDeviceSynchronize 函数是阻塞操作，它会一直阻塞 CPU 主机代码的执行，直到 GPU 设备上的操作完成。

cudaDeviceReset 函数一般在 CUDA 程序的最后调用，用于资源的清理并恢复到初始状态。