c++ - 将视频帧数据移动到 GPU 的最有效方法是什么？

Question

我目前正在我的 GPU (CUDA/C++) 上实现运动跟踪算法，到目前为止我看到了非常强劲的加速。然而，正如人们可能预料的那样，主要瓶颈是帧(图像)数据从 CPU 到 GPU 的实际传输。

照原样，我正在使用 OpenCV 读取测试视频文件。然而，OpenCV 以 RRGGBB RRGGBB ... 的形式将帧作为打包字节返回，或者换句话说，每个像素都与 24 位边界对齐。这不允许我使用 union 内存访问，这会对 GPU 造成严重的性能损失。按原样，我只是使用一些预先生成的测试数据，这些数据 32 位对齐(以 RRGGBB00 RRGGBB00 ... 形式填充零)，但是我现在想开始使用实际的视频数据。

这导致了一些严重的性能损失，所以我有两个主要问题:

1. 虽然我可以在 CPU 上手动预处理感兴趣的像素然后启动传输，但有没有什么方法可以快速将像素数据传输到 GPU，而不是对齐到 32 位边界？ (不过，我认为这与预处理具有相同的性能影响)

2. 我可以使用另一个库来读取不同格式的视频吗？例如，我知道 SDL 表面被打包在 32 位边界内，即使不包含 alpha channel 也是如此。

我们实现的最终目标是与用于机器人控制的摄像头进行实时交互，尽管现在我只想要一些可以有效解码我的测试视频的东西，以测试我们的特征检测和运动跟踪算法定义测试数据。

Answer 1

我尝试编写一个简单的 CUDA 内核，使用共享内存将 24 位值填充为 32 位值。请注意，这不是一个非常整洁的代码(仅适用于 1 个 block ，依赖 int 为 32 位)- 小心使用。我尝试了一个有共享内存原子和没有共享内存原子的版本 - 似乎工作。:

__global__ void pad(unsigned int *data, unsigned int* odata) {
__shared__ unsigned int array[WORK_SIZE];
unsigned int v, high, low;
const int index = (threadIdx.x * sizeof(unsigned int)) / 3;

array[threadIdx.x] = 0;
__syncthreads();

const int shl = threadIdx.x % 3;
const int shr = 3 - shl;

if (threadIdx.x
        < ((WORK_SIZE * 3) + sizeof(unsigned int) - 1)
                / sizeof(unsigned int)) {
    v = data[threadIdx.x];
    high = (v >> (shl * 8)) & ~0xFF;
    low = v << (shr * 8);
#if __CUDA_ARCH__ < 200
    array[index] = high;
}
__syncthreads();
if (threadIdx.x
        < ((WORK_SIZE * 3) + sizeof(unsigned int) - 1)
        / sizeof(unsigned int)) {
    array[index + 1] += low;
#else
    if (high)
        atomicOr(array + index, high);
    if (low)
        atomicOr(array + 1 + index, low);
#endif
}
__syncthreads();

// Do computations!
odata[threadIdx.x] = array[threadIdx.x] + 0xFF;
}