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 中所有线程同时满足条件或者同时不满足条件。