CUDA 学习(2)——CUDA 介绍

GeForce 256 是英伟达 1999 年开发的第一个 GPU，最初用作显示器上渲染高端图形，只用于像素计算。

在早期，OpenGL 和 DirectX 等图形 API 是与 GPU 唯一的交互方式。后来，人们意识到 GPU 除了用于渲染图形图像外，还可以做其他的数学计算，但是 OpenGL 和 DirectX 等图形 API 的交互方式比较复杂，不利于程序员设计 GPU 计算程序，这促成了 CUDA 编程框架的开发，它提供了一种与 GPU 交互的简单而高效的方式。

1 CUDA 环境搭建

必要的条件：

• Nvidia 的 GPU
• Nvidia 的显卡驱动
• 标准的 C 编译器
• CUDA 开发工具

建立好 CUDA 开发环境之后，可以通过以下命令进行检查：

nvidia-smi
nvcc --version

2 CUDA 编程模型简述

2.1 基本概念

• thread：一个 CUDA 的并行程序会被以许多个 thread 来执行
• block: 多个线程组成一个线程块（Block），同一个 block 的线程会被调度到同一个 SM 上，即同一个 block 的 thread 可以进行同步并可用 SM 上的 share memory 通信，不同 block 的 thread 无法通信
• grid: CUDA 的一个函数叫做一个 kernel，一个 kernel 会发起大量执行相同指令的线程

CUDA 编程软件层次：

这三个概念是 CUDA 编程中最核心的，知道这些，就已经可以写 cuda 代码了，进一步了解硬件结构可以帮助我们更好地对 cuda 代码深度优化。

2.2 helloGPU

尝试编写一个 cuda 程序 hello-gpu.cu，让 GPU 输出Hello World！。

#include <stdio.h>void helloCPU() {printf("Hello World!  --From CPU\n");
}__global__ void helloGPU() {printf("Hello World!  --From GPU\n");
}int main() {helloCPU();helloGPU<<<1, 1>>>();cudaDeviceSynchronize();
}

可以看到 cuda 程序和普通的 c 语言非常相似，也存在一些不一样的地方：

• __global__：定义这是一个 cuda 的 kernel 函数，从主机 host 发起并在设备 device 上执行。
• <<<1, 1>>>：定义 block 和 threads，这里表示发起 1 个 block，每个 block 里有 1 个线程
• cudaDeviceSynchronize：与许多 C/C++ 代码不同，核函数启动方式为异步：CPU 代码将继续执行而无需等待核函数完成启动。调用 CUDA 运行时提供的函数 cudaDeviceSynchronize 将导致主机 (CPU) 代码暂作等待，直至设备 (GPU) 代码执行完成，才能在 CPU 上恢复执行。

写好 cuda 代码后，可以使用 nvcc 对代码进行编译与执行：

nvcc -arch=sm_75 -o hello-gpu hello-gpu.cu -run# Hello World!  --From CPU
# Hello World!  --From GPU

说明：

• nvcc 是使用 nvcc 编译器的命令行命令。
• 将 xxx.cu 作为文件传递以进行编译。
• -o标志用于指定编译程序的输出文件。
• arch 标志表示该文件必须编译为哪个架构类型。本示例中，sm_75 将用于专门针对本实验运行的 NVIDIA GeForce GTX 2080 Ti 进行编译。具体参考：https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#options-for-steering-gpu-code-generation
• 为方便起见，提供 run 标志将执行已成功编译的二进制文件。

从上面的程序，可以知道 GPU 的工作任务是由 CPU 触发的，GPU 自身是无法独立工作的。

cuda 程序整体的工作流程是 CPU 将需要执行的任务异步地交给 GPU，再由 GPU 进行调度，最后再将计算结果同步给 CPU。

假设想要 GPU 发送 66 个Hello World，可以简单地修改 blocks 和 ThreadsPerBlock 的数量，即可实现这项功能：

#include <stdio.h>void helloCPU() {printf("Hello World!  --From CPU\n");
}__global__ void helloGPU() {printf("Hello World!  --From GPU\n");
}int main() {helloCPU();helloGPU<<<6, 11>>>();cudaDeviceSynchronize();
}

