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

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

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

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

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<<<1, N>>> 与 kernel<<<N, 1>>>

前者是 kernel 以 N 个线程块启动、每个块只有一个线程，通常用于 kernel 需要执行 N 个独立的任务，每个任务由一个单独的线程块完成。
后者是 kernel 以 1 个线程块启动、这个块有 N 个线程，通常用于 kernel 需要执行一个任务，但这个任务可以被分解为 N 个并行任务，由同一个线程块中的 N 个线程并行完成。

当每个线程执行的任务是独立的，且没有线程间同步的需求时，使用前者可能更合适：

减少线程块内同步/通信、避免共享内存竞争：同一个线程块内的线程可以协同工作，但这也意味着它们可能需要进行线程间同步/通信、可能会竞争访问共享内存。如果任务是独立的，那同步是不必要的，使用单个线程的线程块可以避免。
提高资源利用率：某些情况下，使用单个线程的线程块可以更有效地利用GPU资源？
提高启动效率：启动大量单个线程的线程块可能比启动少量多线程的线程块更有效率，因为每个线程块的启动开销是固定的，而更多的线程块可以更细粒度地利用GPU的并行处理能力？
简化线程索引计算：如果每个线程块内只有一个线程，那每个线程的全局索引可以直接用其线程块索引表示（但这差不了多少？）。
在某些GPU架构中，可能更适合于处理大量小线程块的情况？

内存管理

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

标准的执行流程：

// 1. 分别 CPU 内存、初始化
double *h_x = (double*) malloc(M);  // 还是 cudaMallocHost？
for (int n = 0; n < N; ++n) {  // init
    h_x[n] = a;
}
// 2. 分配 GPU 内存
double *d_x;
cudaMalloc((void **)&d_x, sizeof(double) * N);
// 3. H2D
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
// 4. kernel 计算
kernel<<<...>>>...
// 5. D2H
cudaMemcpy(h_x, d_x, M, cudaMemcpyDeviceToHost);

memcpy

每个 memcpy 都被分成 synchronous 和 asynchronous 版本，但注意同步版本可能异步、异步版本也可能同步：

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 (UMA) 是 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);

cudaMallocHost、cudaMalloc、H2D、kernel launch、D2H 的标准过程太麻烦了，因此用 unified memory 来处理内存、不用管在哪分配、什么时候拷贝。

UM 中 GPU 与 CPU 的虚拟地址 VA 是一致的（GPU 可能多几位）？以便于编程。

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

在 Pascal 架构之前的 GPU（如 K80）不能处理 page fault，因此调用 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 内存和创建页表条目，而是在访问或预取时分配 (demand paging)；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（直接内存访问）时，数据传输更加高效，因为硬件可以直接从已知地址进行操作。

优点：

从 CPU 到 GPU: 使用 pMem 可以显著提高数据主机（CPU）和设备（GPU）之间的传输速度。通常情况下，数据需要被拷贝到一个临时的 pMem，然后再传输到 GPU。但如果一开始数据就在 pMem 中，这个过程就可以省去，从而加速传输。
cuda 不能异步传输 paged 内存(?)中的数据，只有保证数据始终在内存中，即是 pMem 才可异步传输，允许 H2D/D2H/kernel exec 之间并发。
使用 pageable 内存与 GPU 传递数据可能很影响性能？（可在 nsys - cuda HW - show in events view 中确定）
某些 GPU 可将 pMem 映射到 GPU 上的地址空间，消除 host device 间的拷贝（见 Guide - Mapped Memory）。

缺点（不能滥用）：

内存容量限制: 由于 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 是独立执行的任务队列（一系列 kernel 执行、copy 等其它命令），用来管理一组并行执行的线程，与特定的线程块关联。
所有用 cuda 在 GPU 上启动的任务都会在一个 stream 上执行（可能是隐式的默认 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，则会异步拷贝。

stream 与 context

