GPU与图像采集卡编程实战

eGrabber + CUDA 图像采集：三种内存模式技术解析

1. 背景

在工业视觉应用中，图像采集到 GPU 处理是一条完整的数据流水线。eGrabber 负责从采集卡获取图像数据，CUDA 负责在 GPU 上做图像处理。但图像数据最初落在哪里、怎么到达 GPU，直接决定了整条流水线的性能上限。

cpp/nvidia-cuda/示例提供了三种内存模式，供开发者根据自身场景选择最优方案。

2. 缓冲区模型：N 个独立缓冲区

eGrabber 采用的是多个独立缓冲区循环使用的模型，而非一块大内存存储多帧。

┌──────────┐ ┌──────────┐ ┌──────────┐ │ Buffer 0 │ │ Buffer 1 │ │ Buffer 2 │ ← N 个独立缓冲区 │ W×H 像素 │ │ W×H 像素 │ │ W×H 像素 │ 每个大小 = width × height × 像素深度 └────┬─────┘ └────┬─────┘ └────┬─────┘ │ │ │ └──────┬──────┘──────┬──────┘ ↓ ↓ eGrabber 采集 应用处理（显示/保存/分析）

以本示例为例（NB_BUFFERS = 3）：

// main.cpp L19#defineNB_BUFFERS3// main.cpp L106std::vector<unsignedchar*>pinnedMemory(NB_BUFFERS);// 3 个指针// cuda.cpp L60-61size_t size=grabber.getWidth()*grabber.getHeight();// 每帧大小 = W × H（Mono8）for(size_t i=0;i<pinnedMemory.size();++i){// 循环 3 次// 每次分配一个帧大小的缓冲区，独立注册给 eGrabbergrabber.announceAndQueue(UserMemory(ptr,size,devicePtr));}

为什么不用一块大内存？

• 流水线并行：采集卡正在写入 Buffer 2 时，GPU 可以同时处理 Buffer 0，CPU 可以回收 Buffer 1。多缓冲区是采集与处理并行的基础。
• eGrabber 队列机制：缓冲区被放入输入队列等待采集 → 采集完成后进入输出队列 → 应用取出处理 → 处理完放回输入队列。多个缓冲区保证队列不断流。
• 避免数据覆盖：如果只有一块大内存，采集卡写入新帧时会覆盖尚未处理完的旧帧。

缓冲区数量如何选择？

3. 三种模式详解

3.1 模式一：Pinned Host Memory（默认）

适用场景：通用方案，Windows/Linux 均可使用，适合大多数应用。

核心思路

在 Host 内存上采集，然后一次性 DMA 拷贝到 GPU 本地显存，后续处理全部在 GPU 显存上进行。

数据流

采集卡 ──DMA──→ [Host Pinned 内存] ──PCIe DMA──→ [GPU cudaMem] ──Kernel处理──→ [OpenGL 纹理] Host DRAM ~6 GB/s GPU VRAM GPU HBM全速 GPU 内部

代码拆解

Step 1：分配缓冲区

// cuda.cpp L73-77cudaHostAlloc(&ptr,size,cudaHostAllocMapped);// 分配锁页内存（不可被OS交换到磁盘）cudaHostGetDevicePointer(&devicePtr,ptr,0);// 获取 GPU 可访问的映射地址grabber.announceAndQueue(UserMemory(ptr,size,devicePtr));

• ptr：CPU 虚拟地址，采集卡 DMA 引擎将数据写入此地址对应的物理内存
• devicePtr：GPU 虚拟地址，指向同一块物理内存（通过 PCIe 访问）
• cudaHostAllocMapped：关键标志，告诉 CUDA 驱动为这块内存建立 GPU 页表映射

ptr 与 devicePtr 的关系：它们是同一块物理内存的两个不同虚拟地址。CPU 和 GPU 各有独立的 MMU（内存管理单元），需要各自的页表来翻译到同一块物理页。GPU 不能直接使用ptr，因为 GPU 页表中查不到 CPU 虚拟地址。

