2.1 数据并行#

也就是说，在一个百万元素的张量上做 ReLU，第0个元素和第100个元素无任何关系。这种无跨线程数据依赖的特性，意味着它们具备优秀的数据并行性，是 GPU 优化中最容易达到理论峰值带宽的一类算子。

对GPU而言，这就意味着：

2.2 极低的计算访存比#

这是 Element-wise 算子最重要的性能特征，也决定了它所有的优化策略。

在 GPU 上执行一条算术指令（如浮点加法、乘法）的耗时，往往只有几个时钟周期；而访问一次全局显存的耗时，则是几百个时钟周期。

我们来看一个典型的逐元素加法操作：C[i] = A[i] + B[i]。

• 计算量：1 次加法。
• 访存量：读取 A、读取 B、写入 C。共计 3 次显存操作（假设单精度浮点，每次 4 字节）。 此时，计算访存比 = 1 FLOP / 12 Bytes ≈ 0.08 FLOP/Byte。 而主流 GPU 的理论峰值计算访存比通常在 10~20 FLOP/Byte 以上。

[!WARNING] 逐元素算子的瓶颈不在计算能力上，而在显存带宽上。 因此，衡量一个 Element-wise 核函数好坏的指标不是 TFLOPs，而是它是否跑满了显存带宽。

3.1. 访存合并与对齐#

GPU 的显存系统并非以单个字节或者单个浮点数为单位与核函数交互，而是以内存事务（Memory Transaction） 为单位。

以常见的 NVIDIA GPU 架构为例，L2 缓存与显存控制器之间的最小传输单元通常为 32 字节。当 SM发出一次全局内存访问请求时，硬件会尝试将同一 Warp 的所有请求合并成尽可能少的事务。

• 理想情况：一个Warp的32个线程访问连续且对齐的32个float，也就是128字节，那么只需要发送4次32字节的事务，或者1次128字节的事务。
• 最坏情况：32个线程访问的内存地址在显存各处，并不连续，那么一次Warp的访存操作最差为32次独立的事务。

由于单次事务的延迟高达数百个时钟周期，拆分事务意味着带宽利用率断崖式下跌，核函数则需要等待数据。

3.2. 隐藏访存延迟#

既然运算单元在等数据，那就让每一个线程多负责几个元素的处理。

这也就是非常著名的网格跨步循环(Grid-Stride Loop)

1
__global__ void add_kernel_v2(const float* A, const float* B, float* C, int N) {
2
    // 起始索引
3
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
4
    // 总线程数（网格跨度）
5
    int stride = gridDim.x * blockDim.x;
6

7
    // 跨步循环处理多个元素
8
    for (int i = idx; i < N; i += stride) {
9
        C[i] = A[i] + B[i];
10
    }
11
}

启动配置并不需要覆盖整个N，而是可以让gridSize更小，这样每个线程就可以处理更多的数据。

1
int threadsPerBlock = 256;
2
int blocksPerGrid = 128 * 2;   // 故意只启动 SM 数量 × 若干倍，而非 N/256
3
add_kernel_v2<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这种方案实际上已经成为了Element Wise核函数的事实标准，这样无论数据量多大，核函数都能够自适应，无需调整启动参数。

3.3. 向量化加载与存储#

即使使用了 Grid-Stride Loop，每个线程每次循环依然只处理 1 个 float。而 GPU 的显存总线位宽通常是 32 字节（或更高），这意味着硬件有能力在一次事务中搬运更多数据。

因此我们可以在一次事务中批量进行加载数据。 CUDA 提供了内置的向量类型：float2、float4、double2 等。一个 float4 变量包含 4 个连续的 float，共 16 字节。使用它可以将 4 次 32-bit 访存合并为 1 次 128-bit 访存。

那我们要注意的就是原本idx的计算方式，要进行换算。

1
__global__ void add_kernel_v3(const float* A, const float* B, float* C, int N) {
2
    // 向量化：一次处理 4 个元素
3
    const float4* A4 = reinterpret_cast<const float4*>(A);
4
    const float4* B4 = reinterpret_cast<const float4*>(B);
5
    float4* C4 = reinterpret_cast<float4*>(C);
6

7
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
8
    int stride = gridDim.x * blockDim.x;
9

10
    // 以 float4 为单位循环
11
    for (int i = idx; i < N / 4; i += stride) {
12
        float4 a = A4[i];
13
        float4 b = B4[i];
14
        float4 c;
15
        c.x = a.x + b.x;
16
        c.y = a.y + b.y;
17
        c.z = a.z + b.z;
18
        c.w = a.w + b.w;
19
        C4[i] = c;
20
    }
21

22
    // 处理尾部不足 4 的剩余元素（略）
23
}

那通过Grid-Stride Loop与向量化加载结合，这样就能够使得单个线程的计算密度足够高。