context 是线程间可以共享的全局资源，而 stream 是局部的用于管理执行某个任务的线程的对象。
stream 创建后会与指定 ctx 关联，其中的线程都会在上面执行。stream 间可以通过 ctx 进行同步。

Concurrent

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#asynchronous-concurrent-execution

Concurrent Kernel Execution

GPU 可以并发执行多个 kernel，最多可同时执行的 kernel 数取决于计算能力（见 Table 21）。

同一 ctx 中的 kernel 可并发执行；不同 ctx 间的 kernel 不能并发执行 (concurrently)：类似于 CPU 时间片轮转调度，GPU 会为每个 cuda context 分配一定时间片，在每个时间片内，GPU 会执行该 ctx 的任务，然后 switch ctx 切换到下一个 ctx。
（具体见 MPS - 3.1 Background）

因为不同进程使用不同 ctx，所以进程间不能并发执行，除非通过 MPS。
这种调度方式确保了所有进程都能公平地使用 GPU 资源。

时间片本身应该足够长来避免 ctx switch 的开销影响性能（在 kernel 执行时占比大）。

Overlap of Data Transfer and Kernel Execution / Concurrent Data Transfers

H2D/D2H 与 kernel 执行可以并发执行（需要 GPU 的 device property asyncEngineCount>0）。
H2D 与 D2H 之间也可并发执行（需要 GPU 的 asyncEngineCount==2）。
上面涉及的 host memory 必须是 pin mem。

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
最新的架构：Ada Lovelace

一般说的 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 时无需恢复上下文。

gpu util

nvidia 的 gpu util 代表一段时间内有多长时间至少有一个 kernel 在执行。它仅仅反映 gpu 完全空闲的程度，无法反映 gpu 有多忙、有多少 SM 被使用。

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 内存架构

具体可见 Dissecting the NVidia Turing T4 GPU via Microbenchmarking 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)。

scheduling hierarchy

Dissecting the CUDA scheduling hierarchy

https://zhuanlan.zhihu.com/p/713114525

stream scheduler

FIFO：同一流中的操作按FIFO顺序执行，即先提交的先执行。
流隔离：CUDA流与单个应用程序相关联，不同应用程序的流互不干扰。例如，如果应用程序A0正在运行，则应用程序A1的流不会干扰A0。并行执行：不同流中的操作可以并行执行，但同一流中的操作必须顺序执行。
流优先级：从Maxwell GPU架构（例如Jetson TX1嵌入式板）开始，CUDA提供了一个运行时函数调用，用于为流分配优先级。
当前所有测试过的GPU架构（包括Maxwell、Pascal、Volta和Turing）仅支持两个离散的优先级（高和低）。如果低优先级流占用了一个SM的所有计算资源，则后来提交到高优先级流上的内核可以抢占当前运行的内核。

thread block scheduler

寻找空闲SM 映射CUDA 语义所表达的grid/block/thread 结构
在所有内核被分配到一个流时，线程块会通过所有可用的SM进行循环分配（Round-Robin，RR），先分配到偶数ID的SM，然后是奇数ID的SM
在分配线程块到SM之前，线程块调度器会进行一个占用测试，检查每个SM当前的资源利用情况（线程/warps数量，寄存器，共享内存），以确定是否可以容纳新的线程块。此测试的目的是确保当前的占用率能够满足新内核的需求，从而实现线程块到SM的映射

warp scheduler

每个SM有若干个warp调度器和相应的指令分发单元。
例如，在Pascal架构的GPU中，每个SM有两个warp调度器和两个指令分发单元，每个warp调度器每个时钟周期可以调度两条独立的指令；
图灵架构包含4个Warp scheduler 同时对SM 进行了partition，分为4份；
Maxwell, Pascal, Volta和Turing架构中使用的warp调度策略是松散轮询调度（Loose Round Robin, LRR）。
在LRR策略下，warp以轮询方式调度，当一个warp遇到未满足的依赖（如全局内存未命中）时，它会暂停，使下一个准备好的warp被调度。这种调度策略通过足够的ready warp来隐藏内存访问的延迟，并确保warp之间的公平性。

