逐元素操作算子

Element-wise 算子优化

一、核心特性

1. 数据并行

对于 element-wise 算子，输出位置的值只取决于对应位置的输入，不涉及其他位置。因此线程之间不需要交换数据 / 同步指令，也没有额外的寻址开销，具备优秀的数据并行性。这是 GPU 优化中最容易打到理论峰值带宽的一类算子。

为什么向量加法天然适合 GPU：每个元素的「读 → 算 → 写」彼此独立，互不依赖，所以成千上万份操作可以同时重叠进行，把显存带宽和运算单元一起喂饱。

2. 极低的计算访存比（最重要的性能特征）

这是 element-wise 算子最关键的性能特征，决定了所有优化策略。

GPU 上执行一条算术指令只需要几个时钟周期，而访问全局显存的耗时却高达几百个时钟周期。两者差距悬殊。

以加法 $C[i] = A[i] + B[i]$ 为例：

• 计算量：只有 1 次加法
• 访存量：读 A、读 B、写 C，共 3 次显存操作 = 3 × 4 Byte = 12 Byte
• 计算访存比 = 1 FLOP / 12 Byte ≈ 0.083 FLOP/Byte

而主流 GPU 的理论峰值计算访存比在 10 ~ 20 FLOP/Byte 以上。

⚠️ 核心结论：逐元素算子的瓶颈不在计算能力，而在显存带宽。衡量一个 element-wise 核函数好坏的指标不是 TFLOPS，而是能否跑满带宽。

---

二、优化方向

模板代码（通用框架）

__global__ void add_kernel(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 threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;  // 向上取整
add_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

向上取整的写法：(N + threadsPerBlock - 1) / threadsPerBlock。因为 C++ 整数除法向下取整，若直接写 N / threadsPerBlock，当 N 不能整除时会少算一个块、漏掉尾部元素。先加上「几乎一整块」再除，保证只要有余数就进位。
例：N=1000，每块 256 → (1000+255)/256 = 4 块 → 1024 个线程，足够覆盖 1000 个元素（多出的 24 个线程由 if (idx < N) 拦下）。

---

优化点 1：访存合并与对齐（Coalesced Access）

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

以 NVIDIA 架构为例，L2 缓存与显存控制器之间的最小传输单元为 32 字节。当 SM 发出一次全局内存访问请求时，硬件会尝试把同一个 Warp（32 个线程）的请求合并为尽可能少的事务：

• 理想情况：一个 Warp 的 32 个线程访问连续且对齐的 32 个 float（128 字节），只需 4 次 32 字节事务（或一次 128 字节事务）。
• 最坏情况：32 个线程访问的地址散落在显存各处、不连续，最差需要 32 次独立事务。

由于单次事务延迟有几百个时钟周期，拆分事务意味着带宽利用率断崖式下跌，核函数被迫等待数据。

实践要点：让相邻线程访问相邻内存。后面的 Grid-Stride Loop 在循环的每一轮里，恰好保持了「线程 0、1、2、3 访问相邻位置」这个特性，所以天然是合并访存友好的。

---

优化点 2：隐藏访存延迟（Grid-Stride Loop）

既然运算单元在等数据，那就让每个线程多负责几个元素，这就是著名的网格跨步循环（Grid-Stride Loop）。

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

    for (int i = idx; i < N; i += stride) {               // 跨步循环
        C[i] = A[i] + B[i];
    }
}

启动配置不再需要覆盖整个 N，而是给一个更小、更固定的网格，让每个线程处理更多数据：

int threadsPerBlock = 256;
int blocksPerGrid = numSMs * 32;   // 按 SM 数量的若干倍，而非 N / 256
add_kernel_v2<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这种方案已成为 element-wise 核函数的实施标准：无论数据量多大，核函数都能自适应，无需调整启动参数。

它到底优化了什么——延迟隐藏（Latency Hiding）

这里有一个容易误解的点，需要厘清：

❌ 误解：「写循环，计算单元还是要等所有数据加载完毕。」
✅ 实际：计算单元只等下一份要算的数据，而非全部；而且连这点等待，GPU 也会切去跑别的线程填掉。

机制有两层：

