博客园 首页 私信博主 显示目录 隐藏目录 管理 动画

cuda 笔记

cuda

Tags: 笔记

目录

基础

Grid, Block, Thread

thread 对应 SP (cuda core)，是基本的运算单元/执行单元。
thread block 对应 SM，包含许多运算单元，是能完整独立运行的最小单元。
grid 对应 GPU (device) 上执行的一个函数，包含可分配到不同 SM 上的多个 block 执行。

为了便于编程和管理线程，cuda 引入了 网格 (grid)、线程块 (thread block)、线程 (thead)、线程束 (warp) 四个概念。
执行 kernel 时，其使用的所有 thread 组成了一个 grid，会使用 GPU 上的部分计算单元。每个 grid 包含若干可并行执行的 block，每个 block 内包含若干 thread（最大为 1024，32 个 warp）。

同一 grid 中的所有线程共享全局内存空间。

GPU 可以创建的 grid、block 数量取决于其计算能力。

grid 中的 block 可以以一维、二维、三维三种组织方式排列；block 中的 thread 也是如此。
block 的所有线程实际上是一维的（分成若干 warp）；grid 的所有 block 之间也没有空间关系。分三个维度的好处是可以减少一些计算索引时的除法或取余运算，毕竟整数除法和取余的开销较大。

thread

GPU 通过切换到其它线程执行，来隐藏有依赖的指令的延迟。因此要想充分利用 GPU，线程数要远高于核数或指令流水线数。

block

执行同一函数的多个线程会被划分为多个线程块；每次执行一个函数会启动多个线程块。

同一个 block 中的 thread 只会在同一个 SM 中并行执行，它们可以通过同步或共享内存通信进行协作（不同块之间的 thread 不能直接通信，只能通过全局内存和 cooperative group），所以 block 也被称为 Cooperative Thread Array (CTA)。
在执行 kernel、启动 grid 时，该 grid 的 block 就会被分配到可用的 SM 上执行。
一个 block 只会由一个 SM 调度，且一旦被分配好 SM，就会一直在其中执行。不同 block 可由可用的不同 SM 执行。

block 有一些 block 层的资源：

  • Shared Memory (共享内存)：block 层最显著的资源。它的容量可以由编译器决定，也可以运行时动态指定，但是一般有上限（参考 CUDA Programming Guide: Features and Technical Specifications 的 Table 15）。
    Shared Memory 对于整个 block 内的每个 thread 都是可读可写的，block 外则无法访问。
    较新的架构里 Shared Memory 还支持 atomic 操作，效率比用 global memory 做 atomic 操作有很大提升。
  • Synchronize Barrier：block 有一个重要操作就是同步（比如__syncthreads()，参考 Synchronization Functions，或 PTX 中 bar 和 barrier 指令）。
    每个 block 有16个 barrier，每个可以支持独立的同步操作（规定到达线程数，是 arrive 还是 sync 模式等等，具体见 PTX 的 bar 指令。
    注意同步不仅仅是保证 warp 都运行到某个指令的位置，还要求之前的一些操作如 memory load 的 dependency 完成。
  • 其它资源：特殊寄存器（blockIdx.x/y/z 等）；内部状态变量，比如 GPR 在整个 register file 里的起始地址、local memory 的起始地址；debug 用的资源等。这些对用户不完全开放，一般不用太过关注。

Block 是 kernel 运行时进行资源分配的最小完整单元：block 的所有资源限制在同一个 SM 内，启动前必须全部分配完成：block运行时的所有资源（包括 GPR、每个 thread 的私有资源、warp 私有资源，block 一级的 shared memory 等），都必须在 block 启动前就绪。如果一个 SM 无法提供足够资源，则 block 无法在这个 SM 上启动。
如果部分 warp 提前退出，它的资源应该可以先被释放，但 shared memory 只能在当前 block 的所有线程完成后才释放。

block 保证了 warp 和 thread 在运行前能分配到所有需要的资源，且这些资源在运行时随时可用，使得 warp 可以高效切换。

每个 SM 需要同时运行足够多的 warp 才能有效隐藏带有依赖的指令的延迟，而 SM 上的 block 数量、block 中的 warp 数量都有限。因此，调整 block 的资源占用（每个线程的 GPR 数量、warp 数、shared memory 大小等），可以影响同一个 SM 上能容纳的 block 数目，从而调整 SM occupancy。

Block 之间基本相互独立，没有数据交换，理论上能以任意顺序发送到任何空闲的 SM 上。

如果 CUDA 支持 block 部分 warp 先退出则资源可以先回收，那一个 SM 可能容纳小数个 block，后进的完整 block 能与之前 block 残留的 warp 同时运行。

grid

用户可见的 grid 资源，比如 global memory、constant memory、texture/surface reference 或 object 等，都是可以在同一个 context 下的所有 kernel 间继承和共享的。每个 kernel 在 grid 这一层并没有太多的私有资源。

cudaThreadSynchronize

cudaThreadSynchronize()可同步同一个 block 的所有 thread（均在此阻塞直到全部到达）。

内置变量：

  • blockIdx：线程块在对应 grid 中的编号。
    threadIdx：线程在对应线程块中的编号。
    两者都是uint3向量类型，包含 x, y, z 三个 uint 成员。
  • gridDim：grid 的维度，即每一维的 block 数量。对应调用 kernel 时的 gridSize。
    blockDim：block 的维度，即每一维的 thread 数量。对应调用 kernel 时的 blockSize。
    两者都是dim3向量类型，包含 x, y, z 三个 uint 成员。

例：

  • gridDim.x：当前 grid 中线程块的数量（一维）/第一维度的大小（多维）。
  • blockDim.x：线程块中的线程数量（一维）/第一维度的大小（多维）。

计算：

  • 线程块在整个 grid 的编号：
    一维：blockIdx.x。
    多维：blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z。
  • 线程在整个 grid 中的编号：
    一维：blockIdx.x * blockDim.x + threadIdx.x。
    多维：blockId * (blockDim.x * blockDim.y * blockDim.z) + (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x) + threadIdx.x。
  • grid 中的线程总数量：
    一维：blockDim.x * gridDim.x。
    三维：grid_size.x * grid_size.y * grid_size.z * block_size.x * block_size.y * block_size.z。
    常用作线程枚举的步长 (stride)（这种循环称为 grid-stride loop）。