对于 Turing 每个SM 被划分为4个partition，每个partition 一个scheduler，具体来说：
每个SM有4个Warp Scheduler。
每个Warp Scheduler可以在同一时间调度32个线程。
每个时钟周期内，每个SM可以调度 4*32=128个线程。
每个SM最多支持2048个并发线程，但这些线程并不会在同一个时钟周期内同时运行。因此，对于warp scheduler 来说，多个warp 是通过时分复用的方式实现对scheduler 的占用以及指令的发射，多个warp 间在同一时刻如果处于同一个partition，是串行执行（或者等待前一个warp stall/wait 状态），在不同的partion 之间可以实现并行，从编程的角度我们可以利用这一点。

调度器对warp和SM partition（同时也是调度器id）的映射采用如下简单的方式： scheduler_id = warp_id%4 在同一个block中，warp id 是4 的整数倍的warp 会被调度到同一个partion。一个极端的情况，假如一个block里只有2个warp要做计算，其余warp直接退出。如果这两个要做计算的warp（称为active的warp）对4同余，那么就会造成因为4个partition负载不均衡而产生可能 50% 的性能损失。

优化

确定 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

MPS

https://docs.nvidia.com/deploy/mps/index.html#
控制：https://docs.nvidia.com/deploy/mps/index.html#nvidia-cuda-mps-control

https://www.olcf.ornl.gov/wp-content/uploads/2021/06/MPS_ORNL_20210817.pdf

MPS

MPS (Multi-Process Scheduling) 用于让多进程的 cuda 程序利用 Hyper-Q 真正地在同一 gpu 上并行执行，以提高 gpu 利用率、降低耗时。

关闭 MPS 时，多任务通过时间分片的调度方式共享 GPU、无法并发（见 基础 - Concurrent），且有 ctx switch；开启 MPS 后，多任务共享 Server 的 CUDA Context。

意义：

提高 GPU 利用率：允许不同进程的 kernel 执行和 mem copy 在同一个 GPU 上同时(并发)执行，使用不同部分的 GPU。
降低 GPU 资源消耗：没有 MPS 时，每个使用 GPU 的 CUDA 进程会在 GPU 上分配独立的存储和调度资源。而 MPS server 会分配一份由所有 client 共享的 GPU 存储和调度资源。Volta MPS 支持更高的 client 隔离度，因此降低的不太多？
降低 GPU context 切换：没有 MPS 时，当进程共享 GPU 时，它们的调度资源必须在 GPU 上随时间片调度切换。MPS server 在所有 client 之间共享一套调度资源，消除了 client 之间 ctx 切换的开销。

适用于问题规模不变但计算能力（节点/CPU core/GPU 数）可以提升的场景：

单个进程没有足够的工作、无法充分利用 gpu。此时 ctx switch 的数量和耗时占比就会很大。
这种程序的 blocks-per-grid 会很小。
进程的 threads-per-grid 低导致 gpu 占用率低。

volta 起 MPS client 进程之间的 GPU 地址空间完全隔离，因此不会有读写越界导致 UB 的情况？

当 GPU 发生 fatal fault 时，MPS server 会拒绝所有请求，直到受影响的 GPU 上的所有 client 都退出、重新创建 GPU context；并且无法得知导致错误的 client 是哪个。具体见文档 2.3.3.2. Error Containment。

每个 device 上 Volta MPS server 支持最多 48 个 client CUDA contexts，该限制由各个进程共享。如果超出，进程会在创建 ctx 时报错。

MPI

MPI (Message Passing Interface) 中的概念：

MPI 允许创建进程组（每个进程可属于多个组），每个组中的进程通过其 rank 标识（范围为 0~n-1）。
rank 是一个组内逻辑上的 worker，而进程是实际运行的实例（每个进程可以是多个 rank）。

