Introduction to CUDA C++

__global__ 表示在设备上执行

threads 组成 block，blocks 组成 grid

使用 blockIdx.x 访问 block 下标，使用 threadIdx.x 访问 thread 下标

add<<<1, N>>> 表示启用 1 个 block，N 个 thread

启用 N 个 add() 副本:

add<<N/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>(d_a, d_b, d_c);

blockDim.x 表示每个 block 中的 thread 数量

访问每个 block 中的 thread：int index = threadIdx.x + blockIdx.x * blockDim.x

内存管理：cudaMalloc(), cudaMemcpy(), cudaFree()

CUDA Shared Memory

在一个 block 内，thread 通过 shared memory 共享数据
__shared__ 在每个 block 内分配 shared memory，由同一个 block 内的线程共享，对于其它 block 不可见

__syncthreads() 是屏障，在一个 block 内同步所有的 thread

Fundamental CUDA Optimization (Part 1)

一个 block 由一些 warp 组成，每个 warp 有 32 个线程

15.png

一个 warp 是物理并行于 multiprocessor 上的（SIMD）

通过切换线程可以降低延迟，一般

  • GMEM 延迟：> 100 cycles
  • Arithmemtic 延迟：< 100 cycles

为了消除 arithmemtic 延迟

  • 每个 SM（流多处理器）需要 10 个 warp（320 个线程）
  • 或者，如果同一个 warp 中的指令互相独立（输出不依赖前一个指令），只需要 5 个 warp

应最大化全局内存吞吐量，比如每个线程加载更多的数据

应使用大量的线程确保 GPU 忙碌。一般，每个 SM 使用 512+ 线程（目标是 2048，不过 512 -> 2048 时，收益显著减少）

Threadblock 配置

  • 每个 block 中的 thread 数量应该是 warp 大小（32）的倍数。因为指令总是以 warp 为单位发布的，如果线程数不是 32 的倍数的话，一个 warp 中会有些线程处于未激活状态。
  • SM 可以并行执行至少 16 个 block，一般每个 block 中配置 128-256 个线程（目标是 2048 个线程）

CUDA 中有一个占用率计算工具，应尽量让占用率达到 100%

占用率影响因素：线程负载，每个 block 的线程数，每个线程的寄存器数量，共享内存

Fundamental CUDA Optimization (Part 2)

GMEM 优化指南：

  • 内存合并

    • （内存对齐，内存填充）
    • 一个 warp 应访问连续的区域

  • 有足够的并发访问使得总线饱和

    • 每个线程处理多个元素

      • 多个 load 流水化
      • 索引计算通常可以重复使用

    • 启动足够多的线程以最大化吞吐量

      • 通过切换线程（warp）可隐藏延迟
  • 使用所有的缓存

shared memory 组织：32 个 bank，每个 bank 有 4 个字节

16.png

如果把 shared memory 看作一个二维数组，列对应 bank。数组的宽度为 32 bank（128 字节）

shared memory 性能：

  • 一般，在 1 或 2 个 cycles 中，从每个 bank 中获取 4 个字节，所以推荐按行式访问
  • 共享访问在 32 个线程（warp）发布

  • 串行化：如果 N 个线程的 32 次访问位于同一个 bank（完全列式访问），比如线程 1 访问字节 0, 1, 2, 3，线程 2 访问字节 128, 129, 130, 131，会被串行，性能最差，比最佳性能慢 32 倍

    • 避免列式的一个方式：填充一个新的列（逻辑列，对被访问数组填充 ），可以使得存储布局交错，避免所有线程都访问同一个 bank
  • 多播：N 个线程访问同一个位置，比如线程 1 和线程 2 都访问字节 0, 1, 2, 3，这并不是性能问题

Atomics, Reductions, and Warp Shuffle

  • atomicMax/Min
  • atomicAdd/Sub
  • atomicInc/Dec
  • atomicExch/CAS
  • atomicAnd/Or/Xor

Determine my place in an order

获得下一个位置：

int my_position = atomicAdd(order, 1);