SM

由于一个 block 只会在一个 SM 上执行，且 SM 可同时执行多个 block，因此想要充分利用 GPU，需要启动许多（数倍于 SM 数量）block。

由于 Wave Quantization，block 数量最好是 SM 的整数倍（或接近但小于），来降低最后一个 wave (tail wave) 的浪费。

Warp

线程束 (warp) 是 cuda GPU 调度和执行的基本单元（是 SM 执行的基本单元）（是 cuda 软件的概念，但是与硬件直接相关）。
目前 cuda 一个 warp 的大小 (warpSize) 为32个线程（即一个 SM 上只有32个线程能同时执行）。在同一个 warp 中的线程能以任意顺序执行，并会以不同的数据资源执行相同的指令 (SIMT)。

很多内存指令的访问模式都是以 warp 为单位。

thread 本意是线，warp 是经纱（线的集合）。

通过把 warp 中的 thread 同步运行，分摊指令 fetch、decode 等各种开销，可以降低运行调度的复杂度、降低功耗。

当一个 block 被调度到某个 SM 时，block 中的 thread 会被划分为多个 warp 分别被调度执行，每个 warp 中的 thread 会执行相同指令，但拥有独立的 PC、寄存器和数据。所以 block 的大小（三维总大小）最好是 warpSize (32) 的整数倍，以避免最后一个 warp 中包含无用线程。
在软件 (编程) 角度，block 是线程的集合（可以是一维/二维/三维）；在硬件角度，block 是 warp 的集合（一维）。

一个 warp 需要占用一个 SM 运行，多个 warp 需要轮流进入 SM 执行，由 SM 的硬件 warp scheduler 负责调度。
所以任意时刻 GPU 上的活跃线程最多只有 SM 数量 * 32 个，不是所有线程都在物理上同时执行（只是逻辑上并行）。

当一个 warp 空闲时，SM 就可以调度驻留在该 SM 中的另一个可用 warp。
在并发的 warp 之间切换没有什么代价，因为硬件资源早就被分配到所有的 thread 和 block，新调度的 warp 的状态已经存储在 SM 中了。这不同于 CPU：CPU 切换线程需要保存/恢复上下文；而 GPU 为每个 thread 提供物理寄存器、为每个 block 分配资源，无需保存/恢复上下文。

warp 分支 / warp divergence

一个 warp 中的线程会执行相同指令。如果线程执行的代码中出现分支，则需要分别顺序执行每个分支路径 (warp divergence)：对于单个分支，如果存在某个线程求值为 true，则并行执行为 true 的线程，其它线程等待 (stall)；执行完 true 分支后，所有为 false 的线程再执行，其它线程等待。
因此要提高并行性，应尽量避免同一 warp 中的线程进入不同分支。

warp shuffle

Warp 内的各个线程交换数据可以用 warp shuffle，是直接基于寄存器的数据交换，并不需要额外的存储空间。模式可以一个 lane 广播到所有的 lane，也可以有比较复杂的交换 pattern。warp shuffle 与基于 shared memory 的数据交换各有优劣。

SIMT

CUDA 是单指令多线程 (SIMT) 架构，与 SIMD 类似，更灵活、但效率略低。
二者都通过将同样的指令广播给多个执行单元来实现并行。主要的不同点是：SIMD 要求所有的 vector element 在一个统一的同步组里同步执行，而 SIMT 允许线程在一个 warp 中独立执行；SIMT 中的每个 thread 拥有自己的 instruction address counter (PC)、状态寄存器和独立的执行路径（可选）。

kernel 按 SIMT 执行，即线程会用不同数据执行相同指令。

kernel

__global__修饰的函数为核函数。

kernel 的返回类型只能是 void，如果需要返回值则只能传参指针。
kernel 不能是成员函数。

__device__修饰的函数为设备函数，只能被核函数或其它 device 函数调用，只能在 device 中执行。

类似的__host__修饰的函数为主机函数，只能由 host 调用和执行，一般不用。

执行 kernel

调用：kernel_name<<<grid_size, block_size>>>(args)。

<<<grid_size, block_size>>>称为执行配置 (execution configuration)，告诉 CUDA runtime 在该 grid 中使用多少个 block 和 thread 及组织形式。
grid_size 与 block_size 可以是 dim3 类型的结构体（二维或三维）或一个 unsigned int（一维）。

dim3 grid_size(2, 3);
dim3 grid_size(2, 2, 3);

执行/启动一个 kernel 称为 launch。
执行 kernel 默认不会阻塞 CPU，可使用cudaDeviceSynchronize()阻塞直到 device (GPU) 完成。

从异步调用 kernel launch API 到 kernel 真正执行的这段时间，称为 launch latency (或 induction time)。
launch 时间可能比执行 kernel 本身还要长（约几~几十微秒）。所以会通过算子融合、cuda graph、多个 cuda stream、batching 等方式减少 launch 开销。

CPU 调用 kernel launch API 通常非常快，而 GPU 执行相对更慢。所以如果 CPU 连续异步调用数个 kernel，会很快就调用完成进入空闲，此时需安排其它 CPU 任务执行；GPU 需要依次执行每个 kernel，所以每个 kernel 的 launch latency 都会越来越大（第一个较短，第二个要加上第一个 kernel 的执行时间，依次类推）。

从任务被加入对列到任务执行完成的时间，称为 task latency (或 total time)。

编写 kernel

kernel 主要有两部分：确定数据与线程的对应；处理对应数据。

Dynamic Parallelism

CUDA 5.0 引入了 Dynamic Parallelism。此前 kernel 都只能从 host 端启动，grid 和 block 的 dim 必须在 host 端确定好。这对于形状规则易于均匀划分的计算任务是合适的。但有些应用一些区域任务多、一些区域任务少，有时任务大小需要经过复杂计算，并不能一开始就得到。Dynamic Parallelism 允许 kernel 函数内再启动 kernel，由父 kernel 负责计算子 kernel 所需的 grid、block 的 dimension，也包括分配子 kernel 所需要的一些内存资源等（有一套 device 端的 API）。

