Cuda

基础知识

结构

在硬件上，一个 GPU 分为多个 SM，每个 SM 包含本地寄存器文件、统一数据缓存（L1 缓存和共享内存）和若干个执行计算的单元

执行一个 GPU 程序先会从 CPU 启动 Kernel （核函数），每一个函数称为一个 Grid，分为不同的 Block，每个 Block 执行多个线程 Thread，一个 Grid 使用的 Block 和每个 Block 中含有的 Thread 可以通过编程自定义 <<<Block,ThreadPerBlock>>>，Block 之间无法同步和共享内存（无法确定执行顺序），Cluster 将不同 Block 组织起来使其可以通信和访问共享内存

GPU 的执行单位为 Warp （线程束），它将线程分为每 32 为一组，每组执行完全相同的代码（无分支，对于不同分支屏蔽多余分支只执行一个分支）；将 Block 分到多个 SM 中进行计算（一个 SM 可以有多个 Block，但是一个 Block 只能分配到同一个 SM 中进行计算），同一 Cluster 的 Block被分配到同一 GPC 中

Tile and SIMT

SIMT 模型中，程序员按“线程”思考和编程，需要手动管理每个线程的数据访问和控制流，因此会出现 Warp Divergence 等问题。

Tile Programming 模型中，程序员按“数据块（tile）”思考和编程，整个 block 作为一个整体执行统一控制流，编译器自动把 tile 操作映射到线程，因此不再存在 Warp Divergence 的概念。

在 Tile 编程中，只需关注数据在数据块中如何运算，由编译器负责线程的组织。在这种编程方式下有两种数据类型，一个为数组，一个为 Tile （瓦片），瓦片不可变（对瓦片的操作只会产生新的瓦片），每个维度必须是 2 的幂并且编译期可知（不可动态）；数组和瓦片之间的数据可以进行移动，将数组分为多个小块，自带边界策略将超过边界的内容自动填充为 0

Tile 和 SIMT 可在同一程序中同时出现

GPU内存

为了将一个线程块调度到 SM 上，每个线程所需的寄存器总数乘以线程块中的线程数必须小于或等于 SM 中可用的寄存器数量

每个流式多处理器（SM）都包含一个 L1 缓存，该缓存是统一数据缓存的一部分。更大的 L2 缓存则由 GPU 内所有 SM 共享

CPU 内存只能由 CPU 代码访问，而 GPU 内存只能由运行在 GPU 上的内核访问。CUDA API 用于在 CPU 和 GPU 之间复制内存，被用来在正确的时间将数据显式复制到正确的内存位置

Cubin

一个 Cuda 二进制文件包含 CPU 机器码和一个 fatbin 容器，这个容器中包含针对不同版本的 SM 的 cubin 和 RTX，当该二进制程序运行时会从 fatbin 中选择最适合 GPU 的二进制文件 cubin，或者使用 RTX JIT

GPU编程（C++）

头文件

#include <cuda_runtime_api.h>

内核函数

__global__ void (function name)(argv){
    
}

内核函数的启动可以由三尖括号

__global__ void vecAdd(float* a,float*b ,float* c){
    // ...
}

int main(){
    vecAdd<<<1,256>>>(A,B,C);
}

内核函数启动后，默认情况下 CPU 不会等待 GPU 执行完毕

使用 cudaDeviceSynchronize() 让 CPU 等待 GPU 执行完毕

__syncthreads() 使块内线程等待直到所有线程都到达该点

线程块

可以通过变量设置块数和每块线程数

dim3 grid(16,16); // block:256
dim3 block(8,8); // threadPerBlock:64
vecAdd<<<grid,block>>>(A,B,C);

启动配置中未指定的维度将默认为 1

为了精确控制线程的执行，我们具有一些内建变量

  • threadIdx
  • blockDim
  • blockIdx
  • gridDim

这些都是三分量向量（dim3）

对于已知总线程数和每块线程数，要确定块数，可以使用

>#include <cuda/cmath>
>int thread = 256;
>int block - cuda::ceil_div(totalThread,thread);

内存分配

float* A = nullptr;
cudaMallocManaged(&A,size);
// ...
cudaDeviceSynchronize();
cudaFree(A);

为了提升性能

float* A = nullptr;
float* cudaA = nullptr;
cudaMallocHost(&A,size);
cudaMalloc(&cudaA,size);
cudaMemcpy(cudaA,A,size,cudaMemcpyDefault);
// ...
cudaDeviceSynchronize();
cudaMemcpy(A,cudaA,size,cudaMemcpyDefault);
cudaFree(cudaA);
cudaFreeHost(A);

cudaMemcpyDefault 会在以下类型中选择

  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
  • cudaMemcpyDeviceToDevice

错误处理

每个 CUDA API 都会返回一个枚举类型 cudaError_t 的值，在生产应用中，最佳实践是始终检查并管理

#define CUDA_CHECK(expr_to_check) do { \
cudaError_t result = expr_to_check; \
if(result != cudaSuccess) \
{ \
fprintf(stderr, \
"CUDA Runtime Error: %s:%i:%d = %s\n", \
__FILE__, \
__LINE__, \
result,\
cudaGetErrorString(result)); \
} \
} while(0)

cudaGetLastError 返回当前错误状态，然后将其重置为 cudaSuccess

cudaPeekAtLastError 返回错误状态而不重置它

实际上，使用三重尖括号符号启动内核不会返回 cudaError_t

最好手动检查

vecAdd<<<blocks, threads>>>(devA, devB, devC);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

cudaError_t 类型的值 cudaErrorNotReady（可能由 cudaStreamQuery 和 cudaEventQuery 返回）不被视为错误，也不会被 cudaPeekAtLastError 或 cudaGetLastError 报告。

CUDA_LOG_FILE=cudaLog.txt 的设置会让错误记录在日记中

$ env CUDA_LOG_FILE=cudaLog.txt ./test

Device and Host

对于函数

  • __global__ 说明内核函数入口点
  • __device__ 说明只能由内核函数调用
  • __hots__ 说明只能由 CPU 调用

可以同时使用 __device__ 和 __host__

对于变量

  • __device__
  • __constant__
  • __managed__
  • __shared__
  • 在 __device__ 或 __global__ 函数内部声明且未使用说明符则优先分配寄存器

Cluster

块级同步

cluster.sync()

可访问分布式共享内存

函数声明

__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{
}

编译宏

CUDA_ARCH

__host__ __device__
void foo()
{
#ifdef __CUDA_ARCH__
    // GPU code
#else
    // CPU code
#endif
}

GPU编程（python）

与常规 python 文件相比

需要添加包

from numba import cuda

在由 GPU 运行的函数前加上一段标识

@cuda.jit
def function(argv):
    # ...

调用函数时使用 [,] 代替C++中的 <<<,>>>

一样具有索引

  • cuda.threadIdx.[xyz]
  • cuda.blockDim.[xyz]
  • cuda.blockIdx.[xyz]
  • cuda.gridDim.[xyz]

可以仿照在C++常见的计算线程索引的方式

idx = cuda.threadIdx.x+cuda.blockIdx.x*cuda.blockDim.x

也可以使用一个简写语法

idx = cuda.grid(1)