返回该位置上的旧值，并将该位置更新为新值

Reserve space in a buffer

int my_dsize = var;
float local_buffer[my_dsize] = {…};

int my_offset = atomicAdd(buffer_idx, my_dsize);

// buffer_ptr+my_offset now points to the first reserved location, of length my_dsize
memcpy(buffer_ptr+my_offset, local_buffer, my_dsize*sizeof(float));

Reduce 例子
28.png

18.png
Warp shuffle Reduce 例子
允许 warp 内部实现直接的线程间通信，不需要借助共享内存等。

  • __shfl_sync(): 从指定线程（lane ID）复制
  • __shfl_xor_sync(): 通过运用 xor 操作来选择从哪个线程束（warp）复制
  • __shfl_up_sync(): 从向上相邻（更低偏移）的线程复制
  • __shfl_down_sync(): 从向下相邻（更高偏移）的线程复制

注意：在使用以上 API 时，源和目标线程都需要处于同一个 warp

19.png

  • mask 指定了哪些线程参与，0xFFFFFFFF 表示 warp 中的所有 32 个线程都参与
  • lane 指在 warp 中处于第几个线程
  • 部分和存储在该 warp 的第 0 号线程中，将其存入共享内存，索引为 warpID（最大为 32）
  • 最终将所有 warp 中的部分和进行规约

Warp shuffle 好处

  • 降低共享内存使用，避免共享内存成为占用率的限制因素
  • warp shulffle 是一条指令，而非多条指令
  • 避免显式同步
  • 仅用一条指令将单一值广播到整个 warp 中的所有线程，无需借用共享内存、循环或其台机制
  • 原子聚合：跨越一个 warp，意味着将 32 个原子操作放入队列中。由于 warp 的所有 32 个线程都在更新同一位置，如果先将这 32 个元素在 warp 内求和，然后仅由一个线程对内存中的该位置执行一次原子操作，而不是让 32 个线程各自独立进行原子操作，可减轻原子系统的压力，避免原子操作成为瓶颈

Managed Memory 托管内存

CUDA 8+ 引入了 unified memory，可避免使用宿主机指针和设备指针两个指针，只需要一个指针即可。unified memory 和 managed memory 是同一个概念

  • 用 cudaMallocManaged() 代替 malloc()
  • managed data 会在 kernel 启动后被大量移动，即使 kernel 没有显式用到
  • 在 kernel 启动后，cudaDeviceSynchronize() 使得数据对 CPU 代码再次可用
  • 使用 cudaFree() 代替 free() 或 delete

20.png

21.png

Unified memory 目标在于简化编程，而非性能。相反，它可能会让性能变差，这是逐页缺页中断导致的，开销高于批量移动数据。所以，对于大量数据移动，使用类似 cudaMemcpy() 的操作更好些。为了避免这点，引入如下函数

cudaMemPrefetchAsync(ptr, length, destDevice, stream)
__global__ void kernel(float *data){
    int idx = …;
    data[idx] = val;
}

…
int n = 256*256;
int ds = n*sizeof(float);
float *data;
cudaMallocManaged(&data, ds);
cudaMemPrefetchAsync(data, ds, 0);
Kernel<<<256,256>>>(data);
cudaMemPrefetchAsync(data, ds, cudaCpuDeviceId); // copy back to host

CUDA Concurrency

Pinned Memory 固定内存

一种特殊的内存分配器，确保数据一直存在于 CPU 物理 RAM 中，避免被“分页”出去

  • 更快的 Host <-> Device 拷贝
  • 异步于 CPU 和 GPU 的拷贝

用法：

  • cudaHostAlloc / cudaFreeHost 返回 / 释放固定内存
  • cudaHostRegister / cudaHostUnregister 对于已经通过 malloc 分配的内存，将其固定

CUDA Streams

流的语义：

  • 向同一个流发出的两个操作按顺序执行
  • 向两个不同流发出的操作之间无预设的顺序关系

创建并使用两个流例子：