内存管理

kernel 中访问的内存必须是位于 GPU 的内存，host 代码中访问的内存必须是位于 CPU 的内存。

cudaError_t cudaMalloc(void**, uint)分配位于 GPU 上的内存。
需传入一个二级指针d_ptr，分配时会修改它的值。

cudaError_t cudaMemcpy(void*, void*, uint, enum cudaMemcpyKind)在 CPU 内存和 GPU 内存之间进行拷贝。可用来初始化或将结果移动回 CPU。
指针参数分别为d_ptr和h_ptr，前者为 device 上的指针，后者为 host 上的指针。
enum 代表数据传输的方向，包括5种：

  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
  • cudaMemcpyHostToHost
  • cudaMemcpyDeviceToDevice
  • cudaMemcpyDefault

例：

// 在 CPU 上初始化
double *h_x = (double*) malloc(M);
for (int n = 0; n < N; ++n) {
    h_x[n] = a;
}
// 分配 GPU 内存、拷贝
double *d_x;
cudaMalloc((void **)&d_x, sizeof(double) * N);
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);

memcpy

each memcpy function is categorized as synchronous or asynchronous, corresponding to the definitions below.

Synchronous:

  • For transfers from pageable host memory to device memory, a stream sync is performed before the copy is initiated. The function will return once the pageable buffer has been copied to the staging memory for DMA transfer to device memory, but the DMA to final destination may not have completed.
  • For transfers from pinned host memory to device memory, the function is synchronous with respect to the host.
  • For transfers from device to either pageable or pinned host memory, the function returns only once the copy has completed.
  • For transfers from device memory to device memory, no host-side synchronization is performed.
  • For transfers from any host memory to any host memory, the function is fully synchronous with respect to the host.

Asynchronous:

  • For transfers between device memory and pageable host memory, the function might be synchronous with respect to host.
  • For transfers from any host memory to any host memory, the function is fully synchronous with respect to the host.
  • If pageable memory must first be staged to pinned memory, the driver may synchronize with the stream and stage the copy into pinned memory.
  • For all other transfers, the function should be fully asynchronous.

Unified Memory

Unified Memory 是 CPU 和 GPU 都可以访问的内存（也称为 managed memory）。（并且可以在任意 GPU 上访问？）

使用cudaMallocManaged(void**, uint)分配统一内存，使用cudaFree(void*)释放统一内存。
与 malloc、free 类似，而不是 new、delete。

// Allocate Unified Memory (accessible from CPU or GPU)
float *x;
cudaMallocManaged(&x, N * sizeof(float));
cudaFree(x);

当访问 unified memory 中的数据时（称为 CUDA managed data），cuda 软件和/或硬件会将内存页迁移到正在访问的处理器 (CPU/GPU) 内存中 (migration)（不是真的共享，只是把数据迁移过程隐藏了）。
内存迁移时间可在 Unified Memory profiling result 中看到。

在 Pascal 架构之前的 GPU（如 K80），调用 cudaMallocManaged 会立刻在 GPU 上分配指定数量的内存，并为分配涉及的页创建页表条目。当在 CPU 上访问这些完全位于 GPU 的内存页时，会触发 page fault，然后 GPU 驱动会将访问的页从 device 迁移到 CPU 内存 (device to host)。
在 launch kernel 时，由于这些旧的 GPU 没有 page fault 机制，因此需要将先前迁移到 CPU 内存或其它 GPU 中的所有页迁移回来 (host to device)，不管实际是否会用到。因此每次启动 kernel 都可能有不必要的开销。
因此，迁移页在 kernel 运行前完成，迁移时间不会被计入运行时间。

自 Pascal 架构起，调用 cudaMallocManaged 可能不会立刻分配 managed 内存和创建页表条目，而是在访问或预取时分配；GPU 支持49位虚拟地址和按需的内存迁移。
内存页同样需要在 CPU 与 GPU 之间迁移，但是在硬件上支持 page fault 和内存迁移，因此启动 kernel 前不需要迁移页、没有不必要的迁移开销；当访问的页不存在时，GPU 会阻塞对应线程的执行，由 Page Migration Engine 将对应的页迁移到 device。
因此，迁移页在 kernel 运行时进行，迁移时间会被计入运行时间。这可能导致新 GPU 的运行时间反而长于旧 GPU。

Pascal 起，程序可以用cudaMemAdvise()引导驱动迁移内存，用cudaMemPrefetchAsync()显式迁移内存。（见这里）

内存页迁移会影响运行时间。可通过以下方式减少页迁移对运行时间测试的影响：

  • 将数据初始化从 CPU 移动到另一个 kernel 中（即 GPU 中），使数据最初就在 GPU 上。
  • 通过 cudaMemPrefetchAsync 在执行 kernel 前将数据预取到 GPU 中。
  • 运行 kernel 多次，取最小运行时间。

并发访问

自 Pascal 架构起，CPU 和 GPU 可同时访问 managed memory，因为都支持 page fault。但程序需要保证没有 data race（比如通过 cudaDeviceSynchronize 等待后再读取）。

Pascal 和 Volta GPU 支持全局的原子内存操作：可在多个 GPU 上原子地读写值。

在 Pascal 架构之前的 GPU，如果 GPU 的 compute capability 低于6.0，则无法同时在 CPU 和 GPU 上访问 managed memory，因为硬件不支持 page fault、难以保证一致性，如果同时访问会导致 segment fault。

Pinned Memory

Pinned memory（也称为 page-locked memory, pMem, non-pageable, non-swappable）是一种不能被操作系统分页 (swapped out, paged) 的内存。这种内存保存在物理内存中，它的一个显著特性就是不会被操作系统的虚拟内存管理机制移动到磁盘上的交换空间 (swap space)。在某些高性能计算和数据传输场景中很有用，比如 GPU 编程（CUDA 提供了对 pinned memory 的支持）。
因为 pinned memory 不会被分页出内存，因此其物理内存地址是固定的。这使得在进行 DMA（直接内存访问）时，数据传输更加高效，因为硬件可以直接从已知地址进行操作。