communicators 是处理进程间通信的对象；intra-communicator 处理同一组内进程间的通信，inter-communicator 处理两个组之间的通信。

3.1 Background

cuda stream 在 driver 中对应 work queue，worker queue 是表示一个 stream 中任务的有序序列的子集的硬件。
支持 Hyper-Q 的 GPU（Kepler 起都支持）有一个并发调度器来从 worker queue 调度属于同一个 cuda ctx 的任务；同一个 ctx 上的任务可并发执行。

不同 cuda ctx 上的任务不能并发执行：GPU 有一个时间片调度器 time sliced scheduler 来调度 work queue 上不同 ctx 的任务。如果单个 ctx 上的任务太少，可能导致 GPU 计算资源利用率低。此时需要通过 MPS。

2.3.5.2 Volta MPS Execution Resource Provisioning

Volta MPS 可以限制执行资源的提供，即限制 client ctx 只能使用部分的可用线程（限制最大使用量，而非为它预分配）。
意义：

降低 client 内存占用：因为(?)每个 MPS client 都有独立的地址空间，每个 client ctx 都会分配独立的 ctx 存储（如栈）和调度资源。这些资源随着 client 可用线程数量的增加而增加。
默认每个 MPS client 都可以使用所有线程（这允许最大程度的调度自由）。但因为 MPS 通常用来同时运行多个进程，所以让所有线程都对每个 client 可用通常是不必要的，分配完整的 ctx 存储可能浪费内存资源。
- 可以用 nvidia-smi 看每个 client 进程的内存。
提高 QoS：可以限制可用的计算带宽来作为一种 QoS 机制。
降低可用线程比例可以将 client 工作限制在特定 SM 内，降低不同 client 的任务之间的影响 (destructive interference?)。

To provide a per-thread stack, CUDA reserves 1kB of GPU memory per thread
This is (2048 threads per SM x 1kB per thread) = 2 MB per SM used, or 164 MB per client for V100 (221 MB for A100)
CUDA_MPS_ACTIVE_THREAD_PERCENTAGE reduces max SM usage, and so reduces memory footprint
Each MPS process also uploads a new copy of the executable code, which adds to the memory footprint.

限制的方式有两类：

限制活跃线程百分比：可分为两种。设置的值会下取整到最近的硬件支持的线程数限制，可通过 cudaDevAttrMultiProcessorCount 查看。
1. uniform partitioning：在 client 进程启动时设置其活跃线程比。无法在中途修改。
  可以通过 MPS control 的 set_default_active_thread_percentage / set_active_thread_percentage 设置后续新创建的 server 创建 client 时的默认比 / 指定 server 上新创建的 client 的百分比。也可以为 MPS 控制进程或 client 设置 CUDA_MPS_ACTIVE_THREAD_PERCENTAGE。
2. non-uniform partitioning：为每个 client cuda ctx 设置其活跃线程比。可以在进程执行中修改。
  uniform 为进程设置的限制可以与对 ctx 的限制同时生效（取最小值），要设置 CUDA_MPS_ENABLE_PER_CTX_DEVICE_MULTIPROCESSOR_PARTITIONING。
通过编程接口 (programmatic interface) 限制可用的 SM 比例：通过 cuCtxCreate_v3 创建 client cuda ctx 并指定 execution affinity CUexecAffinityParam 可以限制 ctx 能使用的 SM 数量。设置的值会上取整到最近的硬件支持的 SM 数限制，可通过 cuCtxGetExecAffinity 查看。
- 好像只有一种 CUexecAffinityParam：CU_EXEC_AFFINITY_TYPE_SM_COUNT：通过 CUexecAffinitySmCount 限制 ctx 可用的 SM 百分比。
- 比限制线程数更细粒度和灵活。

可参考的限制策略：

