1. 异构编程
CPU+高速总线（PCIE/NVLink）+GPU（用于数学/科学计算，有自己的操作系统）
三步走：
- 从 GPU 内存往 CPU 内存拷贝输入数据
- 加载 GPU 代码并执行，在片上缓存数据来提升性能
- 从 GPU 内存往 CPU 内存拷贝结果
并行计算：以 vector add 为例，计算 output 中第一个元素的过程和计算第二、三个元素的过程独立
2. GPU kernels: device code
函数定义
__global__ void mykernel(void) { }
- CUDA C++关键字/装饰器 global 表示这个函数运行在 GPU 上（给编译器的提示），通过 host code 或其它 device code 来调用
- NVCC：编译器。用于把源码划分为 host 和 device 部分
- device 函数由 NVIDIA 编译器来执行
- host 函数由 gcc，cl.exe 等标准 host 编译器执行
函数调用
mykernel<<<1, 1>>>();
- 三层尖括号表示调用 device code（也叫做 kernel launch），括号内的参数是 CUDA 核执行配置
- 第一个参数是 warp 数，第二个参数是每个 warp 包含的线程数
内存管理
- host 和 device 侧的内存是完全分隔的
- device 侧的指针指向 GPU 内存，通常传递给 device code，不在 host 侧解引用
- host 侧的指针指向 CPU 内存，通常不传给 device code，不在 device 侧解引用
- 特例：pinned pointers，ATS，managed memory
- 用于 device 侧内存管理的 API（用法和 C 语言 API 类似）
- cudaMalloc()
- cudaFree()
- cudaMemcpy()
- 这些 API 使用指针来分配、释放或者复制内存。指针只是一个数，它没有 metadata，不是正式的 C++ 对象。
3. 设备侧向量加法
__global__ void add(int *a, int *b, int *c){ c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x]; }
- grid 是线程和 block 的组合，是一种分层描述。grid - block（warp/warps) - threads
- blockIdx 是一个结构体/内置变量，有三个元素：.x，.y，.z。每个元素的索引从 0 开始，到 N-1 结束，其中 N 是 kernel launch 时传入的值。
#define N 512
int main(void){
int *a, *b, *c; // host copies of a, b, c
int *d_a, *d_b, *d_c; // device copies of a, b, c
int size = N * sizeof(int); // Alloc space for device copies of a, b, c
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_b, size);
cudaMalloc((void **)&d_c, size);
// Alloc space for host copies of a, b, c and setup input values
a = (int *)malloc(size);
random_ints(a, N);
b = (int *)malloc(size);
random_ints(b, N);
c = (int *)malloc(size);
// Copy inputs to device
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
// 可以不加 cudaMemcpyHostToDevice，但是加上相当于提供了一层校验
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
// Launch add() kernel on GPU with N blocks (N copies of add())
add<<<N, 1>>>(d_a, d_b, d_c);
// Copy result back to host
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
// Cleanup
free(a);
free(b);
free(c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
- cuda 开头的 API：cuda runtime API。如果执行错误，会返回错误码
- cuda 符合 C++ 2014 标准，但是不支持 C++ 标准库
4. CUDA 线程
- 术语：一个 block 可以被划分为并行线程
改写 add() 来使用并行线程，而非并行 blocks
__global__ void add(int *a, int *b, int *c){ c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x]; }
使用 threadIdx.x 代替 blockIdx.x，main 函数调用该 kernel 时，也需要修改调用参数
add<<<1, N>>>();
5. 线程块和线程
- 如果每个 block 有 M 个线程，则 int index = threadIdx.x + blockIdx.x * M;
__global__ void add(int *a, int *b, int *c){
int index = threadIdx.x + blockIdx.x * blockDim.x;
c[index] = a[index] + b[index];
}
#define N (2048*2048)
#define THREAD_PRE_BLOCK 512
...
add<<<N/THREAD_PRE_BLOCK, THREAD_PRE_BLOCK>>>(d_a, d_b, d_c);
...
6. 处理任意大小的向量
__global__ void add(int *a, int *b, int *c, int n){
int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < n){
c[index] = a[index] + b[index];
}
}
add<<<(N + M - 1) / M, M>>>(d_a, d_b, d_c, N);
相同块内的线程可以互相沟通和同步。但是不同块中的线程没有这个能力。
Q&A
- cudaMemcpy 通过高速总线实现数据传输。会有传输耗时，它的带宽比 GPU 内部带宽要小，会限制 GPU 的性能。因此尽量减少 host 和 device 侧的数据传输。
- host 和 device 侧的 size 可以确保是一致的，不会有 LP/LLP 差异，如 host 侧的 long 是 64bit，那么 device 侧的 long 也是 64bit。但在 host 侧使用的编译器类型受 cuda 限制。
- grid 的大小通常是根据问题规模来确定的
- threadIdx<1024，没有其他限制
- host 和 device 侧的代码可以放在两个文件中，只要能用 C/C++ 方式编译起来