cuda 不允许异步传输 paged 内存中的数据？只有保证数据始终在内存中才可异步传输。
使用 pageable 内存与 GPU 传递数据可能很影响性能（可以在 nsys - cuda HW - show in events view 中确定）。

优点：

  • 从 CPU 到 GPU: 使用 pMem 可以显著提高数据主机（CPU）和设备（GPU）之间的传输速度。通常情况下，数据需要被拷贝到一个临时的 pMem，然后再传输到 GPU。但如果一开始数据就在 pMem 中，这个过程就可以省去，从而加速传输。

缺点（不能滥用）：

  • 内存容量限制: 由于 pMem 不能被分页，其数量受到物理内存的限制，是稀缺的 os 资源。使用过多的 pMem 会减少可供操作系统和其他应用程序使用的内存、影响系统性能。
  • 内存分配时间: 分配和释放 pMem 的时间比普通的分页内存更长，也取决于 os 可用资源。

使用 cudaMallocHost 分配 pMem（malloc 分配的是 non-pinned memory。cudaHostAlloc 与 cudaFreeHost 是什么？）。
pMem 可以实现真正的异步行为？比如拷贝和计算可异步、同时进行？

如果在分配时 pMem 时指定 cudaHostAllocMapped，则分配的内存会被映射到 GPU 地址空间？此时 GPU 可以直接访问 pMem。
在某些条件下（只要支持 unified addressing？）不指定该 flag（使用 cudaHostAllocDefault）也会进行映射 (Automatic Mapping)？

All host memory allocated through all devices using cudaMallocHost() and cudaHostAlloc() is always directly accessible from all devices that support unified addressing. This is the case regardless of whether or not the flags cudaHostAllocPortable and cudaHostAllocMapped are specified.

因此 pMem 一般可以同时在 device 和 host 上访问而不需拷贝，所以也称为 zero-copy。
但在 device 上的访问速度不如 GPU 本地内存，所以通常还是会拷贝到 GPU。

Stream

stream 是独立执行的任务队列，stream 内的任务可能有依赖关系，但 stream 之间的一定没有。
不同 stream 可在某些条件下并行执行，以提高 GPU 利用率和 warp occupancy：某个 stream 在做内存传输 (h2d/d2h) 时，另一个 stream 可执行 kernel 计算。
将独立的任务划分到独立的 stream 是有必要的。

stream 在 kernel configuration 的第三个位置。
当不指定 stream 时，效果取决于 nvcc 选项 --default-stream 或宏CUDA_API_PER_THREAD_DEFAULT_STREAM。
legacy default stream 会与相同 CUcontext 中的所有 blocking stream（创建时未指定 cudaStreamNonBlocking）同步。
per-thread default stream 不与任何线程同步（除了 legacy default stream），与显式创建的 stream 类似。

为 kernel 指定 stream 一般没有什么影响，除了会在指定 stream 中执行。
例外：当调用device_array.copy_to_host()时（不带参数），拷贝始终是同步的；当传递一个 stream 参数device_array.copy_to_host(stream)时，如果 device_array 不是 pinned 则同步，否则异步拷贝。即如果传递 stream 且是 pinned，则会异步拷贝。

GPU

https://docs.nvidia.com/deeplearning/performance/dl-performance-gpu-background/index.html

各 GPU 计算能力 (compute capability)：https://developer.nvidia.com/cuda-gpus#compute

各架构整体：https://zhuanlan.zhihu.com/p/394352476

一般说的 GPU 计算并不是纯粹由 GPU 完成的，而是由 CPU 调度、GPU 计算共同完成的异构计算任务。
其中：CPU + 内存 被称为 Host (主机)，GPU + 显存 被称为 Device (设备)。

流处理器 (SP, streaming processor)：最基本的运算单元 (cuda core)。指令和任务都是在 SP 上执行。

流式多处理器 (SM, streaming multiprocessors)：多个 SP 和其它资源组成一个 SM（类似多核 CPU 的一个 CPU 核）。
资源包括：warp scheduler、register file (寄存器组)、shared memory/L1 cache、dispatch unit、load/store unit 等。这些资源有限，由 SM 上的所有 thread 共享。
GPU 就是 SM 阵列，包含多个 SM 和许多 SP，由架构和型号决定。

GPU 是数据并行，由每个 thread 执行不同数据的计算（任务是相同的）。
每个 thread 都有自己的物理寄存器组，因此切换 warp 时无需恢复上下文。

A10 有 72 个 SM，T4：40，A100：108。

tensor core

tensor core 是 Volta 及之后的架构中才有的。相比于 CUDA core，可以提供更高效的运算。

索引地址寄存器

每个线程有自己在 block 内的索引地址寄存器SR_TID.X, SR_TID.Y, SR_TID.Z（对应 threadIdx 的 x、y、z 分量）。
如果需要获得这些值，可以用S2R R0, SR_TID.X这种指令。

Predicate

Predicate 是1位的 bool 谓词，每线程有8个, SASS里用 P0~P7 表示，其中 P7=PT 始终为 true，@PT为默认谓词。
Predicate 通常是一些 bool 运算（比如大于、不等于）的输出，用作指令的谓词或操作数。
CUDA 的每个汇编指令都可以用 Predicate 来控制是否真正运行，形式为@!P0 BRA 0xc40。这里@!P0表示 P0 为假时这行才生效，!代表取反。

在指令中，Predicate 由4位编码：3位为索引 (0~7)，1位表示是否取反。

Predicate 与直接 branch 跳转的优点为可以避免 warp divergence，而且开销相比 branch 指令更小（branch 的延迟一般较高，还涉及到指令 Cache 和内部 pipeline 连续性的问题）。
但是，即使 Predicate 为否，这个指令的运行开销也不会被省略，只是指令不写回结果。但由于同一 warp 的线程同时只能发射相同的指令，所以这样也没有影响，反而少了跳转的开销。

bank conflict

此外还有 register bank conflict。

