GPU 架构与 CUDA

GPU 架构与 CUDA

GPU 通过提供数千个核心以支撑大规模并行，彻底改变了 AI。本节讲 GPU 与 CPU 在设计哲学上的差异、GPU 的内存层级、用 CUDA C++ 编程、SIMT 执行模型、访存模式、同步、流（stream）、性能剖析，以及 NVIDIA 各代 GPU——这些都是写出并读懂 GPU 内核所需的知识。

• 想要一套带完整可运行示例的 CUDA 动手教程，可以看配套仓库：github.com/HenryNdubuaku/cuda-tutorials。

• 一颗现代 NVIDIA GPU 有超过 1 万个 CUDA 核心。一颗 CPU 只有 4-128 个核心。这 100-1000 倍的核心优势，正是 GPU 统治 ML 的原因：训练一个 transformer 需要数万亿次乘加，而 GPU 能以 CPU 无法企及的规模并行处理。

• 哪怕你从不自己写 CUDA 内核，理解 GPU 架构也能解释：为什么 batch size 重要（要有足够的工作来填满 GPU）、为什么内存通常是瓶颈（而不是算力）、以及为什么某些操作（scatter、条件分支）在 GPU 上很慢。

GPU vs CPU：根本不同的设计

• CPU 为**延迟（latency）**而设计：最小化完成单个任务的时间。它把大部分晶体管预算花在缓存、分支预测器、乱序执行上——都是让单个线程变快的招数。

• GPU 为**吞吐（throughput）**而设计：最大化每秒完成的任务数。它把大部分晶体管花在执行单元（ALU）上。单个线程很慢，但有成千上万个。

• GPU 的内存带宽优势（10-30 倍）往往比它的算力优势更重要。许多 ML 操作是访存受限的（逐元素操作、归一化、注意力），GPU 的带宽让它能把数据喂给核心的速度够快。

GPU 内存层级

• 理解 GPU 内存至关重要，因为内存访问才是主要瓶颈，而不是计算。

• 寄存器最快但最有限。每个线程有一组私有寄存器（通常最多 255 个）。每个线程用太多寄存器会降低占用率（occupancy）（能同时跑的线程更少）。

• **共享内存（shared memory）是由程序员管理的缓存，被一个块（block）里所有线程共享。它是写出快速 CUDA 内核的关键：把一块数据从慢的全局内存加载到快的共享内存，再在上面做计算。这就是主导 GPU 编程的分块（tiling）**模式。

• 全局内存（HBM）：GPU 的主内存（VRAM）。大但慢（400 周期延迟）。所有数据都从这里开始、也在这里结束。内核优化的目标就是尽量减少全局内存访问。

CUDA 编程模型

• CUDA（Compute Unified Device Architecture，统一计算设备架构）是 NVIDIA 为 GPU 设计的编程模型。你写内核（kernel）：跑在 GPU 上、由数千个线程同时执行的函数。

层级：Grid、Block、Thread

Grid（整个启动）
├── Block (0,0)
│   ├── Thread (0,0)
│   ├── Thread (1,0)
│   ├── Thread (2,0)
│   └── ...（每个块最多 1024 个线程）
├── Block (1,0)
│   ├── Thread (0,0)
│   └── ...
└── ...（可能有上百万个块）
​

• 线程（Thread）：最小单位。每个线程在所在块内有唯一 ID（threadIdx.x）。

• 块（Block）：一组能共享内存并同步的线程。块 ID：blockIdx.x。块大小：blockDim.x（最多 1024 个线程）。

• 网格（Grid）：一次内核启动的所有块。可以是 1D、2D 或 3D。

• 每个线程算自己的全局索引：int idx = blockIdx.x * blockDim.x + threadIdx.x;

你的第一个 CUDA 内核

// vector_add.cu —— CUDA 源文件（.cu 后缀）

#include <stdio.h>

// __global__ 标记这是一个 GPU 内核（从 CPU 调用，在 GPU 上运行）
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {           // 边界检查（网格可能比数据大）
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // 约 100 万元素
    size_t bytes = n * sizeof(float);

    // 分配 host（CPU）内存
    float *h_a = new float[n];
    float *h_b = new float[n];
    float *h_c = new float[n];

    // 初始化
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    // 分配 device（GPU）内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 把数据从 CPU 复制到 GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 启动内核：每块 256 个线程，块数足够覆盖 n 个元素
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;  // 向上取整
    vector_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

    // 把结果从 GPU 复制回 CPU
    cudaMemcpy(h_c, d_a, bytes, cudaMemcpyDeviceToHost);

    // 验证
    printf("c[0] = %f (expected 3.0)\n", h_c[0]);

    // 释放内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    delete[] h_a; delete[] h_b; delete[] h_c;

    return 0;
}
​

