第4节：CUDA实战——并行向量加法(从CPU到GPU的性能飞跃）

引言

纸上得来终觉浅，绝知此事要躬行

前几节我们学习了GPU的硬件架构、内存体系和编程模型。现在，是时候亲手写一个完整的CUDA程序，并亲眼见证GPU带来的性能飞跃了。

向量加法（Vector Addition）是并行计算的“Hello World”，它足够简单，却能揭示许多关键的性能原理：线程配置的影响、内存访问模式、带宽利用率、延迟隐藏等。

今天，我们将：

• 实现CPU版本的向量加法作为基线
• 实现GPU版本并逐步优化
• 使用CUDA事件精确测量性能
• 对比不同数据规模下的速度差异
• 分析为什么GPU这么快，以及什么情况下可能不如CPU

一、问题定义

给定两个长度为 N 的浮点数组 A 和 B，计算 C = A + B，即：

for i in 0..N-1: C[i] = A[i] + B[i]

这是一个数据并行任务——每个元素的计算相互独立，非常适合GPU。

二、CPU基线版本

2.1 实现代码

// vector_add_cpu.cpp#include<stdio.h>#include<stdlib.h>#include<time.h>voidvector_add_cpu(constfloat*a,constfloat*b,float*c,intn){for(inti=0;i<n;i++){c[i]=a[i]+b[i];}}doubleget_time(){structtimespects;clock_gettime(CLOCK_MONOTONIC,&ts);returnts.tv_sec+ts.tv_nsec*1e-9;}intmain(){intn=1<<20;// 1,048,576 个元素size_t bytes=n*sizeof(float);// 分配主机内存float*h_a=(float*)malloc(bytes);float*h_b=(float*)malloc(bytes);float*h_c=(float*)malloc(bytes);// 初始化数据for(inti=0;i<n;i++){h_a[i]=i*1.0f;h_b[i]=i*2.0f;}// 计时doublestart=get_time();vector_add_cpu(h_a,h_b,h_c,n);doubleend=get_time();doubleelapsed=end-start;doublebandwidth=(bytes*2)/elapsed/1e9;// 读取A+B，写入C，共3次内存访问？实际上是读A和B各一次，写C一次，所以总数据量 = 3 * bytes// 但通常向量加法需要读A和B，写C，共3*bytesdoublegflops=n/elapsed/1e9;// 每秒十亿次浮点运算printf("CPU 向量加法 (n = %d):\n",n);printf(" 时间: %.4f s\n",elapsed);printf(" 带宽: %.2f GB/s\n",bandwidth);printf(" 性能: %.2f GFLOP/s\n",gflops);// 验证结果for(inti=0;i<10;i++){printf("c[%d] = %f (应得 %f)\n",i,h_c[i],h_a[i]+h_b[i]);}free(h_a);free(h_b);free(h_c);return0;}

编译运行：

gcc-O3vector_add_cpu.cpp-ovector_add_cpu-lrt./vector_add_cpu

2.2 CPU性能预期

在一颗现代CPU（如AMD Ryzen 9）上，对于1M元素（约4MB数据），预期结果：

• 时间：约0.002-0.005秒
• 带宽：10-20 GB/s
• GFLOPs：0.2-0.5 GFLOP/s

这个性能已经不错，但我们将看到GPU的恐怖之处。

三、GPU基础版本

3.1 核函数实现

// vector_add_gpu_baseline.cu#include<cuda_runtime.h>#include<stdio.h>#include<stdlib.h>#defineCHECK_CUDA(call){\cudaError_t err=call;\if(err!=cudaSuccess){\printf("CUDA错误 at %s:%d - %s\n",__FILE__,__LINE__,\cudaGetErrorString(err));\exit(1);\}\}__global__voidvector_add(float*a,float*b,float*c,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;if(tid<n){c[tid]=a[tid]+b[tid];}}intmain(){intn=1<<20;// 1M 元素size_t bytes=n*sizeof(float);// 主机内存float*h_a,*h_b,*h_c;h_a=(float*)malloc(bytes);h_b=(float*)malloc(bytes);h_c=(float*)malloc(bytes);for(inti=0;i<n;i++){h_a[i]=i*1.0f;h_b[i]=i*2.0f;}// 设备内存float*d_a,*d_b,*d_c;CHECK_CUDA(cudaMalloc(&d_a,bytes));CHECK_CUDA(cudaMalloc(&d_b,bytes));CHECK_CUDA(cudaMalloc(&d_c,bytes));CHECK_CUDA(cudaMemcpy(d_a,h_a,bytes,cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(d_b,h_b,bytes,cudaMemcpyHostToDevice));// 配置线程intthreads_per_block=256;intblocks_per_grid=(n+threads_per_block-1)/threads_per_block;// 创建CUDA事件用于计时cudaEvent_t start,stop;CHECK_CUDA(cudaEventCreate(&start));CHECK_CUDA(cudaEventCreate(&stop));CHECK_CUDA(cudaEventRecord(start));vector_add<<<blocks_per_grid,threads_per_block>>>(d_a,d_b,d_c,n);CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));floatms;CHECK_CUDA(cudaEventElapsedTime(&ms,start,stop));CHECK_CUDA(cudaMemcpy(h_c,d_c,bytes,cudaMemcpyDeviceToHost));doubleelapsed=ms/1000.0;// 转换为秒doublebandwidth=(bytes*3)/elapsed/1e9;// 读A+B，写C，共3*bytesdoublegflops=n/elapsed/1e9;printf("GPU 向量加法 (n = %d, block = %d):\n",n,threads_per_block);printf(" 时间: %.6f s (%.3f ms)\n",elapsed,ms);printf(" 带宽: %.2f GB/s\n",bandwidth);printf(" 性能: %.2f GFLOP/s\n",gflops);// 验证for(inti=0;i<10;i++){printf("c[%d] = %f (应得 %f)\n",i,h_c[i],h_a[i]+h_b[i]);}CHECK_CUDA(cudaFree(d_a));CHECK_CUDA(cudaFree(d_b));CHECK_CUDA(cudaFree(d_c));free(h_a);free(h_b);free(h_c);CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));return0;}

编译运行：

nvcc-O3vector_add_gpu_baseline.cu-ovector_add_gpu ./vector_add_gpu

3.2 GPU性能预期

在A100上，预期：

• 时间：约0.0001秒（0.1ms）
• 带宽：约1400 GB/s（接近A100的峰值1555 GB/s）
• GFLOPs：约10 GFLOP/s

速度对比：GPU比CPU快20-50倍！

四、性能对比分析

4.1 不同数据规模下的对比

我们编写一个测试脚本，对不同的N（从1K到100M）分别运行CPU和GPU版本，记录时间。

重要发现：

• 当数据量很小时，CPU更快（因为GPU有启动开销和数据传输开销）
• 随着数据量增大，GPU优势越来越明显，最终达到10倍以上
• 对于极大规模数据，GPU带宽成为瓶颈，加速比趋于稳定

4.2 为什么GPU在小数据量时慢？

• kernel启动开销：约几微秒到几十微秒
• 数据拷贝开销：即使数据量小，PCIe传输也有固定延迟
• GPU需要预热：首次调用会有额外开销

结论：GPU适合大规模并行计算，小任务交给CPU更合适。

4.3 带宽利用率分析

理论带宽 vs 实际带宽：

• A100理论显存带宽：1555 GB/s
• 我们的向量加法实测：~1400 GB/s（利用率90%）

为什么达不到100%？

• 指令开销
• 内存访问模式（虽然这里是合并访问）
• 内存控制器效率

五、探究线程配置对性能的影响

5.1 改变block大小

固定N=1M，改变threads_per_block，观察性能变化：

分析：

• block太小 → SM中活跃warp少，无法隐藏访存延迟
• block太大 → 每个线程可用寄存器减少，可能溢出到本地内存
• 256是很多简单kernel的甜点值

5.2 改变grid大小（block数量）

理论上，block数量只要足够多，就能充分利用所有SM。但太多block会导致调度开销。

建议：block数至少是SM数量的几倍（如4-8倍），以保证负载均衡。

六、深入剖析：使用NVIDIA Nsight Compute

要真正理解性能，必须使用性能分析工具。

6.1 安装Nsight Compute

# Linuxsudoaptinstallnvidia-nsight-compute

6.2 运行性能分析

ncu--metricssm__throughput.avg.pct_of_peak_sustained_elapsed ./vector_add_gpu

会输出各种指标，如：

• 计算吞吐量
• 内存吞吐量
• 占用率
• stalls原因

6.3 关键指标解读

内存带宽：是否接近峰值？如果不是，说明存在访存问题。

占用率：活跃warp与最大warp的比值。理想情况下应大于50%。

stall原因：分析kernel卡在什么地方（等待内存、计算依赖等）。

七、进阶优化：使用共享内存

虽然向量加法天然就是元素独立，但我们可以尝试用共享内存来合并全局内存访问？实际上，向量加法不需要共享内存，因为每个线程只访问自己的元素，没有数据重用。

但我们可以通过向量化类型（如float2、float4）来增加每个线程的工作量，减少指令数，提高带宽利用率。

7.1 使用float4优化

__global__voidvector_add_float4(float*a,float*b,float*c,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;// 每个线程处理4个元素intidx=tid*4;if(idx+3<n){float4 a4=reinterpret_cast<float4*>(a)[tid];float4 b4=reinterpret_cast<float4*>(b)[tid];float4 c4;c4.x=a4.x+b4.x;c4.y=a4.y+b4.y;c4.z=a4.z+b4.z;c4.w=a4.w+b4.w;reinterpret_cast<float4*>(c)[tid]=c4;}else{// 处理剩余元素for(inti=idx;i<n;i++){c[i]=a[i]+b[i];}}}

这样每个线程做4次加法，但只发一次128位的内存请求（合并更好），可以减少指令数，提高带宽利用率。

实测性能可提升10-20%。

八、完整性能对比代码框架

为了方便你亲自实验，这里提供一个完整的测试框架，可以自动测试不同规模和不同block大小：

// vector_add_benchmark.cu#include<cuda_runtime.h>#include<stdio.h>#include<stdlib.h>#include<vector>#include<algorithm>#defineCHECK_CUDA(call){...}// 同上__global__voidvector_add(float*a,float*b,float*c,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;if(tid<n)c[tid]=a[tid]+b[tid];}doubletest_gpu(intn,intthreads_per_block){size_t bytes=n*sizeof(float);float*h_a,*h_b,*h_c;h_a=(float*)malloc(bytes);h_b=(float*)malloc(bytes);h_c=(float*)malloc(bytes);for(inti=0;i<n;i++){h_a[i]=i*1.0f;h_b[i]=i*2.0f;}float*d_a,*d_b,*d_c;CHECK_CUDA(cudaMalloc(&d_a,bytes));CHECK_CUDA(cudaMalloc(&d_b,bytes));CHECK_CUDA(cudaMalloc(&d_c,bytes));CHECK_CUDA(cudaMemcpy(d_a,h_a,bytes,cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(d_b,h_b,bytes,cudaMemcpyHostToDevice));intblocks=(n+threads_per_block-1)/threads_per_block;cudaEvent_t start,stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start);vector_add<<<blocks,threads_per_block>>>(d_a,d_b,d_c,n);cudaEventRecord(stop);cudaEventSynchronize(stop);floatms;cudaEventElapsedTime(&ms,start,stop);cudaMemcpy(h_c,d_c,bytes,cudaMemcpyDeviceToHost);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);free(h_a);free(h_b);free(h_c);cudaEventDestroy(start);cudaEventDestroy(stop);returnms;}intmain(){std::vector<int>sizes={1000,10000,100000,1000000,10000000,100000000};std::vector<int>block_sizes={64,128,256,512,1024};printf("N, BlockSize, Time(ms), Bandwidth(GB/s)\n");for(intn:sizes){for(intbs:block_sizes){doublems=test_gpu(n,bs);doublebytes=n*sizeof(float)*3.0;// 读+写doublebw=bytes/ms/1e6;// GB/sprintf("%d, %d, %.3f, %.2f\n",n,bs,ms,bw);}}return0;}

运行并将输出重定向到CSV，然后用Python/matplotlib绘制图表。

九、常见问题与调试

9.1 程序没输出或崩溃

• 检查CUDA错误检查宏
• 确认GPU内存足够
• 检查边界条件（if (tid < n)）

9.2 性能远低于预期

• 是否用了Release模式编译（-O3）
• 是否在调试模式下运行（-G会极大降低性能）
• 检查block大小是否合理
• 检查是否开启了ECC（可能降低带宽）
• 用nvidia-smi查看GPU是否处于P0状态（最高性能）

9.3 数据传输瓶颈

• 用cudaMemcpyAsync和流来重叠数据传输和计算（后续章节）

十、本节总结

10.1 关键收获

1. 向量加法是最简单的并行任务，但能揭示GPU性能的基本规律
2. GPU在大规模数据下优势巨大，但小数据量不如CPU
3. 线程配置（block大小）影响性能，需要实验找到最优值
4. 带宽利用率是衡量内存密集型kernel的关键指标
5. 使用性能分析工具才能真正理解瓶颈

10.2 性能优化 checklist

• 确保合并访问
• 选择合适的block大小（通常128-512）
• 隐藏数据传输（异步拷贝）
• 考虑向量化加载（float4）
• 避免寄存器溢出
• 使用快速数学函数（如果精度允许）

10.3 下节预告

下一节我们将挑战更复杂的矩阵乘法，它将暴露更多的性能问题：访存与计算的权衡、共享内存的使用、分块算法等。准备好迎接真正的挑战吧！

十一、面试真题（2024-2026）

Q1：为什么向量加法在GPU上能获得很高的带宽利用率？

考察点：对合并访问的理解

参考答案：
向量加法的内存访问模式是完美的合并访问：线程0读A[0]，线程1读A[1]，…，线程31读A[31]，这些地址连续且对齐。硬件可以将这32次4字节访问合并成一次128字节的事务，最大化利用显存带宽。同时每个线程的计算非常简单，不会成为瓶颈，因此性能主要受限于内存带宽，容易达到接近理论峰值的水平。

Q2：在测试GPU性能时，你发现对于很小的N，GPU比CPU慢，为什么？

考察点：对GPU开销的理解

参考答案：
GPU有固定的启动开销：

• kernel启动延迟（约几微秒到几十微秒）
• 数据从主机到设备的传输延迟（PCIe延迟）
• 设备可能处于低功耗状态，需要唤醒
  当数据量很小时，这些开销占主导，而CPU可以立即执行，所以总时间更短。因此，GPU适合大规模并行计算，小任务应留在CPU。

Q3：如何测量kernel的实际内存带宽？计算公式是什么？

考察点：性能测量方法

参考答案：
使用CUDA事件记录kernel执行时间，然后计算总数据访问量除以时间。对于向量加法，每个元素需要读A、读B、写C，所以总数据量 = 3 * N * sizeof(float)。带宽 = 总数据量 / 时间。

注意：这测量的是应用层带宽，实际硬件可能还有缓存效应，但作为近似足够。

Q4：你提到了使用float4优化，为什么能提升性能？

考察点：对指令级优化的理解

参考答案：
使用float4可以让每个线程一次处理4个元素，带来几个好处：

1. 减少指令数：4次加法可以用向量化指令，减少取指和调度开销
2. 提高内存合并效率：一次128位加载比4次32位加载更高效，减少内存请求次数
3. 增加计算密度：每个线程做更多工作，有助于隐藏延迟
   但要注意边界处理和剩余元素的处理。

思考题：
在你的GPU上运行性能对比代码，画出不同N下CPU和GPU的时间曲线。找到你机器上GPU开始超越CPU的“交叉点”。对于这个交叉点的大小，你有什么发现？为什么不同GPU/CPU的交叉点不同？欢迎在评论区分享你的实验数据。