3.2. 高级核函数编程

3.2.2.1. SIMT 执行模型

每个 SM 都以一组 32 个并行线程为单位创建、管理、调度和执行线程，这样的一组线程称为线程束。组成一个线程束的各个线程从相同的程序地址一起开始执行，但它们具有各自的指令地址计数器和寄存器状态，因此可以自由地独立分支和执行。线程束一词源于编织，这是最早的并行线技术。半线程束是一个线程束的前半部分或后半部分。四分之一线程束是一个线程束的第一个、第二个、第三个或第四个四分之一部分。

线程束一次执行一条共同的指令，因此，当线程束中的全部 32 个线程遵循相同的执行路径时，就能实现完全的执行效率。如果线程束中的线程通过依赖数据的条件分支发生分歧，线程束就会执行所采取的每条分支路径，并禁用不在该路径上的线程。分支分歧仅发生在线程束内部；不同线程束独立执行，无论它们执行的是相同的代码路径还是互不相交的代码路径。

SIMT 架构与 SIMD（单指令、多数据）向量组织方式相似，都是由一条指令控制多个处理单元。一个关键区别在于，SIMD 向量组织方式向软件暴露 SIMD 宽度，而 SIMT 指令则指定单个线程的执行和分支行为。与 SIMD 向量机不同，SIMT 既使程序员能够为独立的标量线程编写线程级并行代码，也能够为相互协调的线程编写数据并行代码。就正确性而言，程序员基本上可以忽略 SIMT 行为；但是，注意使代码很少要求线程束中的线程发生分歧，可以显著提高性能。在实践中，这类似于缓存行的作用：以正确性为目标进行设计时，可以放心地忽略缓存行大小，但以峰值性能为目标进行设计时，则必须在代码结构中考虑缓存行大小。另一方面，向量架构要求软件将加载合并为向量，并手动管理分歧。

3.2.2.2. 硬件多线程

当一个 SM 被分配一个或多个线程块来执行时，它会将这些线程块划分为线程束，每个线程束由一个线程束调度器调度执行。线程块划分为线程束的方式始终相同；每个线程束包含线程 ID 连续递增的线程，第一个线程束包含线程 0。线程层次结构介绍了线程 ID 与线程块内线程索引之间的关系。

线程块中的线程束总数定义如下：

$\text{ceil}\left( \frac{T}{W_{size}}, 1 \right)$

• T是每个线程块中的线程数，

• Wsize是线程束大小，等于 32，

• ceil(x, y) 等于将 x 向上舍入到最近的 y 的整数倍。

图内文字中文对照

Thread Block

线程块

32 threads

32 个线程

SM

流式多处理器

SM 所处理的每个线程束的执行上下文（程序计数器、寄存器等）在该线程束的整个生命周期内都保留在芯片上。因此，在线程束之间切换不会产生任何开销。在每个指令发射周期，线程束调度器选择一个线程束，该线程束中有线程已准备好执行下一条指令（即线程束中的活跃线程），并向这些线程发射该指令。

每个 SM 都有一组在线程束之间分配的 32 位寄存器，以及一块在线程块之间分配的共享内存。对于给定的核函数，能够驻留在 SM 上并被并发处理的线程块数和线程束数，取决于核函数使用的寄存器和共享内存的数量，以及 SM 上可用的寄存器和共享内存的数量。每个 SM 上的驻留线程块数和线程束数也有最大值。这些限制以及 SM 上可用的寄存器和共享内存数量都取决于设备的计算能力，具体规定见计算能力。如果每个 SM 上没有足够的可用资源来处理至少一个线程块，核函数将无法启动。可以通过占用率一节中介绍的几种方式来确定为一个线程块分配的寄存器总数和共享内存总量。

3.2.2.3. 异步执行特性

近几代 NVIDIA GPU 都包含了异步执行能力，以便在 GPU 内部实现更多的数据移动、计算和同步重叠。这些能力使从 GPU 代码中调用的某些操作能够相对于同一线程块中的其他 GPU 代码异步执行。不应将这种异步执行与第 2.5 节中讨论的异步 CUDA API 混淆，后者使 GPU 核函数启动或内存操作能够相对于彼此或相对于 CPU 异步进行。

计算能力 8.0（NVIDIA Ampere GPU 架构）引入了由硬件加速的从全局内存到共享内存的异步数据复制，以及异步屏障（参见NVIDIA A100 张量核心 GPU 架构）。

计算能力 9.0（NVIDIA Hopper GPU 架构）通过张量内存加速器（TMA）单元、异步事务屏障和异步矩阵乘加操作扩展了异步执行特性，其中 TMA 单元可以将大块数据和多维张量从全局内存传输到共享内存，也可以反向传输（详情请参见深入剖析 Hopper 架构博文。）。

CUDA 提供了可由线程从设备代码中调用的 API，以使用这些特性。异步编程模型定义了异步操作相对于 CUDA 线程的行为。