# 用 NVIDIA 的编译器编译
nvcc -O3 -o vector_add vector_add.cu
./vector_add
​

• CUDA 里的关键 C++ 概念：
  • __global__：CUDA 关键字，标记一个内核函数。从 CPU（host）调用，在 GPU（device）上运行。
  • <<<grid_size, block_size>>>：内核启动语法，指定用多少个块和线程。
  • cudaMalloc / cudaFree：分配/释放 GPU 内存（类似 new/delete，但是针对 GPU）。
  • cudaMemcpy：在 CPU 和 GPU 之间复制数据。这往往是最大的瓶颈（PCIe 带宽约 32 GB/s，而 GPU 内存带宽约 3 TB/s）。

Warp 与 SIMT

• GPU 以 32 个线程为一组来执行，称为warp。一个 warp 里所有 32 个线程在同一时刻执行同一条指令（Single Instruction, Multiple Threads —— SIMT，单指令多线程）。这是 GPU 版的 SIMD，只不过是线程层面的。

• **warp 分歧（warp divergence）**发生在同一个 warp 里的线程走了 if 语句的不同分支。GPU 没法在一个 warp 里同时执行两条不同的指令，所以它会把两条分支串行执行，把不该参与的线程屏蔽掉。这会让性能减半（甚至更糟）。

// 坏：warp 分歧（同一 warp 的线程走不同路径）
if (threadIdx.x % 2 == 0) {
    c[idx] = a[idx] + b[idx];    // 偶数线程做这个
} else {
    c[idx] = a[idx] - b[idx];    // 奇数线程做这个（同一 warp，被串行化）
}

// 更好：无分支（没有分歧）
float sign = (threadIdx.x % 2 == 0) ? 1.0f : -1.0f;
c[idx] = a[idx] + sign * b[idx];  // 所有线程执行同一条指令
​

内存合并

• 合并访问（coalesced access）：当相邻线程访问相邻的内存地址时，GPU 会把它们合并成一次内存事务。这对性能至关重要。

// 好：合并 —— 线程 0 读 a[0]，线程 1 读 a[1]，...
c[idx] = a[idx] + b[idx];

// 坏：跨步 —— 线程 0 读 a[0]，线程 1 读 a[stride]，...
c[idx] = a[idx * stride] + b[idx * stride];  // stride > 1 会浪费带宽
​

• 对一个 32 线程的 warp，合并访问一次事务就读 128 字节（32 × 4 字节，float32）。跨步访问则要多次事务，每次都读 128 字节却只用其中一小部分。最坏情况是步长 32：每次事务读 128 字节，却只有一个线程用到 4 字节（利用率 3%）。

共享内存与分块

• **分块模式（tiling pattern）**是最重要的 GPU 优化技巧。思路是：把一块数据从慢的全局内存加载到快的共享内存，在上面计算，再把结果写回去。

