GPU 性能讨论

在计算机体系结构的语境下，性能主要指两个量：时间和吞吐量。

GPU 在进行运算时，比如大模型推理，我们通常希望端到端的执行时间越短越好，比如一次 prefill 或者一次 decode。

在端到端的执行时，首先从单个 GPU 操作视角看，不同的操作都会从头开始执行到结束；从多个操作视角看，多个操作的执行是可以重叠的，GTE 和 CE 的操作可以并行，一个 SM 也可以并发执行多个 kernel 的 block，一个 partition 也可以并发执行一个 kernel 的 cuda core 指令和 tensor core 指令。

下面进行详细讨论。

完整流水线

Host            User   Mode   Driver
                |     |     |     |
                PB    PB    PB    PB
                |     |     |     |
GPU frontend    Host  Interface  Unit
                |     |     |     |
                FIFO  FIFO  FIFO  FIFO
                |     |     |     |
                Command    Processor
                |              |
GPU backend     GTE            CE
                |
                SM - SMSP
                     |  |
                  CUDA  Tensor
                  core  core
                  |  |
              A(FP)  B(FP/INT)

单个 Kernel 性能分析

首先，我们对单个 kernel 的性能进行分析。主要关注的硬件是 GPU 后端的 SM。

Kernel 的理论性能

Kernel 理论性能具体分析

我们可以通过一个 vector add 例子来具体分析一下。

#include <cuda_runtime.h>
#include <iostream>

__global__ void vectorAdd(const float *A, const float *B, float *C,
                          int numElements) {         // line 5
    // 计算全局线程索引
    int i = blockDim.x * blockIdx.x + threadIdx.x;   // line 7

    // 边界检查，防止越界访存
    if (i < numElements) {                           // line 10
        C[i] = A[i] + B[i];
    }
}                                                    // line 13

int main() {
    // 仅为了通过编译，省略了 Host 端的内存分配和 Kernel 调用代码
    return 0;
}

运行指令：

nvcc -arch=sm_86 -O3 -cubin vector_add.cu -o vector_add.cubin -lineinfo
nvdisasm -c -g vector_add.cubin > vector_add.sass

SASS 指令如下：

.text._Z9vectorAddPKfS0_Pfi:
line 5
        /*0000*/                   MOV R1, c[0x0][0x28] ;
line 7
        /*0010*/                   S2R R6, SR_CTAID.X ;
        /*0020*/                   S2R R3, SR_TID.X ;
        /*0030*/                   IMAD R6, R6, c[0x0][0x0], R3 ;
line 10
        /*0040*/                   ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ;
        /*0050*/               @P0 EXIT ;
line 11
        /*0060*/                   MOV R7, 0x4 ;
        /*0070*/                   ULDC.64 UR4, c[0x0][0x118] ;  # 没起作用
        /*0080*/                   IMAD.WIDE R4, R6, R7, c[0x0][0x168] ;
        /*0090*/                   IMAD.WIDE R2, R6.reuse, R7.reuse, c[0x0][0x160] ;
        /*00a0*/                   LDG.E R4, [R4.64] ;
        /*00b0*/                   LDG.E R3, [R2.64] ;
        /*00c0*/                   IMAD.WIDE R6, R6, R7, c[0x0][0x170] ;
        /*00d0*/                   FADD R9, R4, R3 ;             # 算出 C[i]
        /*00e0*/                   STG.E [R6.64], R9 ;           # 写入 C[i]
line 13
        /*00f0*/                   EXIT ;

假设 kernel launch 的时候，block size 是 256（8 个 warp），grid size 是 504。GPU 后端操作如下：

我们只关注单精度浮点运算，可以看到，运算 FADD 和访存 LDG STG 实际上是交错的。而且，必须先执行两次

Kernel 的实测性能

硬件的性能上限与瞬时性能

举个例子，ga102 tensor core fp16mfp16a 的硬件性能上限是 149.7 TFLOPS，即 1s 1740 M 个周期内，最多可以做 149.7 T 的运算。即 1 个周期内，最多可以做 128 x 4 x 84 次运算。

要达到这个值，我们实际上需要 84 个 SM 里面的 84x4 = 336 个 tensor core 都在执行有效运算。而一个 tensor core，要么一个周期进行 128 次有效运算，要么进行 0 次有效运算。

所以，某一个周期内，如果 336 个 tensor core 中，只有 n 个 tensor core 在进行有效运算，那么这个周期内的硬件瞬时性能就是 128 x n。当 n = 336 时，这个周期内的硬件瞬时性能达到硬件性能上限。

需要注意的是，我们说的硬件性能上限，是不可达上界，实际可达的是微基准实测峰值（如 V100 125->96 TFLOPS）。

Kernel 的实测性能具体分析

NCU Roofline 计算

TODO

多个 Kernel 分析

现在，我们对涉及多个 kernel 的几种常见情况进行分析，串联起 CPU、GPU 前端和 GPU 后端。

同一个 stream 里的多次 kernel launch

TODO 实际上更复杂 CWD

kernel_A<<<...>>>();
kernel_B<<<...>>>();

不同 stream 里的多次 kernel launch

kernel_A<<<..., stream_1>>>();
kernel_B<<<..., stream_2>>>();

同步：CPU 等 GPU

这类同步的本质是跨越 PCIe 的物理阻塞。

UMD 必须在 PushBuffer 中强行注入向 Host 内存写状态的指令（Semaphore Write），而 CPU 侧的线程会被挂起或陷入自旋轮询（Spinning），直到总线把完成信号传回内存。

cudaDeviceSynchronize

kernel_A<<<..., stream_1>>>();
kernel_B<<<..., stream_2>>>();
cudaDeviceSynchronize(); // CPU 在此处死等，直到整个 GPU 彻底安静

cudaStreamSynchronize

kernel_A<<<..., stream_1>>>();
kernel_B<<<..., stream_2>>>();
cudaStreamSynchronize(stream_1); // CPU 只等 stream_1，此时 stream_2 如果在跑，CPU 不管

cudaEventSynchronize

kernel_A<<<..., stream_1>>>();
cudaEventRecord(event, stream_1);
kernel_B<<<..., stream_1>>>();

cudaEventSynchronize(event); // CPU 死等这个特定标记被触发

同步：GPU stream 等 GPU stream

// 假设有个 event_ready 创建时带有 cudaEventDisableTiming 标志
kernel_A<<<..., stream_1>>>();
cudaEventRecord(event_ready, stream_1); // A 做完，插旗

kernel_C<<<..., stream_2>>>();
cudaStreamWaitEvent(stream_2, event_ready, 0); // stream_2 看到旗子才能放行
kernel_B<<<..., stream_2>>>(); // 强依赖于 A 的结果

同步的坑：Default Stream

同步的坑：隐式同步

有时候你明明写了一堆异步的流，结果用 Nsys 一抓 Timeline，发现它们还是串行了。这往往是因为你在 CPU 端触发了隐式同步操作。

一旦执行以下操作，驱动 (UMD) 会被迫把手头所有 PushBuffer 里的东西先冲刷 (Flush) 到 GPU 执行完，导致全局阻塞：