uniformly partition：将线程均匀分给每个 client，比如设置 active thread percentage 为 100% / 0.5n（n 为 client 进程数）。0.5 是允许 client 利用可能有的空闲资源。
non-uniformly partition：根据每个 client 的工作负载分配活跃线程比。可以使不同 client 的工作集中到不相交的 SM 上，减少 client 间的影响。
最优策略：在已知每个 client 的执行需求时，精确地限制每个 client 能用的 SM 数量。

Compute Mode

GPU Compute Mode 影响了资源在计算时如何被分配和利用。用来调整 GPU 如何处理多进程。

计算模式有三种：

DEFAULT：多个进程可同时使用 GPU 资源，这些进程的各线程都可以并行向 GPU 提交任务。
多进程可能会因为竞争 GPU 资源导致性能下降。
EXCLUSIVE_PROCESS：同一时间只有一个进程可以使用 GPU，该进程的各线程可以并行向 GPU 提交任务。
可以减少其它进程对 GPU 资源的影响，提升单个进程的性能和降低延迟。
PROHIBITED：都不能用 GPU。

通过 nvidia-smi -c MODE 调整计算模式。

MPS 可以让所有 MPS client 在 exclusive process 下也像 default 一样通过 MPS server 同时使用 GPU。
在 shared system 使用 MPS 时，最好使用 exclusive process mode 以保证只有 MPS server 在使用 GPU。

其它

汇编

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 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，context 管理，module 管理。
CUDA Driver API：目的是为了提供对 CUDA 硬件的低层控制，适合需要精细控制硬件资源的应用场景。提供了更多的灵活性和控制能力，但使用起来复杂。
- 提供了更细粒度的控制，例如 cuCtxCreate 和 cuCtxDestroy 用于上下文管理，cuModuleLoad 和 cuModuleUnload 用于模块加载和卸载，cuLaunchKernel 用于启动内核等。
- runtime 的 kernel 会在初始化时自动加载并始终保留，而 driver 可以动态加载 module、只将目前需要的保留。
- 可以获取 device 更细的信息；需要显式初始化 cuda；可以管理 context 和 module。
- 接口以 cu 开头。

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

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

device id

TF 的分类：https://github.com/tensorflow/tensorflow/blob/5c0ef1423f51306d2464918e2126943d0ec1979e/tensorflow/core/common_runtime/device/device_id.h

Context Management

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

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#context

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、cuCtxGetCurrent切换和获取；一个 ctx 只能被一个线程作为 current ctx。如果需要 ctx 的函数使用了一个非 current 的 ctx 则会返回 CUDA_ERROR_INVALID_CONTEXT。

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

context 不直接暴露在 runtime API。runtime API 默认使用 device 的 primary context，如果使用 driver API 指定了 current ctx 则使用指定的。
即 runtime 下多个线程默认会共享一个 ctx。使用 driver API 创建、切换 ctx 也有很大开销。

每个 ctx 都有引用计数，可通过 cuCtxCreate、cuCtxAttach 增加，通过 cuCtxDestroy、cuCtxDetach 减少，并在计数为 0 时销毁释放。这使第三方库共享 ctx 很方便。

每个 ctx 有自己的地址空间，因此不同 ctx 的 CUdeviceptr 指向不同内存地址。

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的机制）。

一个 context 与一个进程类似，所以在 ctx 之间通常会有数据隔离。但两者也有很多不同：

从显存可见范围来看：虚拟地址的范围依赖于 CPU 进程，不能在进程之间直接访问，而 ctx 之间可以通过 IPC？

ctx 不能直接访问其它 ctx 的内存，但 this could be done for example with host-based IPC to pass the address from one process to another, or perhaps a more "manual" method.

The GPU is a collection of multiple independent engines that can run different contexts simultaneously. The two most common engines are the GR (3D + 2D + compute) and the Asynchronous Copy Engine. Each engine can run 1 context at a time. Each engine has a separate page table pointer via the context. On a context switch the TLBs are invalidated.

CUcontext

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

CUcontext 就是一个 cuda context。

primary context

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DRIVER.html

primary context 是一个特殊的 CUDA context，它是 device 的默认 ctx：每个进程在每个 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 的引用计数，并在计数为 0 时销毁。