为了提高带宽，共享内存在物理上被分为 32 个（等于 wrap size）同样宽度、可同时访问的 bank，每个 bank 中的内存地址可以从 0 开始编号。
除了 kepler 架构中每个 bank 的宽度是 8B，其它架构中每个 bank 的宽度都是 4B。
对于 bank 为 4B 宽度的架构，内存上连续的 128B 由 32 个 bank 上的同一位置的 4B 组成，可以并行访问（类似 CPU 的内存，内存带宽的 64B 由 8 个 8B 的 bank 组成）。

当同一个 wrap 中的不同线程并行访问一个 bank 中的不同地址时，这些访问必须被串行处理，即会发生 bank conflict，大大降低内存带宽。
当同一个 wrap 中的不同线程并行访问一个 bank 中的相同地址时，会发生广播、将数据一次传给所有线程，不会发生冲突。

在一个 wrap 内对同一 bank 的 n 个地址同时访问将导致 n 次内存事务，称为发生了 n 路 bank conflict。

不同 wrap 的线程之间不存在 bank 冲突。

通常可以通过改变共享内存数组大小来减少 bank conflict。

memory hierarchy

整体上，每个 SM 有自己的 L1 cache，多个 SM 共享 on-chip L2 cache，然后 L2 与 DRAM (GPU mem) 相连。

内存

具体划分可见下 Turing 内存架构。

按照存储功能，GPU 内存可以分为：全局内存 (global memory)、局部内存 (local memory)、常量内存 (constant memory)、共享内存 (shared memory)、寄存器 (register)、L1/L2 缓存等。
其中全局内存、局部内存、常量内存都是片下内存 (off-chip)，访问速度一样（较慢），储存在 HBM 上，只是访问方式与可见性不同。

  • 全局内存：能被 GPU 的所有线程访问，全局共享。跟 CPU 架构一样，运算单元不能直接使用全局内存的数据，需要经过 cache。HBM 的大部分都用作全局内存。
  • 局部内存：每个线程独享的内存资源，线程之间不可以相互访问。
    主要是用来应对寄存器不足时的场景，即在线程申请的变量超过可用的寄存器大小时 (register spill)，nvcc 会将一部分数据放到片下内存里，或在内存中暂存寄存器的值。
    它不是线程自己动态申请的资源，而是整个 kernel 启动时为每个线程分配好的固定大小的资源。大小由编译器决定，运行过程中不能改变。
  • 常量内存 (constant memory)：只读内存。会通过特殊的 L1.5 常量缓存 (constant cache) 进行缓存读取，比全局内存快，但依然比片上内存慢。

常量内存

保存的常量包括：__constant__ 数据，kernel 调用参数，立即数。
主要解决一个 warp scheduler 内多个线程访问相同数据时速度太慢的问题。假设所有线程都需要访问一个 const 的常量，在存储介质上 const 的数据只保存了一份，而内存的物理读取方式决定了多个线程不能在同一时刻读取到该变量，只能串行读取影响并行效率。
常量内存支持硬件 broadcasting：当一个 wrap 内的多个线程访问相同地址时，常量内存可将数据同时发送给所有线程；当多个线程访问不同地址时，访问依然会被串行化。

其它内存是 GPU 的片上内存 (on-chip)，都是 SRAM，比片下内存更快。

  • L1 cache：每个 SM 有自己的 L1 cache，由该 SM 的 cuda cores 共享。
  • L1.5 constant cache：constant cache 被分为 L1, L1.5, L2 三部分。
    每个 SM 有自己的 L1.5 cache，速度介于 L1 与 L2 之间，用于缓存常量内存中的数据。
    在 Pascal 之前还有 L1.5 instruction cache。
  • L2 cache：所有 SM（整个 GPU）共享 L2 cache，用于缓存 GPU 上的 DRAM（显存）。
    L2 data cache、L2 instruction cache、L2 constant cache 通常被合并为一个更大的 L2 cache。
  • 共享内存 (shared memory)：在线程块内能共同访问的内存。是一个小容量的 cache，主要缓存一些需要反复读写的数据。
    共享内存的位置与速度与 L1 接近，区别在于共享内存的控制与生命周期管理与 L1 不同：共享内存受用户控制，L1 受系统控制。共享内存更利于线程块之间数据交互。
  • 寄存器：线程能独立访问的资源，用来存储线程的暂存数据。访问最快，但容量较小，而且要被许多线程均分？

此外，在 Volta、Turing 架构中：

  • L0 cache：每个 processing block 即每个 wrap 有自己的 L0 指令 cache。
  • L1 cache 和 shared memory 被合并为一个更大的 L1 data cache/shared memory（Maxwell 和 Pascal 中是分开的，但更早的 Kepler 是合并的）：This design reduces the cache hit latency and improves the L1 bandwidth。

Turing 内存架构

具体可见 https://arxiv.org/pdf/1903.07486 Chapter 3 Memory hierarchy；每部分的大小和访问周期见 Table 3.1。

以 T4 为例，每个 SM 被划分为 4 个 processing blocks，每个处理块有 1 个专门的 warp scheduler 和 dispatch unit。同一个 wrap 的指令会被分到同一个 processing block 中，且只能访问该 block 中的处理单元。
wrap 到 processing block（及 scheduler）的映射很简单：block_id = wrap_id % 4。

因此要充分利用一个 SM，一个 thread block 中的线程数量至少需要是 4 个 wrap（128 个线程）。

内存带宽

测内存带宽：https://github.com/NVIDIA/cuda-samples/blob/v11.8/Samples/1_Utilities/bandwidthTest/README.md（make 然后运行即可）

H2D、D2H、D2D 之间可以并行，但同类之间需要串行。

A10 带宽：

 Device 0: NVIDIA A10
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)	Bandwidth(GB/s)
   32000000			25.2

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)	Bandwidth(GB/s)
   32000000			26.3

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)	Bandwidth(GB/s)
   32000000			467.6

register

Register / GPR

通用寄存器 (General Perpose Register, GPR) 有时也直接叫寄存器 (Register)。GPR 通常按个算，一个是32bit，在CUDA的SASS汇编里一般写成 R0，R123 这种格式。
每个线程使用的具体 GPR 数目是编译器根据需要进行配置的。每个 kernel 的所有线程都保有相同数目的 GPR。
最近几代架构单线程最大可用数目是255，因为指令集里 GPR 编码为8位，且 R255=RZ 被用来做恒零的输入，或当输出时表示抛弃输出。

