1. 这一阶段和前面的不同前三阶段你在从零写 kernel这一阶段基本不写 kernel 了——VPI 已经把去畸变、缩放、格式转换这些 kernel 写好并调优了你的工作变成学会调用这套高层 API可以用 VPI 搭真实流水线。核心心态转变阶段二你会由衷觉得这些 kernel 我自己也能写所以现在看 VPI你不是在学黑魔法而是在确认别人把我会写的东西写得更好了。这正是先打 CUDA 地基的回报。2. VPI 是什么VPIVision Programming Interface 一套统一的计算机视觉 API一份代码可切换多个后端CUDA/GPU、VIC、PVA、CPU运行。VPI 是盖在 CUDA 之上的现成房间它把常见视觉算法去畸变 Remap、缩放 Rescale、格式转换 ConvertImageFormat、光流、立体匹配……封装好你选后端、喂数据、调用即可。选哪个后端跑是代码里逐个算法指定的。3. VPI 六大核心概念重点且大半你已认识VPI 概念是什么阶段三的同名兄弟VPIImageVPI 的图像对象封装像素数据 格式 布局—新概念但就是带格式的显存/内存块VPIStream执行流算法提交到它、异步执行cudaStream_t思想一模一样VPIEvent同步与计时事件标记流上某点是否完成cudaEvent_t一模一样VPIPayload算法预分配的资源/上下文如 Remap 的映射表≈ 你为 kernel 提前malloc好的辅助缓冲VPIBackend后端选择VPI_BACKEND_CUDA/_VIC/_CPU/_PVA≈ 在哪块硬件上跑vpiSubmit*把某个算法提交到 stream 上异步执行≈myKernel...,stream(...)启动看到没VPIStream就是cudaStream、VPIEvent就是cudaEvent、vpiSubmit*就是往流里塞一个 kernel。关于VPIPayload唯一需要额外理解的有些算法如 Remap 去畸变需要一张预先算好的映射表/查找表构建它很慢。VPI 让你在初始化时构建一次 payload之后每帧复用——避免每帧重算。这是预分配资源的思想。4. VPI 程序的标准骨架给出一个示例步骤如下读一张图 → 转成 NV12 → 缩放 → 转回 BGR → 存盘几乎所有 VPI 程序都是这五步骨架初始化建 stream、建/包装VPIImage计算vpiSubmit*把一串算法提交到 stream异步同步vpiStreamSync等 stream 全部跑完取结果vpiImageLockData锁定读出 → 用完vpiImageUnlock清理vpiImageDestroy/vpiStreamDestroy示例伪代码// ① 初始化 ------------------------------------------------- vpiStreamCreate(backend | VPI_BACKEND_CUDA, stream); // 建流可多后端 vpiImageCreateWrapper(cvImage, 0, image); // 把图像包成 VPIImage零拷贝 vpiImageCreate(w, h, VPI_IMAGE_FORMAT_NV12_ER, 0, imageNV12); // 建中间图 vpiImageCreate(w/2, h/3, VPI_IMAGE_FORMAT_NV12_ER, 0, outputNV12); // 建输出图 vpiImageCreate(w/2, h/3, VPI_IMAGE_FORMAT_BGR8, 0, output); // ② 计算三个算法串成一条链全部提交到同一个 stream -------- vpiSubmitConvertImageFormat(stream, VPI_BACKEND_CUDA, image, imageNV12, NULL); // BGR→NV12 vpiSubmitRescale(stream, backend, imageNV12, outputNV12, VPI_INTERP_LINEAR, VPI_BORDER_CLAMP, 0); // 缩放 vpiSubmitConvertImageFormat(stream, VPI_BACKEND_CUDA, outputNV12, output, NULL); // NV12→BGR // ③ 同步等 GPU/VIC 把这条链跑完 -------------------------- vpiStreamSync(stream); // ④ 取结果锁定输出图读出像素写盘 -------------------- vpiImageLockData(output, VPI_LOCK_READ, VPI_IMAGE_BUFFER_HOST_PITCH_LINEAR, outData); // ... 用 outData.buffer.pitch.planes[0] 构造 cv::Mat 存盘 ... vpiImageUnlock(output); // ⑤ 清理 -------------------------------------------------- vpiImageDestroy(image); vpiImageDestroy(imageNV12); ... vpiStreamDestroy(stream);和前文直接对应vpiSubmitConvertImageFormatBGR→NV12 你手写的 RGB→NV12 kernelVPI 替你写好了。vpiSubmitRescale 你手写的缩放 kernelVPI 替你写好且是高质量插值。vpiStreamSync 前面的cudaStreamSynchronize。三个算法提交到同一个 stream → 同一 stream 内按序执行天然形成流水线。这就是写完 kernel 再看 VPI瞬间明白 VPI 只是帮我把这些 kernel 写好调优了。5. 后端选择CUDA / VIC / PVA / CPU5.1 四个后端分别是什么硬件后端硬件擅长备注VPI_BACKEND_CUDAGPUCUDA Core通用并行计算、复杂算法算力最强但和 AI 推理抢 GPUVPI_BACKEND_VICVICVideo Image Compositor固定功能硬件缩放、格式转换、去畸变这类杂活Orin/Thor 独有独立于 GPUVPI_BACKEND_PVAPVA可编程向量加速器特定向量算法如立体匹配、光流低功耗VPI_BACKEND_CPUCPU兜底、无 GPU 环境最慢5.2 边缘端设备分工哲学关键动机一句话GPU 是稀缺资源要留给 AI 推理TensorRT。去畸变、缩放、格式转换这些杂活丢给专门的 VIC 硬件去干GPU 就能腾出来专心跑神经网络。通常情况下图像预处理流水线常选“VIC”后端——不是因为 VIC 比 GPU 快而是分工让每种硬件干它最该干的事整机吞吐最大化。在没有 VIC 的普通 x86 独显机器上就只能退回 CUDA/CPU 后端。5.3 切后端有多简单算法代码一行不改vpiStreamCreate(type, ...)这个type可以指定“CUDA”或“VIC”——算法逻辑、流水线结构完全不用改。这正是 VPI 统一接口、多后端的价值写一次哪个硬件跑由一个参数决定。5.4 注意不是所有算法都支持所有后端VPI 每个算法各自支持哪些后端不同比如某些复杂算法只有 CUDA 后端。官方文档每个算法页都有一张后端支持表。选后端前先查算法是否支持——这也是官方示例常写成backend | VPI_BACKEND_CUDA叠加 CUDA 兜底的原因。6. 互操作的动机为什么要零拷贝回顾笔记三的性能第一课CPU↔GPU、甚至设备内不必要的拷贝都是浪费。设想一条真实链路相机原始帧 → VPI 去畸变/缩放VIC 或 CUDA→ 结果图 → 你的 CUDA kernelHWC→CHW 归一化→ TensorRT 推理如果每一步之间都cudaMemcpy复制一份图像数据几路相机 × 每秒几十帧拷贝开销会吃掉大量带宽和时间。零拷贝互操作的思想VPI 和你的 CUDA kernel 直接操作同一块显存——VPI 写完你的 kernel 就地读中间不复制。这就是vpiImageCreateWrapper刚才见过它包装 host 内存的另一面它也能包装 CUDA 显存指针让VPIImage和裸显存指向同一块内存。7. 两种包装方式Host vs CUDAvpiImageCreateWrapper靠VPIImageData.bufferType区分包装的是哪种内存bufferType包装的内存谁能用VPI_IMAGE_BUFFER_HOST_PITCH_LINEARCPUhost内存指针CPU 后端Tegra 上也可给 VICVPI_IMAGE_BUFFER_CUDA_PITCH_LINEARCUDA 显存指针cudaMalloc来的CUDA 后端VPI_IMAGE_BUFFER_NVBUFFERTegra 的NvBufSurface设备上的各硬件8. 实操 A把 CUDA 显存包成 VPIImage 喂给 VPI#include vpi/Image.h #include vpi/Stream.h #include vpi/algo/Rescale.h #include cuda_runtime.h // 假设你已有一块 CUDA 显存里的 NV12 图比如你上一个 kernel 的输出 uint8_t* d_nv12; // cudaMalloc 得到的显存指针 int w 1920, h 1080; size_t ySize (size_t)w * h; // Y 平面 cudaMalloc(d_nv12, ySize * 3 / 2); // NV12 W*H*3/2 // ---- 把这块显存包成 VPIImage零拷贝不复制像素---- VPIImageData data{}; data.bufferType VPI_IMAGE_BUFFER_CUDA_PITCH_LINEAR; // ★ 关键CUDA 显存 VPIImageBufferPitchLinear buf data.buffer.pitch; buf.format VPI_IMAGE_FORMAT_NV12_ER; buf.numPlanes 2; buf.planes[0].pixelType VPI_PIXEL_TYPE_INVALID; buf.planes[0].width w; buf.planes[0].height h; buf.planes[0].pitchBytes w; buf.planes[0].offsetBytes 0; buf.planes[0].pBase (VPIByte*)d_nv12; // ★ 指向 CUDA 显存 buf.planes[1].width w / 2; buf.planes[1].height h / 2; buf.planes[1].pitchBytes w; buf.planes[1].offsetBytes ySize; // UV 平面接在 Y 之后 buf.planes[1].pBase (VPIByte*)d_nv12; VPIImage vpiSrc nullptr; vpiImageCreateWrapper(data, nullptr, VPI_BACKEND_CUDA, vpiSrc); // 用 CUDA 后端标志 // 之后就能把 vpiSrc 交给任意 VPI CUDA 算法VPI 直接读这块显存无拷贝 // vpiSubmitRescale(stream, VPI_BACKEND_CUDA, vpiSrc, vpiDst, ...);9. 实操 B反向——从 VPI 结果取出显存指针喂给自己的 kernelVPI 缩放完你想接一个自己写的 kernel比如以前的的 HWC→CHW 归一化为 TensorRT 准备输入。用vpiImageLockData以 CUDA 视角锁定就能拿到显存指针// 你自己的 kernel阶段二写过的 HWC→CHW这里演示直接吃 VPI 输出的显存 __global__ void myPostKernel(const uint8_t* d_img, float* d_tensor, int w, int h); // ---- VPI 处理完后取出它输出图的 CUDA 显存指针 ---- VPIImageData out{}; vpiStreamSync(stream); // 先确保 VPI 那条流的算法真的跑完了 vpiImageLockData(vpiDst, VPI_LOCK_READ, VPI_IMAGE_BUFFER_CUDA_PITCH_LINEAR, // ★ 以 CUDA 显存视角锁定 out); uint8_t* d_result (uint8_t*)out.buffer.pitch.planes[0].pBase; // ★ 就是显存指针 int pitch out.buffer.pitch.planes[0].pitchBytes; // 直接把这个显存指针喂给你自己的 CUDA kernel —— 全程没有一次 cudaMemcpy dim3 block(16, 16), grid((w15)/16, (h15)/16); myPostKernelgrid, block(d_result, d_tensor, w, h); cudaDeviceSynchronize(); vpiImageUnlock(vpiDst); // lock/unlock 必须配对d_result是 VPI 写好的显存myPostKernel是你手写的 kernel两者共用同一块显存中间零拷贝。之前的手写 kernel和这里的用 VPI在这里合流了。10. 黄金法则先测量再优化Dont guess, measure前面你反复听到这样更快合并访问快、shared memory 快、多流重叠快、零拷贝快、VIC 省 GPU……但这些到目前为止都还是别人告诉你的结论。下面你要学会两个测量工具Nsight Systems / Nsight Compute。亲手验证前面的结论并学会定位自己代码的瓶颈。GPU 优化最大的坑是凭直觉猜瓶颈然后瞎改。90% 的我以为的瓶颈都是错的。正确流程永远是测量→ 找到真正最慢的那一段瓶颈判断瓶颈类型访存计算搬运同步等待针对性优化那一段再测量确认真的变快了而不是自我感觉11. 三个测量层级从粗到细层级工具回答的问题何时用代码内打点cudaEvent阶段三/vpiEventElapsedTimeMillis这段/这个 kernel 花了几毫秒快速定位哪一步慢全局时间线Nsight Systems (nsys)整个程序时间都花哪了拷贝和计算重叠了吗GPU 有没有闲着先用它看全貌单 kernel 深剖Nsight Compute (ncu)这个 kernel 卡在访存还是计算占用率多少访问合并了吗锁定某个热点 kernel 后深挖正确顺序先用nsys看全局找热点 → 再用ncu剖那个最热的 kernel。别一上来就ncu它很慢。12. Nsight Systemsnsys看整条流水线的时间线12.1 它是什么把程序运行的时间线画出来CPU 在干嘛、GPU 在跑哪个 kernel、H2D/D2H 拷贝在什么时候、多个 stream 是否并行、哪里有空隙gap。这是发现搬运没重叠GPU 在干等这类问题的最佳工具。12.2 怎么用# 采集在有 GPU 的机器上生成 report.nsys-rep nsys profile --statstrue -o report ./multi_stream # 命令行直接看统计摘要--statstrue 已经会打印 # 图形界面看时间线本机也装了 nsys-ui nsys-ui report.nsys-rep12.3 重点看什么对应前面阶段的验证验证笔记三的多流重叠打开你之前写的的multi_stream在时间线上看 4 条 stream 的 H2D 拷贝、kernel、D2H 是不是在时间上叠起来了重叠成功还是一条接一条没重叠——多半是忘了用 pinned 内存。找 GPU 空隙gapGPU 时间线上大段空白 GPU 在等 CPU 或等拷贝。空隙多说明瓶颈在别处CPU、传输、同步。看 H2D/D2H 占比如果拷贝时间比 kernel 计算时间还长 → 印证搬运是头号敌人该想零拷贝/减少搬运。13. Nsight Computencu剖单个 kernel13.1 它是什么对单个 kernel 做显微镜级分析占用率、内存带宽利用、是计算受限还是访存受限、访问是否合并、warp 是否发散……用来回答这个 kernel 为什么慢、怎么改快。13.2 怎么用# 剖析全部指标很详细也很慢会把 kernel 重放很多次 ncu --set full -o report ./grayscale # 只剖某个 kernel、只测第一次启动快很多 ncu --kernel-name rgb2gray --launch-count 1 ./grayscale # 图形界面看报告 ncu-ui report.ncu-rep注意 ncu会反复重放replay同一个 kernel 来采集不同指标所以比正常运行慢几十倍属正常现象。13.3 三个最该先看的指标Achieved Occupancy实际占用率SM 上实际活跃的 warp 占理论最大的比例。太低如 30%说明并行度不够——常见原因block 太小、每线程寄存器/shared memory 用太多。先调 block 大小128/256/512试试。Memory Throughput vs Compute Throughput哪个接近上限就是瓶颈类型。访存接近上限、计算很低 →memory-bound多数图像/向量类简单 kernel 都是优化方向是访存合并、shared memory、减少读写。计算接近上限 →compute-bound优化方向是减少计算量或用更快指令。Memory Coalescing访存合并相关警告ncu会直接提示访问是否合并。14. 实战剖析笔记二的手写 kernel找瓶颈14.1 步骤ncu --set full -o gray_report ./grayscale ncu-ui gray_report.ncu-rep # 看 Occupancy 和 Memory Throughput预期发现多数简单图像 kernel 的典型结论是memory-bound访存受限——因为每个像素只做几次乘加却要读写显存计算单元大部分时间在等数据。这解释了为什么阶段二强调合并访问既然瓶颈是访存那把访存做到合并就是最有效的优化。14.2 亲手做个对比实验强烈推荐写两个版本对比最能建立直觉A合并版阶段二的正常写法threadIdx.x对应连续列。B非合并版故意让每个线程隔很远取数据制造跨越访问。用ncu分别测你会看到 B 的 Memory Throughput 明显更差、耗时更长。14.3 调 block 大小看 occupancy把threadsPerBlock从 64 → 128 → 256 → 512 分别测耗时和 occupancy观察拐点。你会发现太小的 block 占用率低、太大也不一定更好——这就是为什么阶段一说 block 取 32 的倍数、通常 128~512。15. 常见优化手段清单定位瓶颈后按需取用务必先测量定位瓶颈再从下面选对应手段不要无脑全上瓶颈类型优化手段搬运太多H2D/D2H 占比高零拷贝、减少 CPU↔GPU 传输、算法留在 GPU 串起来访存受限、非合并合并访问、用 shared memory 复用、减少冗余读写占用率低调 block 大小、降低每线程寄存器/shared 用量GPU 与 CPU/拷贝互相干等多流并发、计算与传输重叠GPU 被杂活占满、抢不到算力把预处理丢给 VIC/PVA 后端warp 分支发散减少同一 warp 内的 if/else 分歧 CUDA学习笔记（四）——CUDA kernel和VPI库的互操作
1. 这一阶段和前面的不同前三阶段你在从零写 kernel这一阶段基本不写 kernel 了——VPI 已经把去畸变、缩放、格式转换这些 kernel 写好并调优了你的工作变成学会调用这套高层 API可以用 VPI 搭真实流水线。核心心态转变阶段二你会由衷觉得这些 kernel 我自己也能写所以现在看 VPI你不是在学黑魔法而是在确认别人把我会写的东西写得更好了。这正是先打 CUDA 地基的回报。2. VPI 是什么VPIVision Programming Interface 一套统一的计算机视觉 API一份代码可切换多个后端CUDA/GPU、VIC、PVA、CPU运行。VPI 是盖在 CUDA 之上的现成房间它把常见视觉算法去畸变 Remap、缩放 Rescale、格式转换 ConvertImageFormat、光流、立体匹配……封装好你选后端、喂数据、调用即可。选哪个后端跑是代码里逐个算法指定的。3. VPI 六大核心概念重点且大半你已认识VPI 概念是什么阶段三的同名兄弟VPIImageVPI 的图像对象封装像素数据 格式 布局—新概念但就是带格式的显存/内存块VPIStream执行流算法提交到它、异步执行cudaStream_t思想一模一样VPIEvent同步与计时事件标记流上某点是否完成cudaEvent_t一模一样VPIPayload算法预分配的资源/上下文如 Remap 的映射表≈ 你为 kernel 提前malloc好的辅助缓冲VPIBackend后端选择VPI_BACKEND_CUDA/_VIC/_CPU/_PVA≈ 在哪块硬件上跑vpiSubmit*把某个算法提交到 stream 上异步执行≈myKernel...,stream(...)启动看到没VPIStream就是cudaStream、VPIEvent就是cudaEvent、vpiSubmit*就是往流里塞一个 kernel。关于VPIPayload唯一需要额外理解的有些算法如 Remap 去畸变需要一张预先算好的映射表/查找表构建它很慢。VPI 让你在初始化时构建一次 payload之后每帧复用——避免每帧重算。这是预分配资源的思想。4. VPI 程序的标准骨架给出一个示例步骤如下读一张图 → 转成 NV12 → 缩放 → 转回 BGR → 存盘几乎所有 VPI 程序都是这五步骨架初始化建 stream、建/包装VPIImage计算vpiSubmit*把一串算法提交到 stream异步同步vpiStreamSync等 stream 全部跑完取结果vpiImageLockData锁定读出 → 用完vpiImageUnlock清理vpiImageDestroy/vpiStreamDestroy示例伪代码// ① 初始化 ------------------------------------------------- vpiStreamCreate(backend | VPI_BACKEND_CUDA, stream); // 建流可多后端 vpiImageCreateWrapper(cvImage, 0, image); // 把图像包成 VPIImage零拷贝 vpiImageCreate(w, h, VPI_IMAGE_FORMAT_NV12_ER, 0, imageNV12); // 建中间图 vpiImageCreate(w/2, h/3, VPI_IMAGE_FORMAT_NV12_ER, 0, outputNV12); // 建输出图 vpiImageCreate(w/2, h/3, VPI_IMAGE_FORMAT_BGR8, 0, output); // ② 计算三个算法串成一条链全部提交到同一个 stream -------- vpiSubmitConvertImageFormat(stream, VPI_BACKEND_CUDA, image, imageNV12, NULL); // BGR→NV12 vpiSubmitRescale(stream, backend, imageNV12, outputNV12, VPI_INTERP_LINEAR, VPI_BORDER_CLAMP, 0); // 缩放 vpiSubmitConvertImageFormat(stream, VPI_BACKEND_CUDA, outputNV12, output, NULL); // NV12→BGR // ③ 同步等 GPU/VIC 把这条链跑完 -------------------------- vpiStreamSync(stream); // ④ 取结果锁定输出图读出像素写盘 -------------------- vpiImageLockData(output, VPI_LOCK_READ, VPI_IMAGE_BUFFER_HOST_PITCH_LINEAR, outData); // ... 用 outData.buffer.pitch.planes[0] 构造 cv::Mat 存盘 ... vpiImageUnlock(output); // ⑤ 清理 -------------------------------------------------- vpiImageDestroy(image); vpiImageDestroy(imageNV12); ... vpiStreamDestroy(stream);和前文直接对应vpiSubmitConvertImageFormatBGR→NV12 你手写的 RGB→NV12 kernelVPI 替你写好了。vpiSubmitRescale 你手写的缩放 kernelVPI 替你写好且是高质量插值。vpiStreamSync 前面的cudaStreamSynchronize。三个算法提交到同一个 stream → 同一 stream 内按序执行天然形成流水线。这就是写完 kernel 再看 VPI瞬间明白 VPI 只是帮我把这些 kernel 写好调优了。5. 后端选择CUDA / VIC / PVA / CPU5.1 四个后端分别是什么硬件后端硬件擅长备注VPI_BACKEND_CUDAGPUCUDA Core通用并行计算、复杂算法算力最强但和 AI 推理抢 GPUVPI_BACKEND_VICVICVideo Image Compositor固定功能硬件缩放、格式转换、去畸变这类杂活Orin/Thor 独有独立于 GPUVPI_BACKEND_PVAPVA可编程向量加速器特定向量算法如立体匹配、光流低功耗VPI_BACKEND_CPUCPU兜底、无 GPU 环境最慢5.2 边缘端设备分工哲学关键动机一句话GPU 是稀缺资源要留给 AI 推理TensorRT。去畸变、缩放、格式转换这些杂活丢给专门的 VIC 硬件去干GPU 就能腾出来专心跑神经网络。通常情况下图像预处理流水线常选“VIC”后端——不是因为 VIC 比 GPU 快而是分工让每种硬件干它最该干的事整机吞吐最大化。在没有 VIC 的普通 x86 独显机器上就只能退回 CUDA/CPU 后端。5.3 切后端有多简单算法代码一行不改vpiStreamCreate(type, ...)这个type可以指定“CUDA”或“VIC”——算法逻辑、流水线结构完全不用改。这正是 VPI 统一接口、多后端的价值写一次哪个硬件跑由一个参数决定。5.4 注意不是所有算法都支持所有后端VPI 每个算法各自支持哪些后端不同比如某些复杂算法只有 CUDA 后端。官方文档每个算法页都有一张后端支持表。选后端前先查算法是否支持——这也是官方示例常写成backend | VPI_BACKEND_CUDA叠加 CUDA 兜底的原因。6. 互操作的动机为什么要零拷贝回顾笔记三的性能第一课CPU↔GPU、甚至设备内不必要的拷贝都是浪费。设想一条真实链路相机原始帧 → VPI 去畸变/缩放VIC 或 CUDA→ 结果图 → 你的 CUDA kernelHWC→CHW 归一化→ TensorRT 推理如果每一步之间都cudaMemcpy复制一份图像数据几路相机 × 每秒几十帧拷贝开销会吃掉大量带宽和时间。零拷贝互操作的思想VPI 和你的 CUDA kernel 直接操作同一块显存——VPI 写完你的 kernel 就地读中间不复制。这就是vpiImageCreateWrapper刚才见过它包装 host 内存的另一面它也能包装 CUDA 显存指针让VPIImage和裸显存指向同一块内存。7. 两种包装方式Host vs CUDAvpiImageCreateWrapper靠VPIImageData.bufferType区分包装的是哪种内存bufferType包装的内存谁能用VPI_IMAGE_BUFFER_HOST_PITCH_LINEARCPUhost内存指针CPU 后端Tegra 上也可给 VICVPI_IMAGE_BUFFER_CUDA_PITCH_LINEARCUDA 显存指针cudaMalloc来的CUDA 后端VPI_IMAGE_BUFFER_NVBUFFERTegra 的NvBufSurface设备上的各硬件8. 实操 A把 CUDA 显存包成 VPIImage 喂给 VPI#include vpi/Image.h #include vpi/Stream.h #include vpi/algo/Rescale.h #include cuda_runtime.h // 假设你已有一块 CUDA 显存里的 NV12 图比如你上一个 kernel 的输出 uint8_t* d_nv12; // cudaMalloc 得到的显存指针 int w 1920, h 1080; size_t ySize (size_t)w * h; // Y 平面 cudaMalloc(d_nv12, ySize * 3 / 2); // NV12 W*H*3/2 // ---- 把这块显存包成 VPIImage零拷贝不复制像素---- VPIImageData data{}; data.bufferType VPI_IMAGE_BUFFER_CUDA_PITCH_LINEAR; // ★ 关键CUDA 显存 VPIImageBufferPitchLinear buf data.buffer.pitch; buf.format VPI_IMAGE_FORMAT_NV12_ER; buf.numPlanes 2; buf.planes[0].pixelType VPI_PIXEL_TYPE_INVALID; buf.planes[0].width w; buf.planes[0].height h; buf.planes[0].pitchBytes w; buf.planes[0].offsetBytes 0; buf.planes[0].pBase (VPIByte*)d_nv12; // ★ 指向 CUDA 显存 buf.planes[1].width w / 2; buf.planes[1].height h / 2; buf.planes[1].pitchBytes w; buf.planes[1].offsetBytes ySize; // UV 平面接在 Y 之后 buf.planes[1].pBase (VPIByte*)d_nv12; VPIImage vpiSrc nullptr; vpiImageCreateWrapper(data, nullptr, VPI_BACKEND_CUDA, vpiSrc); // 用 CUDA 后端标志 // 之后就能把 vpiSrc 交给任意 VPI CUDA 算法VPI 直接读这块显存无拷贝 // vpiSubmitRescale(stream, VPI_BACKEND_CUDA, vpiSrc, vpiDst, ...);9. 实操 B反向——从 VPI 结果取出显存指针喂给自己的 kernelVPI 缩放完你想接一个自己写的 kernel比如以前的的 HWC→CHW 归一化为 TensorRT 准备输入。用vpiImageLockData以 CUDA 视角锁定就能拿到显存指针// 你自己的 kernel阶段二写过的 HWC→CHW这里演示直接吃 VPI 输出的显存 __global__ void myPostKernel(const uint8_t* d_img, float* d_tensor, int w, int h); // ---- VPI 处理完后取出它输出图的 CUDA 显存指针 ---- VPIImageData out{}; vpiStreamSync(stream); // 先确保 VPI 那条流的算法真的跑完了 vpiImageLockData(vpiDst, VPI_LOCK_READ, VPI_IMAGE_BUFFER_CUDA_PITCH_LINEAR, // ★ 以 CUDA 显存视角锁定 out); uint8_t* d_result (uint8_t*)out.buffer.pitch.planes[0].pBase; // ★ 就是显存指针 int pitch out.buffer.pitch.planes[0].pitchBytes; // 直接把这个显存指针喂给你自己的 CUDA kernel —— 全程没有一次 cudaMemcpy dim3 block(16, 16), grid((w15)/16, (h15)/16); myPostKernelgrid, block(d_result, d_tensor, w, h); cudaDeviceSynchronize(); vpiImageUnlock(vpiDst); // lock/unlock 必须配对d_result是 VPI 写好的显存myPostKernel是你手写的 kernel两者共用同一块显存中间零拷贝。之前的手写 kernel和这里的用 VPI在这里合流了。10. 黄金法则先测量再优化Dont guess, measure前面你反复听到这样更快合并访问快、shared memory 快、多流重叠快、零拷贝快、VIC 省 GPU……但这些到目前为止都还是别人告诉你的结论。下面你要学会两个测量工具Nsight Systems / Nsight Compute。亲手验证前面的结论并学会定位自己代码的瓶颈。GPU 优化最大的坑是凭直觉猜瓶颈然后瞎改。90% 的我以为的瓶颈都是错的。正确流程永远是测量→ 找到真正最慢的那一段瓶颈判断瓶颈类型访存计算搬运同步等待针对性优化那一段再测量确认真的变快了而不是自我感觉11. 三个测量层级从粗到细层级工具回答的问题何时用代码内打点cudaEvent阶段三/vpiEventElapsedTimeMillis这段/这个 kernel 花了几毫秒快速定位哪一步慢全局时间线Nsight Systems (nsys)整个程序时间都花哪了拷贝和计算重叠了吗GPU 有没有闲着先用它看全貌单 kernel 深剖Nsight Compute (ncu)这个 kernel 卡在访存还是计算占用率多少访问合并了吗锁定某个热点 kernel 后深挖正确顺序先用nsys看全局找热点 → 再用ncu剖那个最热的 kernel。别一上来就ncu它很慢。12. Nsight Systemsnsys看整条流水线的时间线12.1 它是什么把程序运行的时间线画出来CPU 在干嘛、GPU 在跑哪个 kernel、H2D/D2H 拷贝在什么时候、多个 stream 是否并行、哪里有空隙gap。这是发现搬运没重叠GPU 在干等这类问题的最佳工具。12.2 怎么用# 采集在有 GPU 的机器上生成 report.nsys-rep nsys profile --statstrue -o report ./multi_stream # 命令行直接看统计摘要--statstrue 已经会打印 # 图形界面看时间线本机也装了 nsys-ui nsys-ui report.nsys-rep12.3 重点看什么对应前面阶段的验证验证笔记三的多流重叠打开你之前写的的multi_stream在时间线上看 4 条 stream 的 H2D 拷贝、kernel、D2H 是不是在时间上叠起来了重叠成功还是一条接一条没重叠——多半是忘了用 pinned 内存。找 GPU 空隙gapGPU 时间线上大段空白 GPU 在等 CPU 或等拷贝。空隙多说明瓶颈在别处CPU、传输、同步。看 H2D/D2H 占比如果拷贝时间比 kernel 计算时间还长 → 印证搬运是头号敌人该想零拷贝/减少搬运。13. Nsight Computencu剖单个 kernel13.1 它是什么对单个 kernel 做显微镜级分析占用率、内存带宽利用、是计算受限还是访存受限、访问是否合并、warp 是否发散……用来回答这个 kernel 为什么慢、怎么改快。13.2 怎么用# 剖析全部指标很详细也很慢会把 kernel 重放很多次 ncu --set full -o report ./grayscale # 只剖某个 kernel、只测第一次启动快很多 ncu --kernel-name rgb2gray --launch-count 1 ./grayscale # 图形界面看报告 ncu-ui report.ncu-rep注意 ncu会反复重放replay同一个 kernel 来采集不同指标所以比正常运行慢几十倍属正常现象。13.3 三个最该先看的指标Achieved Occupancy实际占用率SM 上实际活跃的 warp 占理论最大的比例。太低如 30%说明并行度不够——常见原因block 太小、每线程寄存器/shared memory 用太多。先调 block 大小128/256/512试试。Memory Throughput vs Compute Throughput哪个接近上限就是瓶颈类型。访存接近上限、计算很低 →memory-bound多数图像/向量类简单 kernel 都是优化方向是访存合并、shared memory、减少读写。计算接近上限 →compute-bound优化方向是减少计算量或用更快指令。Memory Coalescing访存合并相关警告ncu会直接提示访问是否合并。14. 实战剖析笔记二的手写 kernel找瓶颈14.1 步骤ncu --set full -o gray_report ./grayscale ncu-ui gray_report.ncu-rep # 看 Occupancy 和 Memory Throughput预期发现多数简单图像 kernel 的典型结论是memory-bound访存受限——因为每个像素只做几次乘加却要读写显存计算单元大部分时间在等数据。这解释了为什么阶段二强调合并访问既然瓶颈是访存那把访存做到合并就是最有效的优化。14.2 亲手做个对比实验强烈推荐写两个版本对比最能建立直觉A合并版阶段二的正常写法threadIdx.x对应连续列。B非合并版故意让每个线程隔很远取数据制造跨越访问。用ncu分别测你会看到 B 的 Memory Throughput 明显更差、耗时更长。14.3 调 block 大小看 occupancy把threadsPerBlock从 64 → 128 → 256 → 512 分别测耗时和 occupancy观察拐点。你会发现太小的 block 占用率低、太大也不一定更好——这就是为什么阶段一说 block 取 32 的倍数、通常 128~512。15. 常见优化手段清单定位瓶颈后按需取用务必先测量定位瓶颈再从下面选对应手段不要无脑全上瓶颈类型优化手段搬运太多H2D/D2H 占比高零拷贝、减少 CPU↔GPU 传输、算法留在 GPU 串起来访存受限、非合并合并访问、用 shared memory 复用、减少冗余读写占用率低调 block 大小、降低每线程寄存器/shared 用量GPU 与 CPU/拷贝互相干等多流并发、计算与传输重叠GPU 被杂活占满、抢不到算力把预处理丢给 VIC/PVA 后端warp 分支发散减少同一 warp 内的 if/else 分歧