1. 请求重叠在途：现代 GPU 允许一个线程在前一个数据还没回来时，就把后面几轮的读请求也发出去。于是「读 i」「读 i+stride」「读 i+2·stride」这些请求并排飞向显存，数据谁先到就先算谁，等待时间被重叠掉而非叠加。这正是「让每个线程多处理几个元素」的真正意义——持续发出更多在途请求，把显存的传输管道填满。

2. 线程间切换（延迟隐藏）：当一个线程因等待数据而停顿时，GPU 立刻切换到另一个数据已就绪的线程去执行。靠在大量线程间飞快切换，「等待」被其他线程的「干活」填满，运算单元几乎不空闲。

类比：厨师（运算单元）在等食材（数据）。若每个服务员（线程）只跑一趟腿就下班，厨师常陷入空等；若每个服务员一趟多带几份、并持续往返补货，传菜口始终堆着备好的食材，厨师几乎不用停。

---

优化点 3：向量化加载与存储（Vectorized Access）

即便使用 Grid-Stride Loop，每个线程每轮循环仍只处理 1 个 float。而 GPU 显存总线位宽通常为 32 字节甚至更高，意味着硬件一次事务能搬运更多数据。

CUDA 提供了内置向量类型：float2、float4、double2 等。一个 float4 包含 4 个 float（共 16 字节）。用它可以把 4 次 32-bit 访存合并为 1 次 128-bit 访存。

注意：原本 idx 的含义要换算——现在一个下标代表一整包（4 个元素）。

__global__ void add_kernel_v3(const float* A, const float* B, float* C, int N) {
    // 用 float4 的"眼镜"重新解读同一段内存（数据本身没动，只换了读法）
    const float4* A4 = reinterpret_cast<const float4*>(A);
    const float4* B4 = reinterpret_cast<const float4*>(B);
    float4*       C4 = reinterpret_cast<float4*>(C);

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

    int n4 = N / 4;                          // 完整的 float4 包数
    // 以 float4 为单位循环
    for (int i = idx; i < n4; i += stride) {
        float4 a = A4[i];                    // 一次读 4 个
        float4 b = B4[i];
        float4 c;
        c.x = a.x + b.x;
        c.y = a.y + b.y;
        c.z = a.z + b.z;
        c.w = a.w + b.w;
        C4[i] = c;                           // 一次写 4 个
    }

    // ---- 尾部处理：N 不是 4 的整数倍时，补算剩下不足 4 个的元素 ----
    // 例：N=10，上面只处理了前 8 个，这里补算第 8、9 个
    int tail_start = n4 * 4;
    for (int i = tail_start + idx; i < N; i += stride) {
        C[i] = A[i] + B[i];
    }
}

⚠️ 尾部处理不能省：N / 4 是整数除法。若 N 不是 4 的倍数，主循环会漏掉末尾不足 4 个的元素，必须用一段标量循环补算，否则结果会少算几个。

通过 Grid-Stride Loop + 向量化加载结合，单个线程的计算密度足够高，同时拥有前者的灵活性与后者的宽访存。

---

三、三个版本对比小结

版本	思路	线程数要求	主要优化
v1 naive	一线程一元素	必须 ≥ N	最易懂的并行写法
v2 grid-stride	一线程跨步处理多个元素	与 N 解耦，可固定	延迟隐藏 + 自适应 + 合并访存
v3 vectorized	在 v2 基础上用 float4 一次搬 4 个	与 N 解耦	进一步减少内存事务数、压满带宽

三者是层层叠加的关系，而非互相替代：v3 同时具备 grid-stride 的灵活性和 float4 的宽访存。优化主线始终围绕一个目标——因为是访存受限，所以一切都为了把显存带宽用满。

---

四、注意事项

• 实际运行时 NVCC 会对代码做编译优化，三版的性能差距可能不像理论上那么大——简单的 v1 模式编译器也能优化得不错。不同型号 GPU 的结果差异往往更明显。
• 想精确验证「带宽利用率」，可用 Nsight Compute 等工具实测每版的 achieved memory bandwidth，比单看耗时更有说服力。
• 衡量这类核函数，始终记住：看的是带宽利用率，不是 TFLOPS。

---