1. GPU简介
GPU与CPU的主要区别在于：
- CPU拥有少数几个快速的计算核心，而GPU拥有成百上千个不那么快速的计算核心。
- CPU中有更多的晶体管用于数据缓存和流程控制，而GPU中有更多的晶体管用于算数逻辑单元。
所以，GPU依靠众多的计算核心来获得相对较高的并行计算性能。
一块单独的GPU无法独立地完成所有计算任务，它必须在CPU的调度下才能完成特定任务，因此当我们讨论GPU计算时，其实指的是CPU+GPU的异构计算。通常将起控制作用的CPU称为主机(host)，起加速作用的GPU称为设备(device)，它们之间一般采用PCIe总线连接。
NVIDIA公司出品的GPU中，支持CUDA(Compute Unified Device Architecture)编程的系列如下：
- Tesla系列：主要用于科学计算。
- Quadro系列：主要用于专业绘图设计。
- GeForce系列：主要用于游戏与娱乐。
- Jetson系列：主要用于嵌入式设备。
每款GPU都有一个计算能力(compute capability)，写为形如X.Y的形式。计算能力决定了GPU硬件所支持的功能，它与性能不是简单的正比关系。下表列出了部分计算能力及其架构代号与发布年份，详细的GPU计算能力信息可以查阅官方网站：https://developer.nvidia.com/cuda-gpus
| 计算能力 | 架构代号 | 发布时间 |
|---|---|---|
| X = 1 | Tesla（特斯拉） | 2006 |
| X = 2 | Fermi（费米） | 2010 |
| X = 3 | Kepler（开普勒） | 2012 |
| X = 5 | Maxwell（麦克斯韦） | 2014 |
| X = 6 | Pascal（帕斯卡） | 2016 |
| X.Y = 7.0 | Volta（伏特） | 2017 |
| X.Y = 7.5 | Turing（图灵） | 2018 |
| X.Y = 8.6 | Ampere（安培） | 2020 |
| X.Y = 8.9 | Ada（阿达） | 2022 |
表征GPU性能的一个重要参数是每秒浮点运算次数(floating-point operations per second，FLOPS)，其数值通常在1012量级，即teraFLOPS(TFLOPS)。浮点运算有单精度和双精度之分，双精度浮点运算速度通常小于单精度浮点运算速度，对于Tesla系列GPU来说其比例一般是1/2左右，对于GeForce系列GPU来说其比例一般是1/32左右。另一个影响GPU性能的重要参数是显存带宽，它限制了显卡芯片与显存之间的数据交换速率。
CUDA官方文档包含了安装指南、编程指南、API手册、工具介绍等内容，网址是：https://docs.nvidia.com/cuda/
安装完CUDA开发工具后，可以在命令行中执行nvidia-smi来查看设备信息。
PS C:\> nvidia-smi
Wed Apr 19 21:53:50 2023
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 531.14 Driver Version: 531.14 CUDA Version: 12.1 |
|-----------------------------------------+----------------------+----------------------+
| GPU Name TCC/WDDM | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+======================+======================|
| 0 NVIDIA GeForce RTX 3060 L... WDDM | 00000000:01:00.0 On | N/A |
| N/A 47C P8 13W / N/A| 737MiB / 6144MiB | 1% Default |
| | | N/A |
+-----------------------------------------+----------------------+----------------------+
+---------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=======================================================================================|
| 0 N/A N/A 9236 C+G C:\Windows\explorer.exe N/A |
+---------------------------------------------------------------------------------------+
2. 运行时API
CUDA提供了两层API供程序员使用，分别是CUDA驱动(driver)API和CUDA运行时(runtime)API。其中，驱动API较为底层，它虽然编程接口更加灵活但编程难度更高，例如cuCtxCreate()等cu开头的函数；运行时API则在驱动API的基础上进行了封装，更加容易使用，例如cudaMalloc()等cuda开头的函数。CUDA运行时API中没有显式初始化设备的函数，在第一次调用一个和设备管理/版本查询功能无关的运行时API时，设备将自动初始化。
下面是一段利用CUDA运行时API进行数组相加的程序，它体现了一个CUDA程序的基本编程框架。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cstdlib>
// CUDA核函数的定义
void __global__ add(const double* x, const double* y, double* z)
{
const int n = blockDim.x * blockIdx.x + threadIdx.x;
z[n] = x[n] + y[n];
}
int main()
{
// 分配主机内存、初始化数据
const int N = 100000000;
const int M = sizeof(double) * N;
double* h_x = (double*)malloc(M);
double* h_y = (double*)malloc(M);
double* h_z = (double*)malloc(M);
for (int n = 0; n < N; ++n)
{
h_x[n] = 1.23;
h_y[n] = 4.56;
}
// 分配设备内存、把主机数据复制到设备中
double* d_x, * d_y, * d_z;
cudaMalloc((void**)&d_x, M);
cudaMalloc((void**)&d_y, M);
cudaMalloc((void**)&d_z, M);
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
// 调用核函数在设备中进行计算
const int block_size = 128;
const int grid_size = N / block_size;
add<<<grid_size, block_size>>>(d_x, d_y, d_z);
// 把设备数据复制到主机中
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
// 释放主机和设备的内存
free(h_x);
free(h_y);
free(h_z);
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_z);
return 0;
}
3. 内存操作
在CUDA中，设备内存的动态分配可由cudaMalloc函数实现。第一个参数p是待分配设备内存指针的地址，第二个参数s是待分配内存的字节数。
cudaError_t cudaMalloc(void **p, size_t s);
用cudaMalloc申请的设备内存需要用cudaFree函数释放。参数p是待释放设备内存的指针。
cudaError_t cudaFree(void *p);
主机内存与设备内存之间的数据传递可以使用cudaMemcpy函数。参数dst是目标地址，src是源地址，count是复制数据是字节数，kind表示数据传递的方向。
enum cudaMemcpyKind
{
cudaMemcpyHostToHost = 0, /**< Host -> Host */
cudaMemcpyHostToDevice = 1, /**< Host -> Device */
cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
4. 核函数
主机对设备的调用是通过核函数(kernel function)来实现的，核函数与C++函数的主要区别是：
- 核函数需要被限定词
__global__修饰。 - 核函数的返回类型必须是
void。
核函数的线程(thread)往往组织为线程块(thread block)，所有线程块构成了一个网格(grid)。网格大小(grid size)是指网格中包含的线程块个数，线程块大小(block size)是指线程块中包含的线程个数。调用核函数时，需要在三括号<<<>>>中指明网格大小以及线程块大小，即<<<网格大小, 线程块大小>>>（也可以理解为<<<线程块个数, 每个线程块包含的线程个数>>>），核函数中的总线程数就等于网格大小乘以线程块大小。
网格大小与线程块大小既可以是一维的，也可以是二维或者三维的。对于多维的情况，需要用dim3结构体来表示，其中x维度在逻辑上是最内层的，即变化最快的。网格大小在x、y、z方向上的最大值分别是231-1、65535、65535；线程块大小在x、y、z方向上的最大值分别是1024、1024、64，并且三者的乘积不能大于1024，也就是说一个线程块最多只能拥有1024个线程。
//简化的uint3结构体定义
struct uint3
{
unsigned int x, y, z;
};
//简化的dim3结构体定义
struct dim3
{
unsigned int x, y, z;
constexpr dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
constexpr dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
constexpr operator uint3(void) const { return uint3{x, y, z}; }
};
//三维的网格与线程块
dim3 grid_size(4, 3, 2);
dim3 block_size(4, 3, 2);
kernel_func<<<grid_size, block_size>>>();
//二维的网格与线程块
dim3 grid_size(4, 3); //等价于dim3 grid_size(4, 3, 1);
dim3 block_size(4, 3); //等价于dim3 block_size(4, 3, 1);
kernel_func<<<grid_size, block_size>>>();
//一维的网格与线程块
dim3 grid_size(4); //等价于dim3 grid_size(4, 1, 1);
dim3 block_size(4); //等价于dim3 block_size(4, 1, 1);
kernel_func<<<grid_size, block_size>>>(); //一维情况下三括号中也可以直接填数字，例如kernel_func<4, 4>();
在核函数内部，可以分别通过dim3类型的内建变量gridDim和blockDim来获取网格大小与线程块大小：gridDim.x、gridDim.y、gridDim.z分别表示网格大小在x、y、z维度上的值；blockDim.x、blockDim.y、blockDim.z分别表示线程块大小在x、y、z维度上的值。
类似地，核函数中也分别定义了uint3类型的内建变量blockIdx和threadIdx来表示当前线程块的标号以及线程的标号，blockIdx.x的取值范围是0到gridDim.x - 1，threadIdx.x的取值范围是0到blockDim.x - 1，y维度和z维度的情况可以以此类推。
此外，还有int型的内建变量warpSize表示线程束(thread warp)的大小。一个线程块中连续warpSize个线程构成一个线程束，具体地说，一个线程块中第0~31个线程属于第0个线程束，第32~63个线程属于第1个线程束。对于目前所有的GPU架构来说，warpSize的值都是32。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cstdio>
__global__ void hello_from_gpu()
{
const int bx = blockIdx.x;
const int tx = threadIdx.x;
const int ty = threadIdx.y;
printf("block-%d and thread-(%d, %d)!\n", bx, tx, ty);
}
int main(void)
{
const dim3 block_size(2, 3);
hello_from_gpu<<<2, block_size>>>();
cudaDeviceSynchronize();
return 0;
}
/*
线程块的计算是相互独立的，以下是一种可能的输出情况，有可能block-0先完成计算，也有可能block-1先完成计算
block-1 and thread-(0, 0)
block-1 and thread-(1, 0)
block-1 and thread-(0, 1)
block-1 and thread-(1, 1)
block-1 and thread-(0, 2)
block-1 and thread-(1, 2)
block-0 and thread-(0, 0)
block-0 and thread-(1, 0)
block-0 and thread-(0, 1)
block-0 and thread-(1, 1)
block-0 and thread-(0, 2)
block-0 and thread-(1, 2)
*/
5. 设备函数
核函数可以调用不带执行配置的自定义函数，这样的自定义函数称为设备函数(device function)。设备函数可以有返回值。
__global__修饰的函数称为核函数，一般由主机调用，在设备中执行。__device__修饰的函数称为设备函数，只能由核函数或其它设备函数调用，在设备中执行。__host__修饰的函数就是主机端的普通C++函数，由主机调用，在主机中执行。对于主机端的函数，该修饰符可以省略。之所以提供这样的修饰符是因为有时可以同时用__host__和__device__修饰同一个函数，使得该函数既是一个普通C++函数又是一个设备函数，这样做可以减少冗余代码，编译器将针对主机和设备分别编译该函数。
double __device__ add_device(const double x, const double y)
{
return x + y;
}
void __global__ add(const double* x, const double* y, double* z, const int N)
{
const int n = blockDim.x * blockIdx.x + threadIdx.x;
z[n] = add_device(x[n], y[n]);
}
不能同时用__device__和__global__修饰一个函数，即不能将一个函数同时定义为设备函数与核函数。同理也不能同时用__host__和__global__修饰一个函数，即不能将一个函数同时定义为主机函数与核函数。
可以使用__noinline__建议一个设备函数为非内联函数，也可以使用__forceinline__建议一个设备函数为内联函数。
6. 错误检测
所有CUDA运行时API函数都以cuda作为前缀，而且都返回一个cudaError_t类型的值表示错误信息，返回值为cudaSuccess时表示成功调用了API函数。可以使用cudaGetErrorString函数来将错误码转换成错误的文字描述。
#define CHECK(call) \
do \
{ \
const cudaError_t error_code = call; \
if (error_code != cudaSuccess) \
{ \
printf("CUDA Error:\n"); \
printf(" File: %s\n", __FILE__); \
printf(" Line: %d\n", __LINE__); \
printf(" Error code: %d\n", error_code); \
printf(" Error text: %s\n", cudaGetErrorString(error_code)); \
exit(1); \
} \
} while (0)
CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));
//这里故意把cudaMemcpyHostToDevice写成cudaMemcpyDeviceToHost，得到的错误信息可能如下：
//CUDA Error:
// File: test.cu
// Line: 42
// Error code: 11
// Error text: invalid argument
由于核函数没有返回值，因此没法直接使用上述方法来捕捉错误。为了捕捉核函数可能发生的错误，可以在调用核函数之后使用cudaGetLastError来获取错误信息。
add<<<256, 1280>>>();
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
//线程块大小的最大值是1024，上面故意写成1280，得到的错误信息可能如下：
//CUDA Error:
// File: test.cu
// Line: 42
// Error code: 9
// Error text: invalid configuration argument
7. NVCC
一般来说，一个CUDA程序既有标准的C++代码，也有不属于标准C++的CUDA代码。CUDA程序编译器nvcc在编译一个CUDA程序时，会将标准C++代码交给C++编译器（例如g++或cl）去处理，它自己则负责编译CUDA代码的部分。CUDA程序源文件的扩展名通常是.cu，不带任何参数选项地使用nvcc编译一个源文件的指令如下：
nvcc hello.cu
nvcc的编译过程分为两个阶段：
- 首先将设备代码编译为一种面向虚拟架构的PTX(parallel thread execution)伪汇编代码。
- 然后将PTX代码编译为面向实际架构的cubin目标代码。
对于nvcc编译器，-arch选项指定了第一阶段使用什么虚拟架构，-code选项指定了第二阶段使用什么实际架构，实际架构的计算能力必须大于等于虚拟架构，例如：
-arch=compute_XY -code=sm_ZW
上述选项生成的可执行文件，只能在计算能力为Z.W的GPU上运行。为了让编译出来的可执行文件能在更多的GPU上运行，nvcc也提供了即时编译(just in time compilation)的机制，可以在运行时从其中保留的PTX代码临时编译出一个cubin目标代码。要在文件中保留PTX代码，就需要用如下方式指定所保留PTX代码的虚拟架构，这里的两个计算能力都是虚拟架构的计算能力，必须完全一致：
-arch=compute_XY -code=compute_XY
nvcc也支持使用-gencode选项来执行多组计算能力，例如：
-gencode arch=compute_35,code=sm_35
-gencode arch=compute_50,code=sm_50
-gencode arch=compute_60,code=sm_60
-gencode arch=compute_60,code=compute_60
上述选项生成的目标文件将会包含：
- 基于
compute_35PTX代码产生的sm_35目标代码 - 基于
compute_50PTX代码产生的sm_50目标代码 - 基于
compute_60PTX代码产生的sm_60目标代码 compute_60PTX代码
在目标文件运行时，若目标代码可直接运行在GPU上，则直接运行目标代码；否则，若文件中包含PTX代码，则显卡驱动会尝试将PTX代码动态编译为目标代码然后执行。
在CMakeLists.txt中添加CUDA支持的示例如下：
cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
enable_language(CUDA) # 也可以在project命令中添加CUDA支持，例如：project(TestCUDA LANGUAGES CXX CUDA)
set(CMAKE_CUDA_ARCHITECTURES 52) # https://cmake.org/cmake/help/latest/variable/CMAKE_CUDA_ARCHITECTURES.html