GPU 的 GPR 实际来自于 SRAM 组成的一大块 Register File，每个线程可以分得其中的一部分。一旦线程创建后，物理上的单个 GPR 和线程里的 R0、R1 等就建立了一一对应关系，不再改变，直到退出。

Turing 有 256 个 regular registers（255 个 general-purpose registers (R0–R254) 和 Zero Register (RZ)）和 64 个 uniform registers。同一个 kernel 最多只能使用 256 个寄存器。

register file, register bank conflict

以 Turing 和 Volta 为例，每个 processing block 有一个包含 16384 个 32 位元素的寄存器文件，线程可见的逻辑寄存器就是该寄存器文件上第 8k 到 8k+7 个元素的不同位置（对 wrap 来说，就是 32*8k 到 32*8k+255 个元素的不同位置）。

每个寄存器文件由 2 个 bank 组成，每个 bank 有 2 个 32 位端口。每个寄存器会根据它的名字 % 2 决定它位于哪个 bank（如 R99 位于 bank 1）。
在一个时钟周期内，每个 32 位端口只能满足一个 32 位读，因此如果某个指令包括 3 个及更多的操作数，且至少有 3 个的源寄存器位于同一个双端口 bank 上，那么就会发生 register bank conflict，导致执行时 stall。

以包含三个操作数的 FFMA (single-precision floating-point fused multiply-and-add) 为例：

  • FFMA R6, R98, R99, RX 始终不会发生 bank conflict，因为 R98、R99 位于不同 bank（目的寄存器 R6 没有影响）。
  • FFMA R6, R98, R100, RX 会在 X 为偶数时发生 bank conflict，因为它们都位于 bank 0。

通过仔细分配寄存器，可以减少 bank conflict。

在 Volta 前的架构，寄存器文件为 4 个单端口 bank，因此要注意避免有两个寄存器映射到同一个 bank。

Uniform Registers

为了提高 main datapath (主数据路径？可处理浮点) 的计算吞吐，Turing 引入了一个可以与 main datapath 并行进行的独立的 uniform datapath（但只能处理整数）。

uniform datapath 主要用来优化计算密集的数组运算，它们的 main datapath 中几乎全是浮点指令（比如 FFMA, HMMA），但也有少量的整数运算（比如更新数组下标、循环变量、循环指针、数组下标检查），这些少量的整数运算会打乱 main datapath（不能形成纯粹的 FFMA 序列），影响计算吞吐。
Turing 中，编译器可以选择使用 uniform datapath instructions 将整数运算移到 uniform datapath 中执行。
常规指令可以访问 regular 和 uniform registers，而 uniform datapath 指令几乎只能访问 uniform registers？

Turing 有 64 个 uniform registers：Uniform Zero Register (URZ) 和 63 个 general-purpose uniform registers (UR0–UR62)。

语法

常用

优化

确定 Warp, Block 大小

通常先确定 block 大小（每维线程数量），然后根据块大小和数据规模计算 warp 大小（每维块数量）。

block 的大小通常需要考虑内核的性能特性和 GPU 的资源限制（比如寄存器和共享内存的大小）。

  • 应尽量避免使用小的线程块，因为无法充分利用硬件资源。
  • CPU 与 GPU 中均以行为方式访问内存，因此应尽量保证相邻内存只由同一 block（即同一 SM）读写。

tensor core, tile quantization, wave quantization

介绍：https://developer.nvidia.com/blog/programming-tensor-cores-cuda-9/
使用时的要求与限制、Dimension Quantization Effects：https://docs.nvidia.com/deeplearning/performance/dl-performance-matrix-multiplication/index.html
参考 2：https://docs.nvidia.com/deeplearning/performance/dl-performance-fully-connected/index.html#tc-guidelines

ll /usr/local/cuda/lib64/libcublas.so* 和 ll .../libcudnn.so* 来查看 cublas 和 cudnn 的版本。
ldd 看程序链接的哪个 cublas，然后再 ll 看版本。

在 cuda 11.0 前，如果不满足要求就不会使用 tensor core，所以必须要注意。

tensor core

requirements：见上链接表格。

typical tile dimensions in cuBLAS and performance

可用的 tile size：

  • 256x128 and 128x256 (most efficient)
  • 128x128
  • 256x64 and 64x256
  • 128x64 and 64x128
  • 64x64 (least efficient)：通常只有 GEMM size 特别小时使用。

更大的 tile 有更多的 data reuse，因此会比小 tile 使用更少的带宽、更高效；但可并行的 tile 数量会更少，并行度会更低，可能导致 gpu 利用率低。
当 GEMM 足够大时，即使使用最大的 tile 也能有足够的任务来充分利用 gpu；当 GEMM 较小时，更小的 tile、更小的并行度会一起导致 gpu 利用率更低。
因此增大 GEMM 会提高性能（耗时会增加，但没有 size 增加的快）。有时可以以很小的耗时增加换取更大的 GEMM。

tile quantization

设 matmul 结果矩阵的大小是 n*m，每个 tile 的大小是 a*b，则 n 等于或接近且小于 a*k、m 等于或接近且小于 b*k' 时效率最高。否则会有 tile 包含大量无效的运算。

wave quantization

同时执行的线程块称为一个 wave。

设 matmul 结果矩阵的大小是 n*m，每个 tile 的大小是 a*b，每个 SM 只能执行一个 tile，gpu 共有 S 个 SM。
则 \(\lceil n/a\rceil \times \lceil m/b\rceil\) 等于或接近且小于 S 时效率最高。否则需要 gpu 为多余的 tile 执行一轮新的计算。

测试 tensor core 使用率

https://docs.nvidia.com/nsight-systems/UserGuide/index.html#launching-gpu-metrics-from-the-cli

If you run nsys profile --gpu-metrics-devices all, the Tensor Core utilization can be found in the GUI under the SM instructions/Tensor Active row.

要先 nsys profile --gpu-metrics-devices=help 查看本机哪些 gpu 支持这个操作。

矩乘实现

https://blog.csdn.net/LostUnravel/article/details/138034380

其它