Step 2：分配 GPU 工作缓冲区

// cuda.cpp L96-97cudaMalloc((void**)&cudaMem,width*height);// 在 GPU VRAM 上额外分配一块工作内存

Step 3：采集后拷贝到 GPU 再处理

// kernel.cu L33-34cudaMemcpy(cudaMem,devicePtr,size,cudaMemcpyDeviceToDevice);// 将数据从 Host 内存一次性拷贝到 GPU 显存（走 PCIe，约 6 GB/s）cudaProcess<<<gdim,bdim>>>(cudaMem,size);// 后续 Kernel 在 GPU 显存上执行，享受 HBM 全速带宽（~900 GB/s）

Step 4：结果写入 OpenGL 纹理

// kernel.cu L41cudaMemcpyToArray(cudaTexture,0,0,cudaMem,size,cudaMemcpyDeviceToDevice);// GPU 内部拷贝，约 14 GB/s

性能特点

• PCIe 仅使用 2 次：采集 DMA 写入 Host + Host→GPU 拷贝
• Kernel 处理完全在 GPU VRAM 上，无 PCIe 瓶颈
• 适合对处理性能有要求、但不需要极致延迟的场景

3.2 模式二：Host Memory 直接处理（cudaHostMemory）

适用场景：GPU 显存紧张、或图像处理逻辑极轻量时，省掉一次显存拷贝。

核心思路

省去 GPU 上的cudaMem工作缓冲区，CUDA Kernel 直接通过映射地址读写 Host 内存。

数据流

采集卡 ──DMA──→ [Host Pinned 内存] Host DRAM ↓ Kernel 直接读写（每次访问都走 PCIe） ↓ 结果直接拷贝到 OpenGL 纹理

代码差异

Step 2 跳过：不分配cudaMem。

// cuda.cpp L96-97if(!options[CUDA_HOSTMEMORY]){// CUDA_HOSTMEMORY == true，不执行cudaMalloc(&cudaMem,...);}

Step 3 Kernel 直接操作 Host 内存：

// kernel.cu L29-30if(params.useHostMemory||params.useRDMA){// useHostMemory == truecudaProcess<<<gdim,bdim>>>(params.devicePtr,params.size);// devicePtr 指向 Host DRAM，Kernel 每次读写都经过 PCIe}

性能特点

• 省了cudaMem这块 GPU 显存，也省了一次 Host→GPU 的整体拷贝
• 但代价更隐蔽：Kernel 中每个像素的读写都是独立的 PCIe 事务
  • 一次*p = 255 - *p操作 = 1 次 PCIe 读 + 1 次 PCIe 写
  • PCIe 延迟约 ~1μs/次，远大于 HBM 的 ~纳秒级
• 当图像处理 Kernel 的访存模式是随机读写或多次遍历时，性能会急剧下降
• 仅适合极其简单的逐像素操作，且 GPU 显存确实不够用的情况

实际建议：除非 GPU 显存实在分配不出额外空间，否则不推荐此模式。模式一的"一次拷贝 + 全速处理"在绝大多数场景下总耗时更短。

3.3 模式三：NVIDIA RDMA（cudaRDMA，仅 Linux）

适用场景：追求极致性能的高帧率工业检测、实时 AI 推理管线。

核心思路

缓冲区直接分配在 GPU VRAM 上，采集卡通过 NVIDIA RDMA（Remote Direct Memory Access）将图像数据直接写入 GPU 显存，完全绕过 CPU 和 Host 内存。

数据流

采集卡 ──RDMA──→ [GPU VRAM] ──Kernel处理──→ [OpenGL 纹理] 直接写入 GPU HBM全速 GPU 内部 零拷贝

代码拆解

Step 1：在 GPU VRAM 上分配缓冲区

