[向量加法] LeetGPU 第一题详解 - 基础 CUDA 核函数与性能分析

LeetGPU 题解

写在前面

本系列是 LeetGPU 平台的刷题记录，每道题我会给出从基础实现到优化的完整思路。如果你是 CUDA 新手，建议先掌握以下基础知识：

题目描述

给定两个浮点数组 A 和 B，将它们对应位置的元素相加，结果存入数组 C 中。数学表达式为：

\[ C_i = A_i + B_i \quad (i = 0, 1, 2, \dots, N-1) \]

输入输出

解法 v0：基础实现


__global__ void vector_add_v0(const float* A, const float* B, float* C, int N) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < N)
        C[idx] = A[idx] + B[idx];
}

代码解读

性能特点

优点缺点
实现简洁，易于理解网格配置与数据规模强绑定
天然并行，无数据依赖盲目启动海量线程增加调度开销
合并访存，连续线程访问连续显存

优化1：跨步循环（Grid-Stride Loop）

v0 的问题在于网格配置与数据规模强绑定。为了达到最佳 Occupancy，我们通常希望根据设备 SM 数量固定启动线程数。跨步循环完美解耦了两者：


__global__ void vector_add_v1(const float* A, const float* B, float* C, int N) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int step = blockDim.x * gridDim.x;  // 所有线程总数
    for (int i = idx; i < N; i += step)
        C[i] = A[i] + B[i];
}

核心思想

优化2：向量化读取

进一步减少循环开销，一次处理 4 个元素：


__global__ void vector_add_v2(const float* A, const float* B, float* C, int N) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int step = blockDim.x * gridDim.x;
    const float4* a4 = (const float4*)A;
    const float4* b4 = (const float4*)B;
    float4* c4 = (float4*)C;
    int N4 = N / 4;

    for (int i = idx; i < N4; i += step) {
        float4 at = a4[i];
        float4 bt = b4[i];
        c4[i] = make_float4(at.x + bt.x, at.y + bt.y, at.z + bt.z, at.w + bt.w);
    }

    // 处理尾部剩余元素
    int tail = 4 * N4;
    for (int i = idx + tail; i < N; i += step)
        C[i] = A[i] + B[i];
}

性能提升

v1 中处理 4 个元素需要 4 次循环、4 次边界判断、4 次步长累加。v2 只需 1 次循环，大幅减轻 ALU 和指令发射器的负担。

性能对比

N = (1 << 28) + 3

版本执行总指令数相对v0指令变化
v0268,435,456100% (基准)
v1139,657,216减少48%
v248,758,784减少82%

由于该核函数完全受访存带宽限制，总指令数降低对执行时间无明显提升。但对于计算密集型核函数，v2 预期能带来显著性能提升。

小结

本题是 GPU 并行计算的入门级题目，关键要点：