对 runtime API 来说，一个 device 与它的 primary ctx 是等价的。

同步异步

API Synchronization behavior

https://docs.nvidia.com/cuda/cuda-driver-api/api-sync-behavior.html#api-sync-behavior

memcpy 和 memset API 有同步和异步两种形式，异步 API 以 Async 为后缀，但所有这些 API 可能表现出同步行为也可能是异步的（取决于实参）。

所有 cuda API 都可能会因为内部资源不可用/竞争等原因导致阻塞或同步。这类行为可能会随版本变化、不能依赖。

memcpy 行为见 cuda 笔记 - memcpy 或文档。
对于 memset：异步版本始终是异步的；同步版本只有在目标内存是 pinned host mem 时才是同步的，否则是异步的。
所有 kernel launch 相对于 host 都是异步的。

posted @ 2024-12-13 20:05 SovietPower 阅读(286) 评论(0) 编辑收藏举报

刷新页面返回顶部

登录后才能查看或发表评论，立即登录或者逛逛博客园首页

相关博文：

· CppQuiz 笔记

· CMake 笔记

· CUDA基础

· CUDA笔记

· GPU入门

阅读排行：
· 震惊！C++程序真的从main开始吗？99%的程序员都答错了
· 【硬核科普】Trae如何「偷看」你的代码？零基础破解AI编程运行原理
· 单元测试从入门到精通
· 上周热点回顾（3.3-3.9）
· winform 绘制太阳，地球，月球运作规律

历史上的今天：
2018-12-13 BZOJ.1312.[Neerc2006]Hard Life(分数规划最大权闭合子图)

公告

找导师和找对象一样，都是迷惑行为。

------------------------------
~~An ~~OIer~~ACMer from SDSH.~~
~~大四退役ACMer~~
刚毕业的元中心牛马
欢迎各位神犇一起开黑 /kel

------------------------------
-> Click here!
（然而图挂了一直懒得修_(:з」∠)_

------------------------------

-----------学长学姐-----------
IIIIIIIIIU
thmyl
TRTTG
ZlycerQan
------------同级orz-----------
lovewhy
hwim大佬
Grary
11101001大佬
自为风月马前卒大佬
------------学弟学妹----------
GTBA
loceaner
Nanjo_Qin~~o_Qo_Q~~
sdgzyorz
xmex
wxyww
ydclyq
yxj
复杂的哈皮狗
----------外校dalao-----------
Milky Way
Ning_Mew

昵称： SovietPower
园龄： 7年11个月
粉丝： 100
关注： 41

+加关注

2025年3月

日

一

二

三

四

五

六

SovietPower

因你而起的一泓喜悲权当年轻留个纪念

cuda 笔记

cuda

基础

Grid, Block, Thread

SM

Warp

SIMT

kernel

内存管理

Unified Memory

Pinned Memory

Stream

Concurrent

GPU

bank conflict

memory hierarchy

register

scheduling hierarchy

优化

确定 Warp, Block 大小

tensor core, tile quantization, wave quantization

矩乘

MPS

其它

汇编

性能测试

volta

kernel

基础

Context Management

同步异步

公告

搜索

最新随笔

随笔分类 (1848)

随笔档案 (868)

SovietPower

因你而起的一泓喜悲 权当年轻 留个纪念

cuda 笔记

cuda

基础

Grid, Block, Thread

SM

Warp

SIMT

kernel

内存管理

Unified Memory

Pinned Memory

Stream

Concurrent

GPU

bank conflict

memory hierarchy

register

scheduling hierarchy

优化

确定 Warp, Block 大小

tensor core, tile quantization, wave quantization

矩乘

MPS

其它

汇编

性能测试

volta

kernel

基础

Context Management

同步 异步

公告

搜索

最新随笔

随笔分类 (1848)

随笔档案 (868)

因你而起的一泓喜悲权当年轻留个纪念

同步异步