// cuda.cpp L63-72cudaMalloc(&ptr,size);// 直接在 GPU VRAM 上分配cuPointerSetAttribute(&flag,CU_POINTER_ATTRIBUTE_SYNC_MEMOPS,...);// 确保 DMA 写入和 GPU 读取之间的内存操作同步devicePtr=ptr;// 同一个地址！VRAM 上的地址对 GPU 天然可见grabber.announceAndQueue(NvidiaRdmaMemory(ptr,size,devicePtr));// ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^// 用 NvidiaRdmaMemory 而非 UserMemory// 告诉 eGrabber 驱动："请通过 NVIDIA RDMA 直接写入 GPU 显存"

与模式一/二的关键区别：

• ptr == devicePtr：内存在 GPU VRAM 上，GPU 访问无需任何映射
• 使用NvidiaRdmaMemory注册：这是 eGrabber 提供的专用类型，触发采集卡驱动的 RDMA 路径
• 采集卡的 DMA 引擎通过 NVIDIA 提供的 GPUDirect RDMA 接口直接寻址 GPU 物理内存

Step 3 Kernel 在 VRAM 上原位处理：

// kernel.cu L29-30cudaProcess<<<gdim,bdim>>>(params.devicePtr,params.size);// 直接在 GPU VRAM 上执行，HBM 全速带宽

性能特点

• 全程零 PCIe 瓶颈（采集直接到 GPU，处理在 GPU 内部）
• 零 CPU 参与（CPU 只负责调度和控制流）
• 最低延迟，最高吞吐

使用限制

4. 三种模式对比

5. 选型指南

需要跨平台（Windows + Linux）？ ├── 是 → GPU 显存充足？ │ ├── 是 → 模式一：Pinned Host Memory（推荐默认） │ └── 否 → 模式二：HostMemory 直接处理（不推荐，仅作备选） └── 否（仅 Linux）→ 追求极致性能？ ├── 是 → GPU 支持 RDMA？ │ ├── 是 → 模式三：NVIDIA RDMA │ └── 否 → 模式一：Pinned Host Memory └── 否 → 模式一：Pinned Host Memory

让我从代码层面再深入验证一下，同时看看标准模式和自定义内存模式的对比。

答案：N 个独立缓冲区，每个存一帧

eGrabber 的缓冲区模型是：分配 N 个独立的缓冲区，每个缓冲区大小 = 一帧图像。不存在"一个大缓冲区存 N 帧"的设计。

代码证据

标准模式（eGrabber 自动管理内存）：

// 100-grabn.cpp L15grabber.reallocBuffers(20);// 申请 20 个独立缓冲区，每个 = 1 帧大小// eGrabber 内部会循环调用 GenTLMemory() 分配 20 块独立内存

自定义内存模式（nvidia-cuda 示例）：

// cuda.cpp L60-77size_t size=grabber.getWidth()*grabber.getHeight();// 每帧大小for(size_t i=0;i<pinnedMemory.size();++i){// 循环 3 次cudaHostAlloc(&ptr,size,cudaHostAllocMapped);// 每次分配 1 帧grabber.announceAndQueue(UserMemory(ptr,size,devicePtr));// 逐个注册}

更高级的用法（221-queue-buffer-ranges）甚至可以声明多组不同配置的缓冲区：

// 第一组：10 个缓冲区BufferIndexRangebufferSet1(grabber.announceAndQueue(GenTLMemory(),10));// 第二组：20 个缓冲区（可以有不同的采集参数）BufferIndexRangebufferSet2(grabber.announceAndQueue(GenTLMemory(),20));

为什么必须是独立的 N 个缓冲区？

因为 eGrabber 内部维护的是一个缓冲区队列，采集和处理是并行的：

时刻 T1: [输入队列] [正在采集] [输出队列] [正在处理] Buffer 2 Buffer 0 Buffer 1 - ↑采集卡写入 (还没取) 时刻 T2: [输入队列] [正在采集] [输出队列] [正在处理] Buffer 0 Buffer 2 - Buffer 1 (处理完归还) ↑采集卡写入 ↑应用读取 时刻 T3: [输入队列] [正在采集] [输出队列] [正在处理] - Buffer 2 Buffer 0 Buffer 1 (新帧到达) (还在处理)

