1 线程束（Warp）¶

1.1 线程束的基本概念¶

当一个线程块（block）被调度到流多处理器（SM）上执行后，线程块中的线程会被进一步划分为线程束。

从逻辑视图来看，一个线程块包含大量线程；从硬件视图来看，这些线程被组织成若干个固定大小的组，每个组就是一个线程束；最终由 SM 的控制逻辑统一调度执行。

由于 GPU 内核的设计原理，这种执行方式可以最大限度地提高吞吐量，从而提高计算性能。在 GPU 编程中，使用线程束可以确保在处理大量数据时最大程度地利用硬件资源。

1.2 获取线程束大小¶

CUDA 提供了内建变量 warpSize，表示当前硬件支持的线程束大小。

#include <iostream>

__global__ void hello_PKU()
{
    printf("warpSize=%d\n", warpSize);
}

int main()
{
    // number of blocks
    int num_blk = 2;
    // number of threads
    int num_thd = 4;

    hello_PKU<<< num_blk, num_thd >>>();

    cudaDeviceSynchronize();

    return 0;
}

运行结果：

warpSize=32
warpSize=32
warpSize=32
warpSize=32
warpSize=32
warpSize=32
warpSize=32
warpSize=32

1.3 分支对线程束效率的影响¶

GPU 中频繁的分支程序会导致利用率的严重下降。

具体来说：

• 不同线程的分支选择：如果有一半线程执行条件为真的分支代码块，另一半线程将执行条件为假的分支代码块
• 由于硬件上每次只能为一个线程束获取一条执行指令，因此近一半的线程将会阻塞，GPU 的硬件利用率将只有 50%

1.4 分支发散示例¶

下面的示例展示了线程束内分支的执行效果。

#include <iostream>
using namespace std;

__global__ void hello_PKU()
{
    const int index = threadIdx.x + 1;

    if (index <= 8)
    {
        printf("id<=8 id=%d\n", index);
    }
    else
    {
        printf("id> 8 id=%d\n", index);
    }
}

int main()
{
    cout << "Hello from CPU!" << endl;

    // number of blocks
    int num_blk = 1;
    // number of threads
    int num_thd = 9;

    // kernel function
    hello_PKU<<< num_blk, num_thd >>>();

    cudaDeviceReset();

    return 0;
}

运行结果：

Hello from CPU!
id> 8 id=9
id<=8 id=1
id<=8 id=2
id<=8 id=3
id<=8 id=4
id<=8 id=5
id<=8 id=6
id<=8 id=7
id<=8 id=8

1.5 回顾矩阵加法中的线程调度¶

回顾一段矩阵加法的 CUDA 程序：

__global__ void add_matrices(
    double *a,
    double *b,
    double *c,
    const int n)
{
    const int index = threadIdx.x;
    c[index] = a[index] + b[index];
    printf("Hello from GPU %d!\n", index);
}

这里有 3 个操作：

1. 读取内存里 a 和 b 的值
2. 把 a 和 b 加起来的操作
3. 把结果 c 存下来的操作

2 线程同步机制¶

2.1 __syncwarp() 与 __syncthreads()¶

__syncwarp() 和 __syncthreads() 都是用来同步线程的函数，但它们作用的范围不同。

2.2 线程块级同步：__syncthreads()¶

下面的示例展示了使用 __syncthreads() 进行线程块级同步。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void block_sync_demo()
{
    __shared__ int sh_data[64];
    int tid = threadIdx.x;

    // 所有线程写入共享内存
    sh_data[tid] = tid;

    // 等待当前 block 内全部 64 个线程执行到此处
    __syncthreads();

    // 安全读取其他线程的数据
    if (tid == 0)
        printf("block 同步完成, sh_data[1] = %d\n", sh_data[1]);
}

int main()
{
    // 1 个 block，64 个线程（包含 2 个 warp）
    block_sync_demo<<<1, 64>>>();
    cudaDeviceSynchronize();
    return 0;
}

2.3 线程束级同步：__syncwarp()¶

下面的示例展示了使用 __syncwarp() 进行线程束级同步。

__global__ void warp_sync_demo()
{
    int tid = threadIdx.x;
    int lane = tid % 32; // 当前线程在 warp 内的编号 0~31

    // 分支内使用 __syncwarp 合法
    if (lane < 16)
    {
        // 仅同步当前这一个 warp 的线程
        __syncwarp();
        printf("warp 内线程 %d 同步完成\n", tid);
    }
}

int main()
{
    // 1 个 block，64 个线程（2 个独立 warp）
    warp_sync_demo<<<1, 64>>>();
    cudaDeviceSynchronize();
    return 0;
}

2.4 __syncthreads() 与 __syncwarp() 的分支对比¶

情况 1：使用 __syncthreads() 在分支内（死锁）

// 情况2: lane < 16
if (lane < 16) {
    __syncthreads(); // 只有前16个线程执行
}

