GPU架构与CUDA

GPU通过提供数千个核心来实现大规模并行，从而变革了AI。本文件涵盖GPU vs CPU的设计哲学、GPU内存层次、CUDA C++编程、SIMT执行模型、内存访问模式、同步、流、性能分析和NVIDIA GPU世代——编写和理解GPU内核所需的知识。

• 如需带有完整示例的CUDA实践教程，参见配套仓库：github.com/HenryNdubuaku/cuda-tutorials。

• 现代NVIDIA GPU有超过10,000个CUDA核心。CPU有4-128个核心。这100-1000倍的核心优势是GPU在ML中占主导地位的原因：训练Transformer需要数万亿次乘加运算，GPU以CPU无法匹敌的规模并行处理它们。

• 即使你从不亲自写CUDA内核，理解GPU架构也能解释：为什么批次大小重要（需要足够的工作来饱和GPU）、为什么内存通常是瓶颈（而非计算），以及为什么某些操作（散列、条件分支）在GPU上很慢。

GPU vs CPU：根本不同的设计

• CPU为延迟设计：最小化完成一个任务的时间。它将大部分晶体管预算用于缓存、分支预测器和乱序执行——所有使单线程快的技巧。

• GPU为吞吐量设计：最大化每秒完成的任务数。它将大部分晶体管用于执行单元（ALU）。单个线程慢，但有数千个。

	CPU	GPU
核心	4-128（复杂、快）	1,000-20,000（简单、慢）
时钟速度	3-5 GHz	1-2.5 GHz
缓存	大（32 MB+ L3）	小（每SM共享内存）
分支预测	复杂精密	无（所有线程沿同一路径）
最佳用途	低延迟、复杂控制流	高吞吐量、数据并行工作
典型FLOPS（FP32）	1-5 TFLOPS	30-80 TFLOPS
内存带宽	50-100 GB/s	1-3 TB/s

• GPU的内存带宽优势（10-30倍）往往比其计算优势更重要。许多ML操作是内存受限的（逐元素操作、归一化、注意力），GPU的带宽使其能以足够的速度向核心馈送数据。

GPU内存层次

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

内存	大小	延迟	带宽	作用域
寄存器	每SM ~256 KB	0周期	最高	每线程
共享内存	每SM 48-228 KB	~5周期	~20 TB/s	每线程块
L1缓存	每SM 128-256 KB	~30周期		每SM
L2缓存	4-96 MB	~200周期	~6 TB/s	全局
全局内存（HBM）	24-192 GB	~400周期	1-3.3 TB/s	全局

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

• 共享内存是程序员管理的缓存，block中所有线程共享。它是编写快速CUDA内核的关键：从慢速全局内存加载一块数据到快速共享内存，然后在其上计算。这是主导GPU编程的分块模式。

• 全局内存（HBM）：主GPU内存（VRAM）。大但慢（400周期延迟）。所有数据从这里开始和结束。内核优化的目标是最小化全局内存访问。

CUDA编程模型

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

层次：Grid、Block、Thread

Grid（整个发射）
├── Block (0,0)
│   ├── Thread (0,0)
│   ├── Thread (1,0)
│   ├── Thread (2,0)
│   └── ...（最多每block 1024线程）
├── Block (1,0)
│   ├── Thread (0,0)
│   └── ...
└── ...（数百万block可能）

• Thread：最小单元。每个线程在其block内有唯一ID（threadIdx.x）。
• Block：可以共享内存和同步的线程组。Block ID：blockIdx.x。Block大小：blockDim.x（最多1024线程）。

• Grid：单次内核发射启动的所有block。可以是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) {           // 边界检查（grid可能大于数据量）
        c[idx] = a[idx] + b[idx];
    }
}

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

    // 分配主机（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;
    }

    // 分配设备（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);

    // 发射内核：每block 256线程，足够的block覆盖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 (期望 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>>>：内核发射语法。指定使用多少block和线程。
  • cudaMalloc / cudaFree：分配/释放GPU内存（类似new/delete但用于GPU）。
  • cudaMemcpy：在CPU和GPU之间复制数据。这通常是最大的瓶颈（PCIe带宽约32 GB/s，而GPU内存带宽约3 TB/s）。

Warp和SIMT

• GPU以32个线程为一组执行，称为warp。同一warp中所有32个线程同时执行相同指令（SIMT——单指令多线程）。这是GPU的SIMD等价物，但是在线程层面。

• Warp发散发生在同一warp的线程走不同分支时。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];  // 所有线程执行相同指令

内存合并

• 合并访问：当连续线程访问连续内存地址时，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字节但只用一小部分。stride为32是最坏情况：每次事务加载128字节但只有一个线程使用4字节（3%利用率）。

共享内存和分块

• 分块模式是最重要的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__：声明block中所有线程共享的内存（快速、片上）。
• __syncthreads()：一个屏障，等待block中所有线程到达此点。在写入共享内存和从共享内存读取之间必须使用（否则一些线程读到过时数据）。
• 为什么分块有效：没有它，每个线程每次乘法都从全局内存加载。有了分块，TILE_SIZE × TILE_SIZE的数据块加载一次到共享内存，被block中所有线程重用。重用因子为TILE_SIZE，将全局内存流量降低该因子。