汇编

CUDA的汇编语言分为两种，一种是 Parallel Thread Execution (PTX)，另一种是 Streaming Assembly (SASS)。
PTX 是一种中间语言，可以在不同的 GPU 上运行；SASS 是一种特定的汇编语言，只能在特定的 GPU 上运行。

性能测试

性能分析工具

Nsight System、Nsight Compute、Nsight VS Code

nsight system 可以获取系统级的各项信息。
nsight compute 可以深入分析某个 cuda kernel。
nvprof 已被前两者取代。

带宽计算

https://developer.nvidia.com/blog/how-implement-performance-metrics-cuda-cc/

volta

kernel

Gemm

Gemm (general matrix-matrix multiplication) 是密集广义矩阵乘法 kernel，形式为 \(C=\alpha AB+\beta C\)。计算密集且内存高效。
Gemv (general matrix-vector multiplication) 是稀疏广义矩阵乘法 kernel，形式为 \(C=\alpha AB+\beta C\) 且设 A 为 n*k B 为 k*m，要么 n=1 要么 m=1，因此其计算密度总会小于 1。所以受限于带宽、内存低效。
因此，Gemm 的计算可以复用数据，会比 Gemv 更快，SM 利用率更高。

Gemmk1, Gemv2T, Gemv2N 是 cuBLAS kernel。T/N 应该指操作矩阵是否转置，2T 代表都转置。
Gemmk1 是 Gemm 的变种，Gemv2T/2N 是 Gemv 的变种。

volta_sgemm_128_32
volta_sgemm_128x32_nn_v1

  • volta: GPU architecture
  • s: 累加器类型 (accumulator type)，s 是 single precision
  • gemm: kernel type: matrix multiplication in this case
  • 128: number of elements per CTA in M dimension of the C matrix
  • 32: number of elements per CTA in N dimension of the C matrix
  • 128x32：指矩阵乘法内核的块大小（block size）。具体来说，128x32可能表示内核在执行矩阵乘法时，每次处理128行和32列的数据块。这种块大小的选择通常是为了优化内存访问和计算效率。
  • nn: storage mode for A and B matrices, respectively: “normal” or “no-transpose” (column-major) in your case.
    矩阵乘法的操作模式。nn 表示两个矩阵都不转置；nt 是第一个矩阵不转置，第二个矩阵转置；tn 是第一个矩阵转置，第二个矩阵不转置；tt 是两个矩阵都转置。
  • v1：可能是该 kernel 的实现版本。

cuda driver API

cuda driver API, cuda runtime API

https://docs.nvidia.com/cuda/cuda-runtime-api/driver-vs-runtime-api.html#driver-vs-runtime-api

都是 NVIDIA 提供的用于开发 CUDA 程序的接口，但有所区别：

  • CUDA Runtime API：目的是为了简化 CUDA 编程，提供更高层次的抽象，使开发者可以更方便地编写和管理 CUDA 代码。一般都是用这类。隐藏了许多底层细节，使代码更加简洁易读。
    • 提供了更高级别的抽象，例如 cudaMalloc 和 cudaMemcpy 用于内存管理，cudaStreamCreate 和 cudaStreamDestroy 用于流控制，cudaEventCreate 和 cudaEventRecord 用于事件同步等。
  • CUDA Driver API：目的是为了提供对 CUDA 硬件的低层控制，适合需要精细控制硬件资源的应用场景。提供了更多的灵活性和控制能力，但使用起来相对复杂。
    • 提供了更细粒度的控制，例如 cuCtxCreate 和 cuCtxDestroy 用于上下文管理，cuModuleLoad 和 cuModuleUnload 用于模块加载和卸载，cuLaunchKernel 用于启动内核等。
    • 可以获取 device 更细的信息；需要显式初始化 cuda；可以管理 context 和 module。
    • 接口以 cu 开头。

runtime api 就是对 driver api 的封装，所以两者性能没有差异，但 runtime 不能提供更细的优化。

在编译方面也有区别：runtime API 需要使用 nvida 的编译器进行编译，并且可以将 CUDA kernel 链接到同一个 executable 中。driver API 则完全可以不依赖 nvida 的编译器，可以通过 NVRTC 库来实现对 kernel code 的在线编译，生成 PTX string。

Context Management

相关 API：https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX

cuda context

cuda context 保存了控制和使用 device 的所有信息数据，包括：所有分配的 memory 信息，加载的 module（即加载的 kernel 代码？），CPU 和 GPU 之间 unified memory 的 mapping。

cuda ctx 是一个运行时环境，用于执行 CUDA 内核和管理 GPU 资源（内存、流、事件、模块等）。每个上下文都有自己的独立状态，包括内存分配、内核配置等。每个 CUDA 应用程序通常需要至少一个 ctx 来与 GPU 进行交互。

cuda ctx 可以绑定到 host 上的特定线程。同一时刻：每个线程只能有一个 current ctx，可通过 cuCtxSetCurrent 切换；一个 ctx 只能被一个线程作为 current ctx。

一个线程中创建的上下文只能在该线程中使用，除非显式地切换上下文：使用 cuCtxPushCurrent、cuCtxPopCurrent、cuCtxSetCurrent 可以在不同的线程之间切换当前 CUcontext。

context 不直接暴露在 runtime API，runtime API 默认使用驱动内部提供的 primary context。

context管理了包括stack memory用于每线程临时变量，stack memory 实际是 global memory，context 内部申请的总显存大小为 SMs * max thread per SM * size per thread，因为kernel可能被调度到不同的SM中的物理线程，这样分配保证多个kernel执行时stack memory不冲突，运行时驱动可以动态拓展增加stack memory的总size；用于多device之间同步的cooperative group memory，其作用于cudaLaunchCooperativeKernel()；在多个device上建立页表的显存; 用于临时搬运paged host memory的staging memory，staging memory是pinned host memory，用于src/dst memory来自paged host的cudaMemcpy系列API；还包括用于管理驱动内部同步关系的host/device memory，用作信号量等。module也是按照context进行管理，module包括用户编写CUDA函数的ISA二进制文件也包括驱动内部可能使用到的用户不可见的程序binary，如果同一device创建了多个context，同一份kernel的binary需要被load到不同的context，占用更多显存（其实也有lazy loading的机制）。

