文章目录引言一、问题定义二、CPU基线版本2.1 实现代码2.2 CPU性能预期三、GPU基础版本3.1 核函数实现3.2 GPU性能预期四、性能对比分析4.1 不同数据规模下的对比4.2 为什么GPU在小数据量时慢4.3 带宽利用率分析五、探究线程配置对性能的影响5.1 改变block大小5.2 改变grid大小block数量六、深入剖析使用NVIDIA Nsight Compute6.1 安装Nsight Compute6.2 运行性能分析6.3 关键指标解读七、进阶优化使用共享内存7.1 使用float4优化八、完整性能对比代码框架九、常见问题与调试9.1 程序没输出或崩溃9.2 性能远低于预期9.3 数据传输瓶颈十、本节总结10.1 关键收获10.2 性能优化 checklist10.3 下节预告十一、面试真题2024-2026Q1为什么向量加法在GPU上能获得很高的带宽利用率Q2在测试GPU性能时你发现对于很小的NGPU比CPU慢为什么Q3如何测量kernel的实际内存带宽计算公式是什么Q4你提到了使用float4优化为什么能提升性能引言纸上得来终觉浅绝知此事要躬行前几节我们学习了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#includestdio.h#includestdlib.h#includetime.hvoidvector_add_cpu(constfloat*a,constfloat*b,float*c,intn){for(inti0;in;i){c[i]a[i]b[i];}}doubleget_time(){structtimespects;clock_gettime(CLOCK_MONOTONIC,ts);returnts.tv_sects.tv_nsec*1e-9;}intmain(){intn120;// 1,048,576 个元素size_t bytesn*sizeof(float);// 分配主机内存float*h_a(float*)malloc(bytes);float*h_b(float*)malloc(bytes);float*h_c(float*)malloc(bytes);// 初始化数据for(inti0;in;i){h_a[i]i*1.0f;h_b[i]i*2.0f;}// 计时doublestartget_time();vector_add_cpu(h_a,h_b,h_c,n);doubleendget_time();doubleelapsedend-start;doublebandwidth(bytes*2)/elapsed/1e9;// 读取AB写入C共3次内存访问实际上是读A和B各一次写C一次所以总数据量 3 * bytes// 但通常向量加法需要读A和B写C共3*bytesdoublegflopsn/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(inti0;i10;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_cpu2.2 CPU性能预期在一颗现代CPU如AMD Ryzen 9上对于1M元素约4MB数据预期结果时间约0.002-0.005秒带宽10-20 GB/sGFLOPs0.2-0.5 GFLOP/s这个性能已经不错但我们将看到GPU的恐怖之处。三、GPU基础版本3.1 核函数实现// vector_add_gpu_baseline.cu#includecuda_runtime.h#includestdio.h#includestdlib.h#defineCHECK_CUDA(call){\cudaError_t errcall;\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){inttidblockIdx.x*blockDim.xthreadIdx.x;if(tidn){c[tid]a[tid]b[tid];}}intmain(){intn120;// 1M 元素size_t bytesn*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(inti0;in;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_block256;intblocks_per_grid(nthreads_per_block-1)/threads_per_block;// 创建CUDA事件用于计时cudaEvent_t start,stop;CHECK_CUDA(cudaEventCreate(start));CHECK_CUDA(cudaEventCreate(stop));CHECK_CUDA(cudaEventRecord(start));vector_addblocks_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));doubleelapsedms/1000.0;// 转换为秒doublebandwidth(bytes*3)/elapsed/1e9;// 读AB写C共3*bytesdoublegflopsn/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(inti0;i10;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_gpu3.2 GPU性能预期在A100上预期时间约0.0001秒0.1ms带宽约1400 GB/s接近A100的峰值1555 GB/sGFLOPs约10 GFLOP/s速度对比GPU比CPU快20-50倍四、性能对比分析4.1 不同数据规模下的对比我们编写一个测试脚本对不同的N从1K到100M分别运行CPU和GPU版本记录时间。N元素个数数据量MBCPU时间msGPU时间ms加速比1,0000.0120.0020.050.0410,0000.120.0150.060.25100,0001.20.120.081.51,000,000121.20.15810,000,000120120.815100,000,0001200120815重要发现当数据量很小时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大小固定N1M改变threads_per_block观察性能变化block大小时间ms带宽GB/s说明320.22650占用率低640.188001280.169002560.15950最佳5120.155920略有下降10240.17840寄存器压力分析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-compute6.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){inttidblockIdx.x*blockDim.xthreadIdx.x;// 每个线程处理4个元素intidxtid*4;if(idx3n){float4 a4reinterpret_castfloat4*(a)[tid];float4 b4reinterpret_castfloat4*(b)[tid];float4 c4;c4.xa4.xb4.x;c4.ya4.yb4.y;c4.za4.zb4.z;c4.wa4.wb4.w;reinterpret_castfloat4*(c)[tid]c4;}else{// 处理剩余元素for(intiidx;in;i){c[i]a[i]b[i];}}}这样每个线程做4次加法但只发一次128位的内存请求合并更好可以减少指令数提高带宽利用率。实测性能可提升10-20%。八、完整性能对比代码框架为了方便你亲自实验这里提供一个完整的测试框架可以自动测试不同规模和不同block大小// vector_add_benchmark.cu#includecuda_runtime.h#includestdio.h#includestdlib.h#includevector#includealgorithm#defineCHECK_CUDA(call){...}// 同上__global__voidvector_add(float*a,float*b,float*c,intn){inttidblockIdx.x*blockDim.xthreadIdx.x;if(tidn)c[tid]a[tid]b[tid];}doubletest_gpu(intn,intthreads_per_block){size_t bytesn*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(inti0;in;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(nthreads_per_block-1)/threads_per_block;cudaEvent_t start,stop;cudaEventCreate(start);cudaEventCreate(stop);cudaEventRecord(start);vector_addblocks,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::vectorintsizes{1000,10000,100000,1000000,10000000,100000000};std::vectorintblock_sizes{64,128,256,512,1024};printf(N, BlockSize, Time(ms), Bandwidth(GB/s)\n);for(intn:sizes){for(intbs:block_sizes){doublemstest_gpu(n,bs);doublebytesn*sizeof(float)*3.0;// 读写doublebwbytes/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 关键收获向量加法是最简单的并行任务但能揭示GPU性能的基本规律GPU在大规模数据下优势巨大但小数据量不如CPU线程配置block大小影响性能需要实验找到最优值带宽利用率是衡量内存密集型kernel的关键指标使用性能分析工具才能真正理解瓶颈10.2 性能优化 checklist确保合并访问选择合适的block大小通常128-512隐藏数据传输异步拷贝考虑向量化加载float4避免寄存器溢出使用快速数学函数如果精度允许10.3 下节预告下一节我们将挑战更复杂的矩阵乘法它将暴露更多的性能问题访存与计算的权衡、共享内存的使用、分块算法等。准备好迎接真正的挑战吧十一、面试真题2024-2026Q1为什么向量加法在GPU上能获得很高的带宽利用率考察点对合并访问的理解参考答案向量加法的内存访问模式是完美的合并访问线程0读A[0]线程1读A[1]…线程31读A[31]这些地址连续且对齐。硬件可以将这32次4字节访问合并成一次128字节的事务最大化利用显存带宽。同时每个线程的计算非常简单不会成为瓶颈因此性能主要受限于内存带宽容易达到接近理论峰值的水平。Q2在测试GPU性能时你发现对于很小的NGPU比CPU慢为什么考察点对GPU开销的理解参考答案GPU有固定的启动开销kernel启动延迟约几微秒到几十微秒数据从主机到设备的传输延迟PCIe延迟设备可能处于低功耗状态需要唤醒当数据量很小时这些开销占主导而CPU可以立即执行所以总时间更短。因此GPU适合大规模并行计算小任务应留在CPU。Q3如何测量kernel的实际内存带宽计算公式是什么考察点性能测量方法参考答案使用CUDA事件记录kernel执行时间然后计算总数据访问量除以时间。对于向量加法每个元素需要读A、读B、写C所以总数据量 3 * N * sizeof(float)。带宽 总数据量 / 时间。注意这测量的是应用层带宽实际硬件可能还有缓存效应但作为近似足够。Q4你提到了使用float4优化为什么能提升性能考察点对指令级优化的理解参考答案使用float4可以让每个线程一次处理4个元素带来几个好处减少指令数4次加法可以用向量化指令减少取指和调度开销提高内存合并效率一次128位加载比4次32位加载更高效减少内存请求次数增加计算密度每个线程做更多工作有助于隐藏延迟但要注意边界处理和剩余元素的处理。思考题在你的GPU上运行性能对比代码画出不同N下CPU和GPU的时间曲线。找到你机器上GPU开始超越CPU的“交叉点”。对于这个交叉点的大小你有什么发现为什么不同GPU/CPU的交叉点不同欢迎在评论区分享你的实验数据。
第4节：CUDA实战——并行向量加法(从CPU到GPU的性能飞跃）
文章目录引言一、问题定义二、CPU基线版本2.1 实现代码2.2 CPU性能预期三、GPU基础版本3.1 核函数实现3.2 GPU性能预期四、性能对比分析4.1 不同数据规模下的对比4.2 为什么GPU在小数据量时慢4.3 带宽利用率分析五、探究线程配置对性能的影响5.1 改变block大小5.2 改变grid大小block数量六、深入剖析使用NVIDIA Nsight Compute6.1 安装Nsight Compute6.2 运行性能分析6.3 关键指标解读七、进阶优化使用共享内存7.1 使用float4优化八、完整性能对比代码框架九、常见问题与调试9.1 程序没输出或崩溃9.2 性能远低于预期9.3 数据传输瓶颈十、本节总结10.1 关键收获10.2 性能优化 checklist10.3 下节预告十一、面试真题2024-2026Q1为什么向量加法在GPU上能获得很高的带宽利用率Q2在测试GPU性能时你发现对于很小的NGPU比CPU慢为什么Q3如何测量kernel的实际内存带宽计算公式是什么Q4你提到了使用float4优化为什么能提升性能引言纸上得来终觉浅绝知此事要躬行前几节我们学习了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#includestdio.h#includestdlib.h#includetime.hvoidvector_add_cpu(constfloat*a,constfloat*b,float*c,intn){for(inti0;in;i){c[i]a[i]b[i];}}doubleget_time(){structtimespects;clock_gettime(CLOCK_MONOTONIC,ts);returnts.tv_sects.tv_nsec*1e-9;}intmain(){intn120;// 1,048,576 个元素size_t bytesn*sizeof(float);// 分配主机内存float*h_a(float*)malloc(bytes);float*h_b(float*)malloc(bytes);float*h_c(float*)malloc(bytes);// 初始化数据for(inti0;in;i){h_a[i]i*1.0f;h_b[i]i*2.0f;}// 计时doublestartget_time();vector_add_cpu(h_a,h_b,h_c,n);doubleendget_time();doubleelapsedend-start;doublebandwidth(bytes*2)/elapsed/1e9;// 读取AB写入C共3次内存访问实际上是读A和B各一次写C一次所以总数据量 3 * bytes// 但通常向量加法需要读A和B写C共3*bytesdoublegflopsn/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(inti0;i10;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_cpu2.2 CPU性能预期在一颗现代CPU如AMD Ryzen 9上对于1M元素约4MB数据预期结果时间约0.002-0.005秒带宽10-20 GB/sGFLOPs0.2-0.5 GFLOP/s这个性能已经不错但我们将看到GPU的恐怖之处。三、GPU基础版本3.1 核函数实现// vector_add_gpu_baseline.cu#includecuda_runtime.h#includestdio.h#includestdlib.h#defineCHECK_CUDA(call){\cudaError_t errcall;\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){inttidblockIdx.x*blockDim.xthreadIdx.x;if(tidn){c[tid]a[tid]b[tid];}}intmain(){intn120;// 1M 元素size_t bytesn*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(inti0;in;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_block256;intblocks_per_grid(nthreads_per_block-1)/threads_per_block;// 创建CUDA事件用于计时cudaEvent_t start,stop;CHECK_CUDA(cudaEventCreate(start));CHECK_CUDA(cudaEventCreate(stop));CHECK_CUDA(cudaEventRecord(start));vector_addblocks_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));doubleelapsedms/1000.0;// 转换为秒doublebandwidth(bytes*3)/elapsed/1e9;// 读AB写C共3*bytesdoublegflopsn/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(inti0;i10;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_gpu3.2 GPU性能预期在A100上预期时间约0.0001秒0.1ms带宽约1400 GB/s接近A100的峰值1555 GB/sGFLOPs约10 GFLOP/s速度对比GPU比CPU快20-50倍四、性能对比分析4.1 不同数据规模下的对比我们编写一个测试脚本对不同的N从1K到100M分别运行CPU和GPU版本记录时间。N元素个数数据量MBCPU时间msGPU时间ms加速比1,0000.0120.0020.050.0410,0000.120.0150.060.25100,0001.20.120.081.51,000,000121.20.15810,000,000120120.815100,000,0001200120815重要发现当数据量很小时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大小固定N1M改变threads_per_block观察性能变化block大小时间ms带宽GB/s说明320.22650占用率低640.188001280.169002560.15950最佳5120.155920略有下降10240.17840寄存器压力分析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-compute6.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){inttidblockIdx.x*blockDim.xthreadIdx.x;// 每个线程处理4个元素intidxtid*4;if(idx3n){float4 a4reinterpret_castfloat4*(a)[tid];float4 b4reinterpret_castfloat4*(b)[tid];float4 c4;c4.xa4.xb4.x;c4.ya4.yb4.y;c4.za4.zb4.z;c4.wa4.wb4.w;reinterpret_castfloat4*(c)[tid]c4;}else{// 处理剩余元素for(intiidx;in;i){c[i]a[i]b[i];}}}这样每个线程做4次加法但只发一次128位的内存请求合并更好可以减少指令数提高带宽利用率。实测性能可提升10-20%。八、完整性能对比代码框架为了方便你亲自实验这里提供一个完整的测试框架可以自动测试不同规模和不同block大小// vector_add_benchmark.cu#includecuda_runtime.h#includestdio.h#includestdlib.h#includevector#includealgorithm#defineCHECK_CUDA(call){...}// 同上__global__voidvector_add(float*a,float*b,float*c,intn){inttidblockIdx.x*blockDim.xthreadIdx.x;if(tidn)c[tid]a[tid]b[tid];}doubletest_gpu(intn,intthreads_per_block){size_t bytesn*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(inti0;in;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(nthreads_per_block-1)/threads_per_block;cudaEvent_t start,stop;cudaEventCreate(start);cudaEventCreate(stop);cudaEventRecord(start);vector_addblocks,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::vectorintsizes{1000,10000,100000,1000000,10000000,100000000};std::vectorintblock_sizes{64,128,256,512,1024};printf(N, BlockSize, Time(ms), Bandwidth(GB/s)\n);for(intn:sizes){for(intbs:block_sizes){doublemstest_gpu(n,bs);doublebytesn*sizeof(float)*3.0;// 读写doublebwbytes/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 关键收获向量加法是最简单的并行任务但能揭示GPU性能的基本规律GPU在大规模数据下优势巨大但小数据量不如CPU线程配置block大小影响性能需要实验找到最优值带宽利用率是衡量内存密集型kernel的关键指标使用性能分析工具才能真正理解瓶颈10.2 性能优化 checklist确保合并访问选择合适的block大小通常128-512隐藏数据传输异步拷贝考虑向量化加载float4避免寄存器溢出使用快速数学函数如果精度允许10.3 下节预告下一节我们将挑战更复杂的矩阵乘法它将暴露更多的性能问题访存与计算的权衡、共享内存的使用、分块算法等。准备好迎接真正的挑战吧十一、面试真题2024-2026Q1为什么向量加法在GPU上能获得很高的带宽利用率考察点对合并访问的理解参考答案向量加法的内存访问模式是完美的合并访问线程0读A[0]线程1读A[1]…线程31读A[31]这些地址连续且对齐。硬件可以将这32次4字节访问合并成一次128字节的事务最大化利用显存带宽。同时每个线程的计算非常简单不会成为瓶颈因此性能主要受限于内存带宽容易达到接近理论峰值的水平。Q2在测试GPU性能时你发现对于很小的NGPU比CPU慢为什么考察点对GPU开销的理解参考答案GPU有固定的启动开销kernel启动延迟约几微秒到几十微秒数据从主机到设备的传输延迟PCIe延迟设备可能处于低功耗状态需要唤醒当数据量很小时这些开销占主导而CPU可以立即执行所以总时间更短。因此GPU适合大规模并行计算小任务应留在CPU。Q3如何测量kernel的实际内存带宽计算公式是什么考察点性能测量方法参考答案使用CUDA事件记录kernel执行时间然后计算总数据访问量除以时间。对于向量加法每个元素需要读A、读B、写C所以总数据量 3 * N * sizeof(float)。带宽 总数据量 / 时间。注意这测量的是应用层带宽实际硬件可能还有缓存效应但作为近似足够。Q4你提到了使用float4优化为什么能提升性能考察点对指令级优化的理解参考答案使用float4可以让每个线程一次处理4个元素带来几个好处减少指令数4次加法可以用向量化指令减少取指和调度开销提高内存合并效率一次128位加载比4次32位加载更高效减少内存请求次数增加计算密度每个线程做更多工作有助于隐藏延迟但要注意边界处理和剩余元素的处理。思考题在你的GPU上运行性能对比代码画出不同N下CPU和GPU的时间曲线。找到你机器上GPU开始超越CPU的“交叉点”。对于这个交叉点的大小你有什么发现为什么不同GPU/CPU的交叉点不同欢迎在评论区分享你的实验数据。