流和并发

• 默认情况下，CUDA操作是顺序的：CPU发射一个内核，等待完成，然后发射下一个。流实现重叠：

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

• 要看什么：
  • 占用率：SM被利用的分数。低占用率（< 50%）意味着线程太少无法隐藏内存延迟。原因：每线程太多寄存器、每block太多共享内存。
  • 内存吞吐量：与峰值带宽比较。如果达到峰值< 50%，内存访问模式效率低（非合并访问、bank冲突）。
  • 计算吞吐量：与峰值FLOPS比较。如果内存和计算吞吐量都低，内核是延迟受限的（并行度不够）。

高级优化技术

• 除了合并访问和共享内存分块的基础之外，高性能GPU（和CPU）代码使用几种高级技术：

数据布局：AoS vs SoA

• 结构数组（AoS）：每个元素把其所有字段存在一起。[{x,y,z}, {x,y,z}, {x,y,z}]。
• 数组结构（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

• SoA对于数据并行工作负载（SIMD、GPU）几乎总是更快。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);
    // ...
}

• 预取指令是一个提示：如果数据已在缓存中，是no-op。如果不在，CPU在执行其他指令时开始在后台获取它。预取距离（此例中提前64个元素）应根据内存延迟和循环迭代时间调优。

内核融合

• 内核融合将多个操作组合到单个内核中，避免将中间结果写入内存。这是对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

• 对于内存受限操作（偏置加法、ReLU、LayerNorm），内存流量主导执行时间。融合完全消除流量。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小2倍，所以内存带宽翻倍（常见的瓶颈），在缓存中可容纳2倍更多数据。Tensor Cores以CUDA Cores的8-16倍速率处理FP16。这就是混合精度训练（第6章）以最小精度损失提供2-3倍加速的原因。

内存池分配器

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

• 内存池（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是峰值（池可能持有比当前使用更多的内存）。

基于性能分析的优化

• 不要盲目优化。先性能分析，找到瓶颈，优化瓶颈，然后再次分析。Roofline模型（文件01）告诉你瓶颈是内存还是计算：

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

• 大多数ML工作负载是内存受限的。令人惊讶的含义：更快的GPU（更多FLOPS）通常无济于事。更快的内存（HBM3 vs HBM2e）帮助更大。这就是为什么A100→H100升级不只是关于FLOPS——H100还有2倍的内存带宽。

NVIDIA GPU世代

世代	年份	关键创新	AI相关性
Pascal (P100)	2016	HBM2, NVLink	第一代严肃的深度学习GPU
Volta (V100)	2017	Tensor Cores（混合精度矩阵乘法）	启用FP16训练，125 TFLOPS TF32
Ampere (A100)	2020	TF32, 稀疏性, 第3代Tensor Cores	312 TFLOPS TF32，结构化稀疏2:4
Hopper (H100)	2022	Transformer Engine（FP8）, HBM3	989 TFLOPS FP8，动态精度切换
Blackwell (B200)	2024	第2代Transformer Engine, NVLink 5	2.5 PFLOPS FP4，多die设计

• Tensor Cores是专用矩阵乘法单元。一条Tensor Core指令在一个周期内计算一个4×4矩阵乘法（D = A×B + C）。常规CUDA Cores需要64次FMA操作。Tensor Cores是混合精度训练（float16计算，float32累积）快速的原因。

• Transformer Engine（Hopper+）在单层内动态切换FP8和FP16精度，只在需要处选择更高精度。这在最大化吞吐量的同时不牺牲模型质量。它专门为现代AI的主导架构Transformer（注意力+MLP）设计。

编程任务（用nvcc编译）

1. 编写一个CUDA内核，对数组应用ReLU。测量包含内存传输的时间。这教授内核编写、cudaMalloc/cudaMemcpy和主机↔设备传输瓶颈。

   // 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;  // ~16M元素
       size_t bytes = N * sizeof(float);
   
       // 分配主机内存
       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;  // 正负混合
       }
   
       // 分配设备内存
       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("时间（含传输）: %.2f ms\n", ms);
       printf("带宽: %.1f GB/s\n", 2.0 * bytes / ms / 1e6);  // 读 + 写
       printf("错误: %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
   
   // 朴素矩阵乘法：每个线程计算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;
       }
   }
   
   // 分块矩阵乘法：使用共享内存减少全局内存访问
   __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("朴素:  %.2f ms, %.1f GFLOPS\n", naive_ms/10, gflops_naive);
       printf("分块:  %.2f ms, %.1f GFLOPS\n", tiled_ms/10, gflops_tiled);
       printf("加速比: %.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("发散:  %.2f ms\n", div_ms / 100);
       printf("无分支: %.2f ms\n", nodiv_ms / 100);
       printf("加速比:    %.2fx\n", div_ms / nodiv_ms);
   
       cudaFree(d_data);
       return 0;
   }