• __syncthreads() 硬性要求：当前线程块内，每一个线程都必须执行到这条指令，严禁分支
• 前 16 个线程：卡在 __syncthreads() 等待块内全部线程汇合，发生死锁！

情况 2：使用 __syncthreads() 在 warp 全参与的分支内（合法）

// 情况1: lane < 32
if (lane < 32) {
    __syncthreads();
}

• block 内所有线程同步，可以运行！
• 只要求参与同步的线程走到指令，允许部分线程跳过，分支里随便用
• 注意：lane < 32 对所有线程都成立，等价于没有分支

3 延迟（Latency）¶

3.1 延迟的基本概念¶

一般来说，延迟越低，意味着 GPU 的效率越高。

3.2 算术指令延迟的影响因素¶

1. 硬件性能：算术指令延迟主要受 GPU 内部处理器的性能影响。处理器的核心数量、运行频率、指令集架构等因素都会影响算术指令的执行速度
2. 指令复杂度：不同的算术指令具有不同的复杂度。例如，加法和减法通常比乘法和除法更简单，因此执行速度更快。指令复杂度越高，算术指令延迟就越大
3. 指令依赖：在某些情况下，一条算术指令需要等待其他指令执行完成后才能开始执行。这种情况下，算术指令延迟会受到依赖指令的影响。如果依赖关系较多，算术指令延迟可能会增大
4. 资源争用：在 GPU 中，多个核心和线程可能会竞争有限的资源，如寄存器、缓存等。如果资源争用严重，可能导致算术指令执行效率降低，从而增加延迟
5. 软件优化：软件层面的优化也会影响算术指令延迟。编译器、驱动和应用程序可以采用各种优化策略来减小算术指令延迟，如指令调度、循环展开、向量化等

3.3 GPU 中的线程调度与延迟隐藏¶

GPU 通过 线程调度 来隐藏访存延迟。当一个线程束因为等待内存数据而被挂起时，SM 会立即切换到其他就绪的线程束继续执行。

调度周期 0：

• 先把线程束排序，SM 依次运行
• 执行 0 号线程束时需读取 a 和 b 数组数据，这个读取操作通常会被合并一起执行
• 线程束 0 进入执行队列，线程束 1~3 在就绪队列等待

3.4 线程挂起¶

线程挂起通常发生在以下场景：

1. 系统资源不足，导致 GPU 线程无法继续执行
2. GPU 线程等待其他线程完成任务，以便继续执行依赖的操作
3. 操作系统或程序中的某种机制导致线程挂起，例如调试、故障排查等
4. GPU 驱动程序或硬件故障导致，问题解决后恢复执行

3.5 线程调度的完整过程¶

调度周期 1：

• 线程束 0 由于访存需要时间（访问了全局内存里的 a 和 b），因此等待时会被挂起
• SM 执行就绪队列中的其他线程束 1
• 内存请求队列中记录了线程束 0 请求的地址（0~31）

调度周期 8：

• 当所有线程束都在等待数据的时候（发出内存请求之后），这个时候 GPU 闲置
• 线程束 0~3 全部进入挂起队列
• 内存请求队列中记录了所有线程束的地址请求

调度周期 9：

• 拿到访存数据之后，线程束 0 将再次排队准备执行下一个指令
• 没有指令时线程束将消亡
• 当所有线程束都结束任务之后，kernel 函数结束，控制权返回 CPU 端

3.6 利特尔法则估算所需线程束数量¶

4 CPU 与 GPU 应对延迟的策略对比¶

4.1 CPU 针对指令延迟的方法¶

CPU 应对指令延迟的策略如下：

• 在一个时钟周期内执行多个指令来提高吞吐量（超标量执行）
• 将指令执行过程划分为多个阶段，使得每个阶段可以在不同的时钟周期内并行执行，从而提高执行速度（流水线）
• 预测程序中的条件分支，以便在实际执行分支指令之前提前获取和执行后续指令（分支预测）
• 根据指令之间的依赖关系，在不影响结果正确性的前提下调整指令的执行顺序，以充分利用处理器资源（乱序执行）
• 通过在 CPU 内部设置多级缓存来减少内存访问延迟，提高指令和数据的访问速度（多级缓存）

4.2 GPU 的延迟隐藏方法¶

GPU 应对指令延迟的策略与 CPU 不同：

• 并行架构：采用多核和多线程技术，同时处理大量的并行任务提高吞吐量
• 单指令多数据（SIMD）：在一个时钟周期内，使用单指令操作多数据
• 本地共享内存：减少全局内存访问的延迟和带宽需求
• 异步计算和流处理：将任务分为多个并行执行的流（流水线），允许不同流之间的任务同时运行，以提高资源利用率
• 内核调度和上下文切换：在不同任务之间动态调度和切换执行资源，以减少执行延迟和提高吞吐量

5 总结¶