如果用一块大内存：

• 采集卡写入帧 2 时，帧 0 还没处理完 →数据被覆盖
• 无法实现"采集和处理并行"
• 队列机制无法工作（队列管理的是离散的 buffer 对象）

缓冲区生命周期的 4 个状态

┌──────────┐ announce ┌──────────┐ start ┌──────────┐ │ 未声明 │ ────&Queue────→│ 输入队列 │ ──采集完成──→│ 输出队列 │ │ (刚分配) │ │ (等待填充) │ │ (等待处理) │ └──────────┘ └──────────┘ └─────┬────┘ ↑ │ │ push 归还 │ pop 取出 └───────────────────────────────────────────────────────┘

• announce: 向 GenTL 驱动声明"我有这块内存可以用来采集"
• Queue: 放入输入队列等待采集卡填充
• 采集完成: 自动移入输出队列，等待应用取走
• pop: 应用取出处理
• push: 处理完后归还，重新放入输入队列

每个缓冲区在这 4 个状态之间独立循环，互不干扰。这就是为什么必须是 N 个独立缓冲区。

缓冲区数量怎么选？

关键原则：缓冲区数 = 你希望同时"在途"（在采集 + 在队列 + 在处理）的最大帧数。

概念说明：

OpenGL 纹理（Texture）是什么

简单来说，纹理就是一块存在 GPU 显存里的图像数据，可以被 GPU 高效读取和渲染到屏幕上。

为什么需要纹理？

GPU 渲染图像时，需要把像素数据"贴"到一个几何图形上（比如一个矩形面片）。但像素数据不能直接放在普通的内存数组里让 GPU 去渲染——GPU 有专门的硬件单元来处理纹理采样，它要求数据存放在纹理对象中。

普通数组: GPU 可以读，但渲染管线不认 纹理对象: GPU 渲染管线专用的数据格式，硬件加速采样

在 nvidia-cuda 示例中的角色

整个显示流程是这样的：

相机图像 → CUDA 处理 → 写入纹理 → GPU 渲染管线 → 屏幕像素 ↑ 这一步必须经过纹理 因为 OpenGL 渲染管线 从纹理中读取像素

具体代码对应（[opengl.cpp](file:///d:/Git/gentl/egrabber-sample-programs/cpp/nvidia-cuda/src/opengl.cpp#L164-L193)）：

// 1. 创建纹理对象（在 GPU 显存中分配一块图像存储空间）glGenTextures(1,&tex);glBindTexture(GL_TEXTURE_2D,tex);glTexImage2D(GL_TEXTURE_2D,0,GL_LUMINANCE,width,height,0,// 宽 × 高GL_LUMINANCE,GL_UNSIGNED_BYTE,// 8-bit 灰度图NULL);// 初始无数据// 2. 创建一个四边形（quad），作为"画布"glGenBuffers(1,&quad);// 4 个顶点 + 纹理坐标，构成一个矩形面片// 3. 渲染时：把纹理贴到四边形上glBindTexture(GL_TEXTURE_2D,tex);// 绑定纹理glDrawArrays(GL_QUADS,0,4);// 画四边形，GPU 自动从纹理取像素glutSwapBuffers();// 刷到屏幕

纹理 vs 普通内存的区别

类比

把纹理理解为GPU 的"专用画布"：

• 你的图像数据是"颜料"
• CUDA Kernel 负责调颜料（图像处理）
• 纹理就是那块"画布"——GPU 渲染引擎只认识画布上的画，不认散装颜料
• cudaMemcpyToArray就是把调好的颜料刷到画布上的动作

所以在 nvidia-cuda 示例中，不管哪种 DMA 模式，最后一步都是：

// kernel.cu — 把处理后的图像数据写入纹理cudaMemcpyToArray(cudaTexture,0,0,...);

这是为了让 OpenGL 渲染管线能拿到处理后的图像，显示到屏幕上。如果不需要显示（比如只做离线分析），这步就可以省掉。