22.png

  • cudaStreamCreate 相当于通知 CUDA，将使用这些命名的流作为流
  • cudaMemcpy(D2H, H2D) 会阻塞 cpu 线程，cudaMemcpyAsync(D2H, H2D) 不会，但要求主机指针指向固定内存，这就是引入固定内存的原因
  • cudaStreamQuery 会返回一个 bool，表明流是否空闲
  • cudaStreamSynchronize 阻塞 cpu 线程，等待此流中的所有 cuda 活动完成
  • cudaStreamDestroy 销毁流

使用流的例子：将整个串行步骤分解为多个并行步骤

23.png

  • H2D 复制的 绿色条长度大约为橙色条的两倍，橙色条与蓝色条的长度大致相同
  • 一般只需 2 到 3 个流即可，否则，可能没有以最佳方式使用流
  • 将一次性开销从性能循环中剔除出去，比如创建流，调用 cudaMalloc 等
  • 在复杂的并发场景中避免使用默认流，应使用命名流或创建的流。cuda 7 后，引入更改默认流的能力，将其转为普通的命令流

cudaLaunchHostFunc() 是流回调，取代 cudaStreamAddCallback()

托管内存下的拷贝-计算并行

使用托管内存时，无需使用 cudaMemcpy 或 cudaMemcpyAsync，使用 cudaMemPrefetchAsync，相当于 cudaMemcpyAsync，实现数据迁移与内核调并行。

细节上，cudaMemPrefetchAsync 比 cudaMemcpyAsync 执行更多的工作，涉及了页面的换入换出，除了数据迁移外，源和目标的页表均需更新。所以，迁移调用可能比普通的异步复制操作耗时更长。

不过这种开销的影响可能会根据工作流的繁忙成多有所不同，反直觉的是，繁忙的流可能有更低的开销，称之为页面表修复的开销

CUDA Event

cudaEvent 是流中的一个标记，本身不会导致任何事情发生。

cudaEvent 基本用来计时：

24.png

  • cudaEventCreate 注册 cudaEvent 名称
  • cudaEvemtRecord 记录 cudaEvent，将其发布到流中
  • cudaEventSynchronize 会阻塞 cpu 线程，直到流的处理达到记录 stop event 的点，即核函数执行完成
  • cudaEventElapsedTime 计算两个 event 之间的时间流逝，以 ms 为单位存放到 float_var 中

Multi GPU

设备管理

有的系统中可能在单个节点内集成多个 gpu，应用可以查询、选择 gpu：

  • cudaGetDeviceCount(int *count) 获取 gpu 数量
  • cudaSetDevice(int device) 默认使用 0 号 gpu，如果使用其它 gpu，需明确设备编号
  • cudaGetDevice(int *device) 获取获取当前在使用哪个 gpu
  • cudaGetDeviceProperties(cudaDeviceProp *prop, int device) 获取设备属性

多个 host 线程可共享一个 gpu，比如每个 host 线程都可以设置 cudaSetDevice(0)

一个 host 线程可以管理多个 gpu

  • cudaMemcpyPeerAsync(...) 支持将数据直接从一个 gpu 拷贝到另一个 gpu

流

  • 流和 cudaEvent 都会自动关联到设备。这意外着，当创建一个流时，该线程中最近通过 cudaSetDevice 设置的设备，将成为该流关联的设备
  • cudaStreamWaitEvent 可以等待另外一个流发出的 event，不过尽量少调用，会让事情变的复杂
  • cudaEventQuery 可以检测另外一个流发出的 event 是否完成

在两个 gpu 上同时运行代码：

25.png

GPU 之间的数据拷贝

cudaMemcpyPeerAsync 可以避免经过主机内存，直接实现 gpu 之间的数据拷贝。但这要求系统拓扑结构支持此特性，即两块 gpu 之间可以建立对等关系，位于同一总线上，无论是 PCI Express 还是 NVLink.

26.png

  • cudaDeviceCanAccessPeer 测试两块 gpu 是否能建立对等关系
  • cudaDeviceEnablePeerAccess 将一个 gpu 设置为另一个的 peer。上图中启用了两个方向的对等访问，先设置了 0 到 1 的访问，然后设置了 1 到 0 的访问，以实现双向传输
  • 一个设备可访问的对等设备数量是有限的