异步操作是由一个 CUDA 线程发起、但以异步方式执行的操作，就像由另一个线程执行一样，我们将这个线程称为异步线程。在格式良好的程序中，一个或多个 CUDA 线程会与异步操作同步。发起异步操作的 CUDA 线程不必属于参与同步的线程。异步线程始终与发起该操作的 CUDA 线程相关联。

异步操作使用同步对象来发出其完成的信号，该对象可以是屏障，也可以是流水线。高级同步原语详细介绍了这些同步对象，异步数据复制则展示了它们在执行异步内存操作中的作用。

3.2.4.1. 限定作用域的原子操作

第 5.4.5 节概述了 CUDA 中可用的原子函数。在本节中，我们将重点介绍支持C++ 标准原子内存语义的限定作用域的原子操作，这些操作可通过libcu++库或编译器内建函数使用。限定作用域的原子操作提供了在 CUDA 线程层次结构的适当层级实现高效同步的工具，从而兼顾复杂并行算法的正确性和性能。

#include <cuda/atomic>

__global__ void block_scoped_counter() {
    // Shared atomic counter visible only within this block
    __shared__ cuda::atomic<int, cuda::thread_scope_block> counter;

    // Initialize counter (only one thread should do this)
    if (threadIdx.x == 0) {
        counter.store(0, cuda::memory_order_relaxed);
    }
    __syncthreads();

    // All threads in block atomically increment
    int old_value = counter.fetch_add(1, cuda::memory_order_relaxed);

    // Use old_value...
}

__global__ void block_scoped_counter() {
    // Shared counter visible only within this block
    __shared__ int counter;

    // Initialize counter (only one thread should do this)
    if (threadIdx.x == 0) {
        __nv_atomic_store_n(&counter, 0,
                            __NV_ATOMIC_RELAXED,
                            __NV_THREAD_SCOPE_BLOCK);
    }
    __syncthreads();

    // All threads in block atomically increment
    int old_value = __nv_atomic_fetch_add(&counter, 1,
                                          __NV_ATOMIC_RELAXED,
                                          __NV_THREAD_SCOPE_BLOCK);

    // Use old_value...
}

__global__ void producer_consumer() {
    __shared__ int data;
    __shared__ cuda::atomic<bool, cuda::thread_scope_block> ready;

    if (threadIdx.x == 0) {
        // Producer: write data then signal ready
        data = 42;
        ready.store(true, cuda::memory_order_release);  // Release ensures data write is visible
    } else {
        // Consumer: wait for ready signal then read data
        while (!ready.load(cuda::memory_order_acquire)) {  // Acquire ensures data read sees the write
            // spin wait
        }
        int value = data;
        // Process value...
    }
}

__global__ void producer_consumer() {
    __shared__ int data;
    __shared__ bool ready; // Only ready flag needs atomic operations

    if (threadIdx.x == 0) {
        // Producer: write data then signal ready
        data = 42;
        __nv_atomic_store_n(&ready, true,
                            __NV_ATOMIC_RELEASE,
                            __NV_THREAD_SCOPE_BLOCK);  // Release ensures data write is visible
    } else {
        // Consumer: wait for ready signal then read data
        while (!__nv_atomic_load_n(&ready,
                                   __NV_ATOMIC_ACQUIRE,
                                   __NV_THREAD_SCOPE_BLOCK)) {  // Acquire ensures data read sees the write
            // spin wait
        }
        int value = data;
        // Process value...
    }
}

3.2.4.2. 异步屏障

异步屏障与典型的单阶段屏障（__syncthreads()）的区别在于，线程发出其已到达屏障的通知（“到达”）与等待其他线程到达屏障的操作（“等待”）是分离的。这种分离允许线程执行与屏障无关的其他操作，从而更有效地利用等待时间，提高执行效率。异步屏障可用于通过 CUDA 线程实现生产者-消费者模式，也可以让复制操作在完成时向屏障发出信号（“到达”屏障），从而支持内存层次结构中的异步数据复制。

计算能力为 7.0 或更高的设备支持异步屏障。计算能力为 8.0 或更高的设备为共享内存中的异步屏障提供硬件加速，并且允许对线程块内 CUDA 线程的任意子集进行硬件加速同步，从而显著提升了同步粒度。此前的架构只能在整个线程束（__syncwarp()）或整个线程块（__syncthreads()）级别加速同步。

CUDA 编程模型通过cuda::std::barrier提供异步屏障，它是libcu++库中提供的、符合 ISO C++ 标准的屏障。除了实现std::barrier之外，该库还提供了 CUDA 专用扩展，可选择屏障的线程作用域以提升性能，并公开了更底层的cuda::ptx API。cuda::barrier可以通过使用friend函数cuda::device::barrier_native_handle()获取屏障的原生句柄并将其传递给cuda::ptx函数，来与cuda::ptx互操作。CUDA 还为线程块作用域下共享内存中的异步屏障提供了原语 API。

