文章目录第 1 张CUDA 内存管理与数据传输总览第 2 张核心 API 与代码流程第 3 张实验结果总结第 4 张工程启示与优化建议实验代码第 1 张CUDA 内存管理与数据传输总览这张图主要讲清楚CPU 内存和 GPU 显存是两个不同的空间。核心要点是CPU 端的数据通常在普通 C 数组或std::vector里例如h_a、h_b、h_cGPU 端的数据在cudaMalloc申请出来的显存里例如d_a、d_b、d_c。两边不能直接共享数据必须通过cudaMemcpy搬运。典型流程是CPU 准备数据 → 拷贝到 GPU → GPU 执行 kernel → 拷回 CPU → 校验结果这张图的核心结论是CUDA 优化不仅要优化 kernel还要优化数据流。第 2 张核心 API 与代码流程这张图主要总结 用到的几个 CUDA API。核心 API 是cudaMalloc 在 GPU 显存上申请空间 cudaMemcpy 在 Host 和 Device 之间拷贝数据 cudaFree 释放 GPU 显存 cudaEvent 测量 GPU 操作耗时图中还强调了一个很重要的计时流程warmup → 正式计时 → 重复多次 → 求平均值原因是第一次 CUDA 调用可能包含额外初始化开销直接计时容易不准重复多次取平均可以减少偶然误差。这张图还给出了带宽计算公式Bandwidth(GB/s) bytes / (ms / 1000) / 1e9核心结论是不只看 kernel 时间还要看数据拷贝时间。第 3 张实验结果总结这张图是核心实验数据。实测结果显示H2D / D2H约 4.5–5.0 GB/s D2D 约 120 GB/s cudaMemset约 200–232 GB/s也就是说CPU 和 GPU 之间的数据传输速度比较慢而 GPU 显存内部操作快得多。最关键的对比是CPU↔GPU 跨设备传输约 5 GB/s GPU 显存内部拷贝约 120 GB/s GPU 显存内部写入约 230 GB/s这说明GPU 内部显存操作远快于 CPU↔GPU 跨设备数据传输。这也解释了为什么第一节课里vector_add的 kernel 很快但整体总时间并不快。第 4 张工程启示与优化建议这张图主要讲第二课的实践意义。上一节vector_add总时间更慢不是因为 GPU 算得慢而是因为GPU total time H2D Kernel D2H其中H2D和D2H数据传输占了大部分时间。低效做法是CPU → GPU → 小 kernel → GPU → CPU 反复多次更好的做法是CPU → GPU 一次 在 GPU 上连续执行多个 kernel 最后 GPU → CPU 一次优化原则可以概括为减少 Host↔Device 往返拷贝 数据搬上 GPU 后尽量留在 GPU 上 多个简单算子尽量融合 分析性能时区分 H2D / Kernel / D2H 同时关注 kernel time 和 end-to-end time核心金句是真正的 CUDA 优化不只是让 kernel 快而是让整个数据流更高效。实验代码#includecuda_runtime.h#includechrono#includecstdlib#includeiomanip#includeiostream#includevector/* * CHECK_CUDA 是一个错误检查宏。 * * CUDA API 调用失败时不会自动终止程序。 * 比如 cudaMalloc 失败、cudaMemcpy 失败 * 如果不检查错误后面可能出现更难定位的问题。 * * 所以建议所有 CUDA API 都包一层 CHECK_CUDA。 */#defineCHECK_CUDA(call)\do{\cudaError_terrcall;\if(err!cudaSuccess){\std::cerrCUDA error at __FILE__:__LINE__\ codestatic_castint(err)\ messagecudaGetErrorString(err)std::endl;\std::exit(EXIT_FAILURE);\}\}while(0)/* * 使用 cudaEvent 测量一段 CUDA 操作的时间。 * * 注意 * cudaEvent 适合测量 GPU 相关操作例如 * 1. kernel 执行时间 * 2. cudaMemcpy 时间 * 3. cudaMemset 时间 * * 它返回的时间单位是毫秒 ms。 */floatmeasureMemcpy(void*dst,constvoid*src,size_tbytes,cudaMemcpyKind kind,intrepeat){cudaEvent_tstart,stop;CHECK_CUDA(cudaEventCreate(start));CHECK_CUDA(cudaEventCreate(stop));/* * 先 warmup 一次。 * * 第一次 CUDA 操作可能包含额外初始化开销 * 所以不把第一次拷贝计入正式统计。 */CHECK_CUDA(cudaMemcpy(dst,src,bytes,kind));CHECK_CUDA(cudaDeviceSynchronize());/* * 正式计时。 * * 为了让结果更稳定不只测一次而是重复 repeat 次。 */CHECK_CUDA(cudaEventRecord(start));for(inti0;irepeat;i){CHECK_CUDA(cudaMemcpy(dst,src,bytes,kind));}CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));floattotal_ms0.0f;CHECK_CUDA(cudaEventElapsedTime(total_ms,start,stop));CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));/* * 返回单次平均耗时。 */returntotal_ms/repeat;}/* * 测量 cudaMemset 的耗时。 * * cudaMemset 常用于把 GPU 显存初始化为 0。 * 例如 * cudaMemset(d_a, 0, bytes); */floatmeasureMemset(void*dst,intvalue,size_tbytes,intrepeat){cudaEvent_tstart,stop;CHECK_CUDA(cudaEventCreate(start));CHECK_CUDA(cudaEventCreate(stop));// warmupCHECK_CUDA(cudaMemset(dst,value,bytes));CHECK_CUDA(cudaDeviceSynchronize());CHECK_CUDA(cudaEventRecord(start));for(inti0;irepeat;i){CHECK_CUDA(cudaMemset(dst,value,bytes));}CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));floattotal_ms0.0f;CHECK_CUDA(cudaEventElapsedTime(total_ms,start,stop));CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));returntotal_ms/repeat;}/* * 根据数据量和时间计算带宽。 * * bandwidth 数据量 / 时间 * * bytes 是字节数。 * ms 是毫秒。 * * GB/s bytes / (ms / 1000) / 1e9 */doublecalcBandwidthGBs(size_tbytes,floatms){returnstatic_castdouble(bytes)/(ms/1000.0)/1e9;}intmain(intargc,char**argv){/* * 默认测试 64M 个 float。 * * 1 个 float 4 bytes * 64M float 256 MB * * 这个数据规模足够大能够比较明显地体现传输带宽。 */intn126;/* * repeat 表示重复测试次数。 * 对于很小的数据单次时间太短误差会比较大。 * 重复多次取平均更稳定。 */intrepeat20;if(argc2){nstd::atoi(argv[1]);}if(argc3){repeatstd::atoi(argv[2]);}size_tbytesstatic_castsize_t(n)*sizeof(float);doublembbytes/1024.0/1024.0;std::coutCUDA Memory Copy Benchmarkstd::endl;std::coutElements : nstd::endl;std::coutSize : mb MBstd::endl;std::coutRepeat : repeatstd::endl;/* * h_a 和 h_b 是 CPU 端数组。 * * std::vector 默认分配在普通 CPU 内存中。 * 这种内存也叫 pageable memory。 */std::vectorfloath_a(n);std::vectorfloath_b(n);/* * 初始化 CPU 数据。 */for(inti0;in;i){h_a[i]static_castfloat(i%100);h_b[i]0.0f;}/* * d_a 和 d_b 是 GPU 端指针。 * * 注意 * 这里只是声明了指针还没有真正申请 GPU 显存。 */float*d_anullptr;float*d_bnullptr;/* * 在 GPU 显存中申请空间。 * * cudaMalloc 的第一个参数是二级指针 * 因为它需要修改 d_a / d_b 的值 * 让它们指向申请好的 GPU 显存地址。 */CHECK_CUDA(cudaMalloc(d_a,bytes));CHECK_CUDA(cudaMalloc(d_b,bytes));/* * 测试 Host to Device。 * * 数据方向 * CPU 内存 h_a - GPU 显存 d_a */floath2d_msmeasureMemcpy(d_a,h_a.data(),bytes,cudaMemcpyHostToDevice,repeat);/* * 测试 Device to Host。 * * 数据方向 * GPU 显存 d_a - CPU 内存 h_b */floatd2h_msmeasureMemcpy(h_b.data(),d_a,bytes,cudaMemcpyDeviceToHost,repeat);/* * 测试 Device to Device。 * * 数据方向 * GPU 显存 d_a - GPU 显存 d_b * * 这个测试不经过 CPU 内存 * 通常会比 Host-Device 拷贝快很多。 */floatd2d_msmeasureMemcpy(d_b,d_a,bytes,cudaMemcpyDeviceToDevice,repeat);/* * 测试 cudaMemset。 * * 把 d_b 对应的 GPU 显存全部设置为 0。 */floatmemset_msmeasureMemset(d_b,0,bytes,repeat);doubleh2d_bwcalcBandwidthGBs(bytes,h2d_ms);doubled2h_bwcalcBandwidthGBs(bytes,d2h_ms);doubled2d_bwcalcBandwidthGBs(bytes,d2d_ms);doublememset_bwcalcBandwidthGBs(bytes,memset_ms);/* * 输出结果。 */std::coutstd::fixedstd::setprecision(4);std::cout\n[Result]std::endl;std::coutHost to Device : std::setw(10)h2d_ms ms, std::setw(10)h2d_bw GB/sstd::endl;std::coutDevice to Host : std::setw(10)d2h_ms ms, std::setw(10)d2h_bw GB/sstd::endl;std::coutDevice to Device: std::setw(10)d2d_ms ms, std::setw(10)d2d_bw GB/sstd::endl;std::coutcudaMemset : std::setw(10)memset_ms ms, std::setw(10)memset_bw GB/sstd::endl;/* * 简单校验 Device to Host 拷贝是否成功。 * * h_b 是从 d_a 拷贝回来的。 * 如果拷贝正确h_b[i] 应该等于 h_a[i]。 */bool oktrue;for(inti0;in;i){if(h_a[i]!h_b[i]){okfalse;std::cerrMismatch at index i, h_ah_a[i], h_bh_b[i]std::endl;break;}}std::cout\n[Check]std::endl;std::coutCopy check: (ok?PASS:FAIL)std::endl;/* * 释放 GPU 显存。 * * 有 cudaMalloc就应该有 cudaFree。 */CHECK_CUDA(cudaFree(d_a));CHECK_CUDA(cudaFree(d_b));returnok?0:1;}