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

写在前面

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

CUDA 核函数编写（__global__）
线程层次结构（Grid / Block / Thread）
显存分配与传递

题目描述

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

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

输入输出

输入：设备指针 A, B 和数据规模 N
输出：设备指针 C

解法 v0：基础实现

1
2
3
4
5
__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];
}

代码解读

每个线程处理一个元素的加法运算
使用一维线程块布局：blockDim.x * blockIdx.x + threadIdx.x 计算全局索引
边界检查 if (idx < N) 防止越界访问

性能特点

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

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

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

1
2
3
4
5
6
__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];
}

核心思想

每个线程处理多个数据，步长 = 总线程数
任意启动配置都能处理全部数据
甚至可以退化为 <<<1, 1>>> 串行执行

优化2：向量化读取

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

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
__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指令变化
v0	268,435,456	100% (基准)
v1	139,657,216	减少48%
v2	48,758,784	减少82%

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

小结

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

索引计算 — 掌握一维线程索引 mapping
边界检查 — 防止越界访问
合并访问 — 连续线程访问连续内存
跨步循环 — 解耦网格配置与数据规模
向量化 — 减少循环开销