// 用共享内存分块的矩阵乘法（简化版）
__global__ void matmul_tiled(const float* A, const float* B, float* C,
                              int M, int N, int K) {
    // 共享内存：A 的一块瓦片和 B 的一块瓦片
    __shared__ float tile_A[TILE_SIZE][TILE_SIZE];
    __shared__ float tile_B[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    // 遍历各块瓦片
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 把 A 和 B 各加载一块到共享内存
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            tile_A[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            tile_A[threadIdx.y][threadIdx.x] = 0.0f;

        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            tile_B[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tile_B[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();  // 等所有线程加载完

        // 从这块瓦片计算部分点积
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tile_A[threadIdx.y][k] * tile_B[k][threadIdx.x];
        }

        __syncthreads();  // 加载下一块之前要等一下
    }

    if (row < M && col < N)
        C[row * N + col] = sum;
}
​

• __shared__：声明被块内所有线程共享的内存（快，在片上）。
• __syncthreads()：一个屏障，等块内所有线程都到达这一点。在写共享内存与读共享内存之间必须用它（否则有些线程会读到旧数据）。
• 为什么分块有效：不分块的话，每个线程每次乘法都要从全局内存加载。分块之后，一整块 TILE_SIZE × TILE_SIZE 的数据只加载一次到共享内存，然后被块内所有线程复用。复用因子是 TILE_SIZE，全局内存流量随之降低这么多倍。

流与并发

• 默认情况下，CUDA 操作是串行的：CPU 启动一个内核，等它完成，再启动下一个。**流（stream）**能实现重叠：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 这些操作可以重叠：不同流并发执行
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);

kernel1<<<grid, block, 0, stream1>>>(d_a, d_c);
kernel2<<<grid, block, 0, stream2>>>(d_b, d_d);
​

• 流让数据传输与计算重叠：当一个流的内核在跑时，另一个流在复制数据。这隐藏了 PCIe 传输延迟，让 GPU 保持忙碌。

剖析 CUDA 代码

# NVIDIA Nsight Compute：内核级剖析
ncu --set full ./my_program

# NVIDIA Nsight Systems：系统级时间线
nsys profile ./my_program

# 快速指标
ncu --metrics sm__throughput,dram__throughput ./my_program
​

• 要看什么：
  • 占用率（Occupancy）：SM 容量被用到的比例。低占用率（< 50%）意味着线程太少，藏不住内存延迟。原因：每个线程寄存器太多、每个块共享内存太多。
  • 内存吞吐：与峰值带宽对比。如果低于峰值的 50%，说明访存模式低效（未合并、bank 冲突）。
  • 算力吞吐：与峰值 FLOPS 对比。如果内存吞吐和算力吞吐都低，说明内核是延迟受限的（并行度不够）。

高级优化技巧

• 在合并访问和共享内存分块这些基础之上，高性能的 GPU（和 CPU）代码还会用到几招进阶技巧：

数据布局：AoS vs SoA

• 结构数组（Array of Structures，AoS）：每个元素把所有字段放在一起。[{x,y,z}, {x,y,z}, {x,y,z}]。
• 数组结构（Structure of Arrays，SoA）：每个字段存在自己连续的数组里。{[x,x,x], [y,y,y], [z,z,z]}。

// AoS：对 SIMD/GPU 不利（访问所有 x 值会碰到不连续的内存）
struct Particle { float x, y, z, mass; };
Particle particles[N];
// particles[0].x、particles[1].x 之间隔了 16 字节

// SoA：对 SIMD/GPU 友好（所有 x 值都是连续的）
struct Particles {
    float x[N], y[N], z[N], mass[N];
};
// x[0]、x[1] 之间隔 4 字节 —— 完美适合合并访问和 SIMD
​

• 对数据并行工作负载（SIMD、GPU），SoA 几乎总是更快。只有当你总是同时访问一个元素的所有字段时，AoS 才更好（在数值代码里很少见）。PyTorch 张量天然就是 SoA：每个特征都是一个连续的维度。

软件预取

• 可以让 CPU 在需要数据之前就开始加载，从而隐藏内存延迟：

#include <xmmintrin.h>  // 用于 _mm_prefetch

for (int i = 0; i < n; i += 4) {
    _mm_prefetch((char*)(a + i + 64), _MM_HINT_T0);  // 提前 64 个元素预取
    // 用 SIMD 处理 a[i:i+4]
    __m128 va = _mm_load_ps(a + i);
    // ...
}
​

• 预取指令只是一个提示：如果数据已经在缓存里，它就是空操作；如果不在，CPU 就在后台开始取，同时执行其他指令。预取距离（本例中提前 64 个元素）要根据内存延迟和循环迭代时间来调。

内核融合

• **内核融合（kernel fusion）**把多个操作合成一个内核，避免把中间结果写回内存。这是 ML 中影响最大的 GPU 优化：

// 未融合：3 次内核启动，3 次全局内存往返
y = matmul(x, W)     // 把 y 写到全局内存
z = y + bias          // 读 y，写 z
out = relu(z)         // 读 z，写 out

// 融合：1 次内核启动，1 次全局内存写入
out = fused_matmul_bias_relu(x, W, bias)  // y 和 z 从不离开 SRAM
​

• 对访存受限的操作（bias 加法、ReLU、layer norm），内存流量主导执行时间。融合能彻底消除这些流量。PyTorch 的 torch.compile 和 Triton 能自动或在极少工作量下实现融合。

混合精度内核

• 用低精度（FP16、BF16、INT8）做计算，用高精度（FP32）做累加，能兼得两者之长：

// Tensor Core：FP16 矩阵相乘，FP32 累加
// 每条 Tensor Core 指令：D (FP32) = A (FP16) × B (FP16) + C (FP32)
nvcuda::wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
​

• FP16 比 FP32 小一半，所以它把内存带宽（通常是瓶颈）翻倍，并在缓存里塞下两倍的数据。Tensor Core 处理 FP16 的速度是 FP32 CUDA 核心的 8-16 倍。这就是为什么混合精度训练（第 6 章）能在精度几乎无损的情况下带来 2-3 倍加速。

内存池分配器

• cudaMalloc 很慢（每次约 1 ms），因为它要与 GPU 同步。在每次迭代都分配临时缓冲区的训练循环里，这会累积起来。

• 内存池（memory pool）（PyTorch 的缓存分配器、CUDA 内存池）会预分配一大块 GPU 内存，并从中子分配，无需系统调用：

# PyTorch 自动就做了这件事——但理解为什么很重要
# 每次 torch.empty() 都从池子里复用内存，不调用 cudaMalloc
temp = torch.empty(1024, 1024, device='cuda')  # 微秒级，不是毫秒级
​

• 这就是为什么 PyTorch 的 torch.cuda.memory_allocated() 和 torch.cuda.max_memory_allocated() 不一样：allocated 是当前在用的，max 是峰值（池子可能持有比当前在用更多的内存）。

性能剖析导向的优化

• 不要盲目优化。先剖析，找到瓶颈，优化它，再重新剖析。屋顶线模型（第 01 节）告诉你瓶颈是内存还是算力：

  • 访存受限（算术强度低）：优化数据布局（SoA）、融合内核、用更低精度、预取。
  • 算力受限（算术强度高）：用 Tensor Core、增加并行度、用更快的指令（FMA）。
  • 延迟受限（并行度不够）：提高占用率、减少寄存器使用、启动更多线程。

• 大多数 ML 工作负载是访存受限的。一个出人意料的推论：更快的 GPU（更多 FLOPS）往往没用；更快的内存（HBM3 vs HBM2e）才有用。这就是为什么 A100→H100 的升级不只是 FLOPS 的提升——H100 的内存带宽也是 2 倍。

NVIDIA GPU 各代

• Tensor Core是专用的矩阵乘单元。一条 Tensor Core 指令在一个周期内完成一次 4×4 矩阵乘（D = A×B + C）。普通 CUDA 核心要 64 次 FMA。Tensor Core 正是混合精度训练（float16 计算、float32 累加）快的原因。

• Transformer Engine（Hopper+）能在单层内动态切换 FP8 和 FP16 精度，只在需要的地方用更高精度。这在不牺牲模型质量的前提下最大化吞吐。它专为 transformer 架构（注意力 + MLP）设计，而 transformer 主导了现代 AI。

编程练习（用 nvcc 编译）

1. 写一个对数组做 ReLU 的 CUDA 内核。测量包含内存传输在内的时间。这会教你写内核、用 cudaMalloc/cudaMemcpy，以及 host↔device 传输瓶颈。

// task1_relu.cu
// 编译：nvcc -O3 -o task1_relu task1_relu.cu

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

__global__ void relu_kernel(const float* input, float* output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = input[idx] > 0.0f ? input[idx] : 0.0f;
    }
}