下表概述了可用于在不同线程作用域下进行同步的异步屏障。

线程作用域	内存位置	到达屏障	在屏障处等待	硬件加速	CUDA API
线程块	本地共享内存	允许	允许	是（8.0+）	cuda::barrier、cuda::ptx、原语
簇	本地共享内存	允许	允许	是（9.0+）	cuda::barrier, cuda::ptx
簇	远程共享内存	允许	不允许	是（9.0+）	cuda::barrier, cuda::ptx
设备	全局内存	允许	允许	否	cuda::barrier
系统	全局内存/统一内存	允许	允许	否	cuda::barrier

同步在时间上的拆分

如果不使用异步到达-等待屏障，线程块内的同步可通过__syncthreads()实现，或者在使用协作组时通过block.sync()实现。

#include <cooperative_groups.h>

__global__ void simple_sync(int iteration_count) {
    auto block = cooperative_groups::this_thread_block();

    for (int i = 0; i < iteration_count; ++i) {
        /* code before arrive */

         // Wait for all threads to arrive here.
        block.sync();

        /* code after wait */
    }
}

线程会在同步点（block.sync()）处阻塞，直到所有线程都到达该同步点。此外，同步点之前发生的内存更新保证在同步点之后对线程块中的所有线程可见。

此模式有三个阶段：

• 同步之前的代码执行内存更新，这些更新将在同步之后被读取。

• 同步点。

• 同步之后的代码能够看到同步之前发生的内存更新。

如果改用异步屏障，在时间上拆分的同步模式如下所示。

#include <cuda/barrier>
#include <cooperative_groups.h>

__device__ void compute(float *data, int iteration);

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  using barrier_t = cuda::barrier<cuda::thread_scope_block>;
  __shared__ barrier_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // Initialize barrier with expected arrival count.
    init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    barrier_t::arrival_token token = bar.arrive();

    compute(data, i);

    // Wait for all threads participating in the barrier to complete bar.arrive().
    bar.wait(std::move(token));

    /* code after wait */
  }
}

#include <cuda/ptx>
#include <cooperative_groups.h>

__device__ void compute(float *data, int iteration);

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  __shared__ uint64_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // Initialize barrier with expected arrival count.
    cuda::ptx::mbarrier_init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    uint64_t token = cuda::ptx::mbarrier_arrive(&bar);

    compute(data, i);

    // Wait for all threads participating in the barrier to complete mbarrier_arrive().
    while(!cuda::ptx::mbarrier_try_wait(&bar, token)) {}

    /* code after wait */
  }
}

#include <cuda_awbarrier_primitives.h>
#include <cooperative_groups.h>

__device__ void compute(float *data, int iteration);

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  __shared__ __mbarrier_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // Initialize barrier with expected arrival count.
    __mbarrier_init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    __mbarrier_token_t token = __mbarrier_arrive(&bar);

    compute(data, i);

    // Wait for all threads participating in the barrier to complete __mbarrier_arrive().
    while(!__mbarrier_try_wait(&bar, token, 1000)) {}

    /* code after wait */
  }
}

在此模式中，同步点被拆分为到达点（bar.arrive()）和等待点（bar.wait(std::move(token))）。线程在首次调用bar.arrive()时开始参与cuda::barrier。当线程调用bar.wait(std::move(token))时，它将被阻塞，直到参与线程完成了预期次数的bar.arrive()，该次数就是传递给init()的预期到达次数参数。参与线程调用bar.arrive()之前发生的内存更新，保证在这些线程调用bar.wait(std::move(token))之后对它们可见。请注意，调用bar.arrive()不会阻塞线程，线程可以继续执行其他工作，只要这些工作不依赖于其他参与线程调用bar.arrive()之前发生的内存更新即可。

到达和等待模式有五个阶段：

• 到达之前的代码执行内存更新，这些更新将在等待之后被读取。

• 具有隐式内存栅栏的到达点（即等效于cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_block)）。

• 到达与等待之间的代码。

• 等待点。

• 等待之后的代码能够看到到达之前执行的更新。

有关如何使用异步屏障的完整指南，请参阅异步屏障。

3.2.4.3. 流水线

CUDA 编程模型提供了流水线同步对象，作为一种协调机制，将异步内存复制按顺序组织成多个阶段，从而便于实现双缓冲或多缓冲的生产者-消费者模式。流水线是一个具有头端和尾端的双端队列，按先进先出（FIFO）的顺序处理工作。生产者线程将工作提交到流水线的头端，而消费者线程从流水线的尾端提取工作。

流水线通过libcu++库中的cuda::pipeline API 以及原语 API提供。下表描述了这两种 API 的主要功能。

cuda::pipeline API 的接口更丰富，限制更少，而原语 API 仅支持跟踪从全局内存到共享内存的异步复制，并且有特定的大小和对齐要求。原语 API 提供的功能等效于使用cuda::thread_scope_thread的cuda::pipeline对象。

有关详细的使用模式和示例，请参阅流水线。