CUcontext

typedef CUctx_st* CUcontext：A regular context handle.（这个类型就是一个指针，可以赋值或比较 nullptr）CUcontext 就是一个 cuda context。

primary context

primary context 是一个特殊的 CUDA context，它是 device 的默认 ctx：每个 CUDA device 在启动时都会自动创建一个 primary ctx（不需要显式调用 cuCtxCreate，也不能显式销毁），它在 device 的整个生命周期内始终存在。
可以被多个 host 线程共享，但同其它 ctx 同时只能被一个线程用作 current ctx。

通过cuDevicePrimaryCtxRetain(CUcontext *pctx, CUdevice dev) 获取某个 device 上的 primary ctx，通过cuDevicePrimaryCtxRelease(CUdevice dev) 释放某个 device 的 primary ctx。它们会更新该 primary ctx 的引用计数。

context 栈

CUDA 驱动内部为每个 CPU 线程维护了内部 context 栈，线程活跃 context (current ctx) 为当前栈顶的 context；同一个 context 可以被进程内所有线程使用。

修改栈顶 context 的 API：

  • cuCtxPushCurrent, cuCtxPopCurrent：push/pop。
  • cuCtxSetCurrent：相当于 pop + push。
  • cuCtxCreate：把新创建的 context 入栈。
  • cuCtxDestroy：如果当前线程的栈顶是被销毁的 context 则进行 pop。

注意如果一个 ctx 被放入多个线程的栈中，在某个线程里 destroy，它依然会保留在其它线程 ctx 栈中；并且对于当前线程，如果它不在当前栈顶，也会保留在栈中。之后相关线程必须正确 pop 避免使用它，否则会 CUDA_ERROR_CONTEXT_IS_DESTROYED。

primary ctx 的 cuDevicePrimaryCtxRetain 和 cuDevicePrimaryCtxRelease 不会对 ctx 栈产生影响。

一个 ctx 在同一线程中可以被多次入栈，栈中可以保存多个相同 ctx，因此需要用户控制 ctx 出入栈成对使用，且保证调用 destroy 的逻辑正确（有且仅有一次，destroy 后不再使用）。

driver API 和 runtime API 可以混用，cudaSetDevice 就是将指定 device 的 primary ctx 绑定到当前线程的 current ctx（放到栈顶），相当于 cuDevicePrimaryCtxRetain（不一定需要） + cuSetCurrentCtx。
runtime API 对 driver API 修改 ctx 栈的操作没有感知，因此可以通过 device API 修改 runtime 当前的使用栈，但要保证逻辑正确。如果只使用 runtime API 开发，primary ctx 其实就是对 Device 的抽象。

memory

cudaMalloc/cuMemAlloc 申请的显存和 cudaMallocHost/cuMemAllocHost 申请的 host memory，由当前栈顶的current ctx 管理，其物理位置属于对应 ctx 的 device，但同一 device 创建的所有其它 ctx 都可见。比如：ctx 1 创建的 stream 可以 copy ctx 0 创建的 memory（copy API 不会限制来源）。

如果不同 device 间存在 link (nvlink 或 pcie)，则可通过 cuCtxEnablePeerAccess 允许其它 device 访问 local device memory，即不同 device 创建的 ctx 之间申请的显存也能是相互可见的。

从显存可见范围来看，ctx 与 CPU 进程有很大不同：虚拟地址的范围依赖于 CPU 进程没法在进程之间直接访问。

flag

可以通过 cuCtxSetFlags、cuDevicePrimaryCtxSetFlags 为 cur ctx / primary ctx 设置 flag。
flag 的低三位 (three LSBs, 3 Least Significant Bit) 用于控制有 cuda ctx 的线程在等待 GPU 返回结果时，如何被 os 调度。可以设置 8 种中的一个：

  • CU_CTX_SCHED_SPIN：自旋，延迟低但占 CPU。
  • CU_CTX_SCHED_YIELD：让出 CPU，延迟高但省 CPU。
  • CU_CTX_SCHED_BLOCKING_SYNC：在一个同步原语上阻塞该线程。
  • CU_CTX_SCHED_AUTO：flag 被设为 0 时的默认方式，用一个启发式的方式自行决定。一般就是：如果进程中活跃的 cuda ctx 数量大于 os 逻辑处理器数量就 yield，否则 spin。
  • CU_CTX_LMEM_RESIZE_TO_MAX：让 CUDA 在为某个 kernel 调整 local memory 大小后，不要将本地内存减少到之前的大小。可以避免在启动多个需要大量 local mem 的内核时，要频繁分配和释放本地内存，从而减少内存抖动 (thrashing)。
    该 flag 已 deprecated，现在它是默认行为，且不能被禁用以避免抖动。
  • CU_CTX_COREDUMP_ENABLE：如果没通过 cuCoredumpSetAttributeGlobal 或环境变量全局启用 GPU coredump，可以在创建 ctx 设置该 flag 让 cuda 在该 ctx 抛出异常时创建 coredump。
    coredump 输出设置默认来自 global settings，当该 ctx 成为 current ctx 时可以通过 cuCoredumpSetAttribute 设置。
  • CU_CTX_USER_COREDUMP_ENABLE：如果没通过 cuCoredumpSetAttributeGlobal 或环境变量全局启用 GPU coredump，可以在创建 ctx 设置该 flag 让 cuda 在数据被写入内核空间的特定 pipe 时创建 coredump。
    在创建该 flag 的 ctx 前必须通过 cuCoredumpSetAttributeGlobal 设置管道名称（如 /tmp/cuda_coredump_pipe？），cuda 会监听它并在有写入时触发 coredump。
    设置该 flag 蕴函 CU_CTX_COREDUMP_ENABLE 被设置。coredump 的输出配置也与其一致。
  • CU_CTX_SYNC_MEMOPS：确保在该 ctx 中初始化的同步内存操作始终是同步的（同步操作可能导致异步行为，见 API Synchronization behavior）。
posted @ 2024-12-13 20:04  SovietPower  阅读(21)  评论(0)  编辑  收藏  举报