int main() {
    const int N = 1 << 24;  // 约 1600 万
    size_t bytes = N * sizeof(float);

    // 分配 host 内存
    float* h_input  = (float*)malloc(bytes);
    float* h_output = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) {
        h_input[i] = (float)(i % 100) - 50.0f;  // 正负混合
    }

    // 分配 device 内存
    float *d_input, *d_output;
    cudaMalloc(&d_input, bytes);
    cudaMalloc(&d_output, bytes);

    // 对完整流水线计时：复制到 GPU、计算、复制回来
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);
    cudaMemcpy(d_input, h_input, bytes, cudaMemcpyHostToDevice);

    int block_size = 256;
    int grid_size = (N + block_size - 1) / block_size;
    relu_kernel<<<grid_size, block_size>>>(d_input, d_output, N);

    cudaMemcpy(h_output, d_output, bytes, cudaMemcpyDeviceToHost);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);

    // 验证
    int errors = 0;
    for (int i = 0; i < N; i++) {
        float expected = h_input[i] > 0.0f ? h_input[i] : 0.0f;
        if (h_output[i] != expected) errors++;
    }

    printf("Time (including transfers): %.2f ms\n", ms);
    printf("Bandwidth: %.1f GB/s\n", 2.0 * bytes / ms / 1e6);  // 读 + 写
    printf("Errors: %d / %d\n", errors, N);

    cudaFree(d_input); cudaFree(d_output);
    free(h_input); free(h_output);
    return 0;
}
​