流优先级

目前只有 2 个优先级（但没有明确说明）。

例子：

27.png

GPU Performance Analysis

编写高性能代码原则：

  • 高效利用内存子系统

    • 全局内存合并访问，即希望同一个 wrap 中的相邻线程能够读取内存中的相邻元素
    • 利用可用的缓存资源，共享内存

  • 充分展现并行性或工作量，充分利用机器的性能并隐藏延迟

    • 尽量使用大量线程
    • 占用率
    • 调整每个线程的工作量

分析驱动的优化过程是一个循环往复的过程：

  • 运用分析器确定性能瓶颈
  • 分析哪些代码与该瓶颈有关
  • 优化代码

性能瓶颈：

  • 内存受限，说明正在大量使用内存，正在使内存系统达到饱和
  • 计算受限，说明正在大量使用计算资源执行算术任务
  • 延迟受限，gpu 上的空闲周期数增大，说明未能妥善隐藏延迟，严重影响了机器性能
  • 通常一次分析一小段代码，比如一次只分析一个内核

确定性能瓶颈的指标（收集一些指标，排查哪个比较高）：

  • 延迟指标

    • SM（流多处理器）效率指标（低说明受限于延迟）：smsp__cycles_active.avg.pct_of_peak_sustained_elapsed

  • 内存指标

    • DRAM 利用率（可能受限于 DRAM 带宽）： dram__throughput.avg.pct_of_peak_sustained_elapsed
    • L2 利用率：lts__t_sectors.avg.pct_of_peak_sustained_elapsed
    • 共享内存利用率：l1tex__data_pipe_lsu_wavefronts_mem_shared.avg.pct_of_peak_sustained_elapsed

  • 计算指标

    • DP 利用率：smsp__inst_executed_pipe_fp64.avg.pct_of_peak_sustained_active
    • SP（单精度浮点数）利用率： smsp__pipe_fma_cycles_active.avg.pct_of_peak_sustained_active
    • HP 利用率：smsp__inst_executed_pipe_fp16.avg.pct_of_peak_sustained_active
    • TC 利用率：sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_active
    • 整数利用率：smsp__sass_thread_inst_executed_op_integer_pred_on.avg.pct_of_peak_sustained_active

内存受限：

  • 当代码受限于内存带宽或延迟时，可视为内存受限。措施是为 gpu 提供更多的并行工作，以提高占用率，使用大量线程解决延迟隐藏不足的问题
  • 确定内存带宽问题所在的位置，是全局内存还是共享内存。在数据重用场景，利用多级内存，比如 L2 缓存，共享内存
  • 计算出数据总量，除以内存带宽，可得到理想情况下加载和存储所有数据一次所需的时间长度，该值为内核性能上限，从而得到内核实际吞吐量和该内核峰值理论吞吐量之间的差距。当测量性能接近理论峰值时，说明性能优化完成了

计算受限：

  • 使用单精度浮点运算计算理论峰值性能，比较与实际性能之间的差距
  • 也可以用利用率，如果利用率高，说明代码正在以接近机器处理器极限的速度发出浮点运算指令
  • gpu 可能有独立的单精度浮点运算单元、双精度浮点运算单元，用于整数处理的独立单元，如果代码受限与整数运算，可以将这部分计算转换为浮点运算

延迟受限：

  • 说明现有的并行工作无法使 gpu 忙碌

  • 一般的策略是增加并行工作

    • 启动大量的线程
    • 提高每个线程的工作量，比如增加循环
    • 使用向量加载来使得一个线程处理多个输入
    • 关注占用率，即为 SM 加载了多少资源

占用率：

  • 指单个 SM 上驻留了多少线程，与机器饱和状态密切相关
  • cuda 提供了占用率计算表格工具
  • 受限于每个线程使用的寄存器数量，每个 block 中的线程数量，共享内存的使用

参考

cuda-training-series

cuda-c-programming-guide