CUDA学习笔记（三）——CUDA 内存模型与异步执行

本阶段定位连接 CUDA 和 VPI 的桥梁。VPI 的VPIStream/VPIEvent概念直接源于这里。学完你会明白GPU 性能第一课到底是什么。1. 从上一篇文章的痛点说起之前每个程序都是这个节奏一步等一步CPU 准备数据 → [②拷进GPU] → [③kernel计算] → [④拷回CPU] → CPU 用结果 等它拷完 等它算完 等它拷完两个明显的浪费搬运本身很贵②④走的是 CPU↔GPU 之间的 PCIe 总线或车载平台的内存通道带宽有限。很多真实场景里搬数据的时间比算的时间还长。全程串行干等拷的时候计算单元闲着算的时候拷贝引擎闲着。如果要处理很多张图明明可以一边拷下一张、一边算这一张却在排队等。本文就解决这两件事① 怎么让搬运更少更快内存模型② 怎么让搬运和计算并行起来异步 stream。2. 三种内存Malloc / Pinned / Managed2.1cudaMalloc 普通malloc你已经在用CPU 侧malloc出的是可分页内存pageable——操作系统可能把它换到磁盘地址不固定。GPU 侧cudaMalloc出的是显存global memory。两者之间用cudaMemcpy搬。这是默认、最通用的方式。局限从可分页内存做拷贝时驱动得先偷偷把数据拷到一块锁定的中转区再传给 GPU多一道手续而且无法真正异步。2.2cudaMallocHostPinned / 锁页内存—— 异步的前提float* h_data; cudaMallocHost(h_data, bytes); // 分配 pinned host 内存 // ... 用法和普通 malloc 的指针一样 ... cudaFreeHost(h_data); // 对应的释放Pinned锁页内存告诉操作系统这块 CPU 内存别换出去地址钉死。好处拷贝更快省掉上面那道中转手续。是cudaMemcpyAsync真正异步的前提后续内容。代价占用物理内存、分配慢别滥用——只给频繁参与 H2D/D2H 传输的缓冲区用。2.3cudaMallocManagedUnified Memory / 统一内存—— 一个指针两头用float* data; cudaMallocManaged(data, bytes); // CPU 和 GPU 都能用同一个指针 ​ for (int i 0; i N; i) data[i] 1.0f; // CPU 直接写 myKernelg, b(data, N); // GPU 直接用同一个指针不用 cudaMemcpy cudaDeviceSynchronize(); // 等 GPU 用完再回 CPU 读 printf(%f\n, data[0]); // CPU 直接读 cudaFree(data);统一内存让 CPU 和 GPU 共享同一个指针数据在谁需要时由系统自动迁移代码里不用手写cudaMemcpy。好处代码简洁、少出错适合入门和原型。代价自动迁移在独显上仍走 PCIe省心不一定省时间性能敏感处仍需手动管理。在NVIDIA jetson平台上意义重大CPU 和 GPU 共享同一块物理内存统一内存可以做到近乎零拷贝——这正是后续省 memcpy的基础先埋个伏笔。2.4 三者怎么选入门期默认继续用cudaMalloccudaMemcpy打基础要玩异步重叠时host 侧改用cudaMallocHostpinned想写得省心试cudaMallocManaged。3. GPU 性能第一课减少 CPU↔GPU 搬运3.1 为什么搬运是头号敌人GPU 算力极强但数据要先从 CPU 内存翻过 PCIe 总线才能到显存。打个比方GPU 是一座算力惊人的工厂但原料要靠一条窄马路PCIe运进来。工厂再快马路堵死也没用。真实数据感受显存内部带宽常达每秒数百 GB 上 TB而 PCIe 单向带宽只有每秒十几几十 GB——差一个数量级。3.2 实践准则能不搬就不搬一串连续操作去畸变→缩放→格式转换应该全部留在 GPU 上一路做完中间结果不要来回倒腾回 CPU。搬运不可避免时让它和计算重叠。在NVIDIA jetson上用共享内存特性免去物理拷贝。记住这句优化 GPU 程序先想能不能少搬数据再想kernel 怎么写快。顺序不能反。4. 同步 vs 异步先认清默认行为kernel 启动是异步的CPU 发出命令后立刻返回往下走不等 GPU 算完。所以之前才要cudaDeviceSynchronize()。cudaMemcpy是同步的CPU 会阻塞等待拷贝完成才继续。cudaMemcpyAsync是异步的发出即返回拷贝在后台进行需要 pinned 内存 stream 才能真正和计算并行。理解这个谁等谁是玩转 stream 的基础。5. Stream流GPU 的异步任务队列5.1 什么是 stream一个 stream 就是一条任务队列你把拷贝、kernel 这些操作按顺序塞进去GPU 按塞入顺序执行。两条铁律同一个 stream 内操作严格按提交顺序执行前一个没完后一个不开始。不同 stream 之间互相独立可以并发硬件允许时同时进行。你之前没指定 stream所有操作都进了默认流default stream编号 0自然全是串行的。想并发就得开多条流。5.2 创建与使用cudaStream_t stream; cudaStreamCreate(stream); ​ // 把操作提交到这条 stream注意 Async 版 最后一个参数是 stream cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream); myKernelgrid, block, 0, stream(...); // 第 4 个启动参数是 stream cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream); ​ cudaStreamSynchronize(stream); // 等这条 stream 全部完成 cudaStreamDestroy(stream);注意 kernel 启动的完整四参数grid, block, sharedMemBytes, stream。前两个你已熟悉第三个是动态 shared memory 大小不用就填 0第四个就是 stream。6. 计算与传输重叠Overlap—— 异步的价值所在6.1 为什么能重叠现代 GPU 有独立的拷贝引擎Copy Engine和计算引擎Compute Engine能同时干活。于是可以串行默认流 [拷入图1][算图1][拷出图1][拷入图2][算图2]... 总时间长 ​ ↓ 用多 stream 重叠 ↓ ​ 重叠多 stream[拷入图1] [算图1 ][拷入图2] [算图2 ][拷入图3][拷出图1] 拷贝和计算在时间上叠起来总时间显著缩短6.2 前提条件要实现真正重叠缺一不可host 内存必须是pinnedcudaMallocHost——否则 Async 退化成同步。用cudaMemcpyAsync多个非默认 stream。把大任务切成多块或把多张图分到不同 stream才有东西可以互相重叠。6.3 典型套路多张图分到多条流并发处理 4 张图开 4 条 stream每条独立走拷入→计算→拷出硬件自动重叠const int NSTREAM 4; cudaStream_t streams[NSTREAM]; for (int i 0; i NSTREAM; i) cudaStreamCreate(streams[i]); ​ for (int i 0; i NSTREAM; i) { cudaMemcpyAsync(d_in[i], h_in[i], bytes, cudaMemcpyHostToDevice, streams[i]); myKernelgrid, block, 0, streams[i](d_in[i], d_out[i], ...); cudaMemcpyAsync(h_out[i], d_out[i], bytes, cudaMemcpyDeviceToHost, streams[i]); } for (int i 0; i NSTREAM; i) cudaStreamSynchronize(streams[i]);7. Event事件给 GPU 计时 流间同步7.1 为什么 GPU 计时不能用 CPU 计时器因为 kernel 是异步的如果你这样写// ✘ 错误kernel 还没算完CPU 就记了结束时间测出来几乎是 0 auto t0 clock(); myKernelg, b(...); auto t1 clock(); // kernel 其实还在 GPU 上跑CPU 发完命令立刻往下走测到的只是发命令的时间不是 GPU 真正的计算时间。正确做法是用cudaEvent在 GPU 的时间线上打时间戳。7.2 用cudaEvent计时cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); ​ cudaEventRecord(start); // 在 GPU 时间线上记开始 myKernelgrid, block(...); // 被计时的 kernel cudaEventRecord(stop); // 记结束 ​ cudaEventSynchronize(stop); // 等 stop 事件真正发生即 kernel 跑完 float ms 0; cudaEventElapsedTime(ms, start, stop); // 算出两事件间的毫秒数 printf(kernel 耗时 %.3f ms\n, ms); ​ cudaEventDestroy(start); cudaEventDestroy(stop);7.3 event 还能做流间同步除了计时event 能让一条 stream 等待另一条 stream 的某个点完成cudaStreamWaitEvent实现跨流的依赖协调。入门先掌握计时用法同步用法知道有即可。8. 验证练习练习 A · 给你写的kernel计时把前面的的向量加法或灰度化用cudaEvent包起来测时间#include cstdio ​ #define CUDA_CHECK(call) do { \ cudaError_t e (call); \ if (e ! cudaSuccess) { \ fprintf(stderr, CUDA %s:%d %s\n, \ __FILE__, __LINE__, cudaGetErrorString(e)); \ return 1; \ } \ } while(0) ​ __global__ void vecAdd(const float* a, const float* b, float* c, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) c[i] a[i] b[i]; } ​ int main() { int N 1 22; size_t bytes N * sizeof(float); ​ float *d_a, *d_b, *d_c; CUDA_CHECK(cudaMalloc(d_a, bytes)); CUDA_CHECK(cudaMalloc(d_b, bytes)); CUDA_CHECK(cudaMalloc(d_c, bytes)); ​ cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); ​ int tpb 256, blocks (N tpb - 1) / tpb; cudaEventRecord(start); vecAddblocks, tpb(d_a, d_b, d_c, N); // 被计时对象 cudaEventRecord(stop); cudaEventSynchronize(stop); ​ float ms 0; cudaEventElapsedTime(ms, start, stop); printf(vecAdd(N%d) 耗时 %.3f ms\n, N, ms); ​ cudaEventDestroy(start); cudaEventDestroy(stop); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; }玩法改 N 大小、改 tpb128/256/512看耗时怎么变——开始建立什么因素影响性能的手感。练习 B · 多 stream 并发处理多张图用 pinned 内存 多 stream对比单流串行和多流并发的耗时差异#include cstdio ​ #define CUDA_CHECK(call) do { \ cudaError_t e (call); \ if (e ! cudaSuccess) { \ fprintf(stderr, CUDA %s:%d %s\n, \ __FILE__, __LINE__, cudaGetErrorString(e)); \ return 1; \ } \ } while(0) ​ // 模拟一张图的处理每个像素做点计算故意多算几轮让 kernel 有耗时 __global__ void process(const float* in, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { float v in[i]; for (int k 0; k 200; k) v v * 1.0001f 0.5f; out[i] v; } } ​ int main() { const int NIMG 4; // 4 张图 const int N 1 20; // 每张图 100 万像素 size_t bytes N * sizeof(float); int tpb 256, blocks (N tpb - 1) / tpb; ​ // pinned host 内存异步重叠的前提 float *h_in[NIMG], *h_out[NIMG]; float *d_in[NIMG], *d_out[NIMG]; for (int i 0; i NIMG; i) { CUDA_CHECK(cudaMallocHost(h_in[i], bytes)); CUDA_CHECK(cudaMallocHost(h_out[i], bytes)); CUDA_CHECK(cudaMalloc(d_in[i], bytes)); CUDA_CHECK(cudaMalloc(d_out[i], bytes)); for (int j 0; j N; j) h_in[i][j] 1.0f; } ​ cudaStream_t streams[NIMG]; for (int i 0; i NIMG; i) cudaStreamCreate(streams[i]); ​ cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); ​ // ---- 多流并发每张图独立走 拷入→算→拷出硬件自动重叠 ---- cudaEventRecord(start); for (int i 0; i NIMG; i) { cudaMemcpyAsync(d_in[i], h_in[i], bytes, cudaMemcpyHostToDevice, streams[i]); processblocks, tpb, 0, streams[i](d_in[i], d_out[i], N); cudaMemcpyAsync(h_out[i], d_out[i], bytes, cudaMemcpyDeviceToHost, streams[i]); } for (int i 0; i NIMG; i) cudaStreamSynchronize(streams[i]); cudaEventRecord(stop); cudaEventSynchronize(stop); ​ float ms 0; cudaEventElapsedTime(ms, start, stop); printf(多流并发处理 %d 张图总耗时 %.3f ms\n, NIMG, ms); ​ for (int i 0; i NIMG; i) { cudaStreamDestroy(streams[i]); cudaFreeHost(h_in[i]); cudaFreeHost(h_out[i]); cudaFree(d_in[i]); cudaFree(d_out[i]); } cudaEventDestroy(start); cudaEventDestroy(stop); return 0; }进阶对照实验把上面循环改成全用默认流去掉 stream 参数、用同步cudaMemcpy测串行耗时和多流版对比——直观感受重叠带来的加速。9. 呼应后面的高层库——VPI本阶段 CUDA 概念VPI 对应cudaStream_t异步队列VPIStreamVPI 的执行流思想一模一样cudaEvent_t计时/同步VPIEventVPI 的同步事件操作提交到流、异步执行vpiSubmit*系列把算法提交到VPIStream减少 CPU↔GPU 搬运VPIImage包装显存指针做 zero-copy10. 常见坑异步却不同步就读结果cudaMemcpyAsync/kernel发出后没等它完成就读数据 → 读到旧值。要cudaStreamSynchronize或 event 同步。用了 Async 但 host 是普通malloc不会真正异步悄悄退化为同步重叠失效。异步必须配 pinnedcudaMallocHost。所有操作都进了默认流默认流会和其它流串行化有特殊同步语义。想并发显式建多条非默认流。用 CPU 计时器测 kernel测出接近 0 的假时间。GPU 计时用cudaEvent。cudaEvent忘了cudaEventSynchronize(stop)直接ElapsedTime可能拿到未完成的结果。pinned 内存分配过多占满物理内存拖慢系统。只给传输缓冲用。11. 通往下一步你现在具备GPU 并行模型、手写 kernel、内存模型、异步/stream/event——这些是能迁移到任何 GPU 库的硬本事。接下来重心转向应用库层。后面不再从零写 kernel而是学会调用 VPI 这套现成的高层视觉 API你会发现VPI 的VPIStream/VPIEvent→ 就是这一阶段的cudaStream/cudaEventVPI 的 Rescale/ConvertImageFormat → 就是你手写过的那些 kernel别人替你写好调优了减少搬运、算法留在 GPU 上串起来 → 就是本文里的性能第一课在 VPI 流水线里的体现。