以上代码则发起了 6 个 block，每个 block 里有 11 个线程。当然，也可以改成helloGPU<<<1, 66>>>();，发起了一个 block，这个 block 里有 66 个线程。

3 Warp

具体怎么设置发起 blocks 和 ThreadsPerBlock 完全由程序员自己设置，而发起后这些 block 和线程在 GPU 中如何调度则由 GPU 内部硬件控制，不被程序员所操作。为了更合理地设置 blocks 和 ThreadsPerBlock，还需要了解 GPU 中的调度策略。

• 首先是 blocks 的调度：同一个 blocks 会被调度到同一个 SM，不同的 blocks 不保证在同一 SM。

为了更好地进行调度，blocks 数可以设置为 GPU 中 SM 的整数倍。由于 SM 上的计算单元是有限的，同一个 blocks 中的 threads 会被划分成多个 warp，一个 warp 才是 GPU 调度与执行的基本单元。

一般来说，一个 warp 是 32 个线程（尽量是每个 SM 中的流处理器数量的整倍数？），所以 ThreadsPerBlock 一般会设置成 32 的整数倍，可以让资源利用率更高。

了解了 GPU 中的调度逻辑，编写 cuda 程序时我们就可以根据手中的 GPU 硬件配置，合理地设置 blocks 和 ThreadsPerBlock 这两个参数。当前 GPU 硬件配置有很多内容，在初学 CUDA 编程中应该关注到的是 GPU 上 SM 数量，warp size，每个 block 的最大线程数，每个 SM 最大 block 数。通过这段代码将 GPU 硬件信息打印出来：

#include <stdio.h>
#include <iostream>int main() {int dev = 0;cudaDeviceProp devProp;cudaGetDeviceProperties(&devProp, dev);std::cout << "使用 GPU device " << dev << ": " << devProp.name << std::endl;std::cout << "SM 的数量：" << devProp.multiProcessorCount << std::endl;int warpSize = devProp.warpSize;std::cout << "Warp size: " << warpSize << std::endl;std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;std::cout << "每个线程块的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;std::cout << "每个 SM 的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;std::cout << "每个 SM 的最大 block 数：" << devProp.maxThreadsPerMultiProcessor / warpSize << std::endl;std::cout << "每个 SM 的寄存器数量：" << devProp.regsPerMultiprocessor << std::endl;
}

编译梦并运行：

nvcc -o get_gpu_hwinfo get_gpu_hwinfo.cu -run使用 GPU device 0: NVIDIA GeForce RTX 2080 Ti
SM 的数量：68
Warp size: 32
每个线程块的共享内存大小：48 KB
每个线程块的最大线程数：1024
每个 SM 的最大线程数：1024
每个 SM 的最大 block 数：32
每个 SM 的寄存器数量：65536

举一个简单的例子来说明如何根据硬件配置合理分配资源：

假设一个 SM 上有 8192 个寄存器，程序员每个 block 设置了 256 个线程。

假设每个线程会占用 10 个寄存器，那么一个 block 中的线程会占用 256*10=2560 个寄存器，8192/2560=3.2，即一个 SM 可以同时加载 3 个 block 正常运行。

假设每个线程会占用 11 个寄存器，那么一个 block 中的线程会占用 256*11=2816 个寄存器，8192/2816=2.9，即一个 SM 只能加载 2 个 block，一个 SM 上硬件资源就跑不满，会造成资源浪费。

blocks 调度到 SM 上：

block 被切分成 wrap：

由于 GPU 没有复杂的控制单元，在 warp 中所有线程都会执行相同的指令，这意味着在遇到分支时，warp 需要一些特殊的处理。

如下图所示，当遇到分支时，warp 中 32 个线程也许有些线程满足条件，有些线程不满足条件，但一个 warp 中所有线程执行指令的时序是一致的，不满足分支条件的线程必须等待需要执行指令的其他线程，这也意味着分支指令会影响 GPU 的运行效率，在程序设计时应该尽量少用，或者在写分支条件时尽可能保证一个 warp 中所有线程同时满足条件或者同时不满足条件。