2. 在 CUDA 里用共享内存写一个分块的矩阵乘。与一个朴素（未分块）版本做性能对比。这会教你共享内存、__syncthreads，以及为什么分块重要。

// task2_matmul.cu
// 编译：nvcc -O3 -o task2_matmul task2_matmul.cu

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

#define TILE 16

// 朴素 matmul：每个线程算 C 的一个元素
__global__ void matmul_naive(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

// 分块 matmul：用共享内存减少全局内存访问
__global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float sA[TILE][TILE];
    __shared__ float sB[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        sA[threadIdx.y][threadIdx.x] = (row < N && t*TILE+threadIdx.x < N)
            ? A[row * N + t*TILE + threadIdx.x] : 0.0f;
        sB[threadIdx.y][threadIdx.x] = (t*TILE+threadIdx.y < N && col < N)
            ? B[(t*TILE + threadIdx.y) * N + col] : 0.0f;

        __syncthreads();
        for (int k = 0; k < TILE; k++)
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        __syncthreads();
    }

    if (row < N && col < N)
        C[row * N + col] = sum;
}

int main() {
    const int N = 1024;
    size_t bytes = N * N * sizeof(float);

    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes); cudaMalloc(&d_B, bytes); cudaMalloc(&d_C, bytes);

    // 用 1 初始化（便于验证：C 应该全是 N）
    float* h_A = new float[N*N];
    for (int i = 0; i < N*N; i++) h_A[i] = 1.0f;
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_A, bytes, cudaMemcpyHostToDevice);

    dim3 block(TILE, TILE);
    dim3 grid((N+TILE-1)/TILE, (N+TILE-1)/TILE);

    // 基准测试 朴素
    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);

    cudaEventRecord(start);
    for (int i = 0; i < 10; i++)
        matmul_naive<<<grid, block>>>(d_A, d_B, d_C, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float naive_ms; cudaEventElapsedTime(&naive_ms, start, stop);

    // 基准测试 分块
    cudaEventRecord(start);
    for (int i = 0; i < 10; i++)
        matmul_tiled<<<grid, block>>>(d_A, d_B, d_C, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float tiled_ms; cudaEventElapsedTime(&tiled_ms, start, stop);

    double gflops_naive = 2.0 * N * N * N * 10 / naive_ms / 1e6;
    double gflops_tiled = 2.0 * N * N * N * 10 / tiled_ms / 1e6;

    printf("Naive:  %.2f ms, %.1f GFLOPS\n", naive_ms/10, gflops_naive);
    printf("Tiled:  %.2f ms, %.1f GFLOPS\n", tiled_ms/10, gflops_tiled);
    printf("Speedup: %.1fx\n", naive_ms / tiled_ms);

    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    delete[] h_A;
    return 0;
}
​

3. 演示 warp 分歧。写一个内核，让同一 warp 的线程走不同分支，再与一个无分支版本做对比。

// task3_divergence.cu
// 编译：nvcc -O3 -o task3_diverge task3_divergence.cu

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

// 坏：warp 分歧 —— 偶数/奇数线程走不同路径
__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (idx % 2 == 0) {
            data[idx] = data[idx] * 2.0f + 1.0f;
        } else {
            data[idx] = data[idx] * 0.5f - 1.0f;
        }
    }
}

// 好：无分支 —— 所有线程执行同一条指令
__global__ void branchless_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float scale = (idx % 2 == 0) ? 2.0f : 0.5f;
        float offset = (idx % 2 == 0) ? 1.0f : -1.0f;
        data[idx] = data[idx] * scale + offset;
    }
}

int main() {
    const int N = 1 << 24;
    float* d_data;
    cudaMalloc(&d_data, N * sizeof(float));
    cudaMemset(d_data, 0, N * sizeof(float));

    int block = 256, grid = (N + block - 1) / block;

    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);

    // 分歧版
    cudaEventRecord(start);
    for (int i = 0; i < 100; i++)
        divergent_kernel<<<grid, block>>>(d_data, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float div_ms; cudaEventElapsedTime(&div_ms, start, stop);

    // 无分支版
    cudaEventRecord(start);
    for (int i = 0; i < 100; i++)
        branchless_kernel<<<grid, block>>>(d_data, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float nodiv_ms; cudaEventElapsedTime(&nodiv_ms, start, stop);

    printf("Divergent:  %.2f ms\n", div_ms / 100);
    printf("Branchless: %.2f ms\n", nodiv_ms / 100);
    printf("Speedup:    %.2fx\n", div_ms / nodiv_ms);

    cudaFree(d_data);
    return 0;
}
​