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

Question

我目前在我的GPU（CUDA/C++）上实现了一个运动跟踪算法，并且目前看到非常强大的加速。然而，正如人们可以预料的那样，主要的瓶颈是帧（图像）数据从CPU到GPU的实际传输。

因为我使用OpenCV来读取测试视频文件。但是，OpenCV以RRGGBB RRGGBB ...的形式将封包以字节形式返回，或者换句话说，每个像素都对齐到24位​​边界。这不允许我使用合并内存访问，这会在GPU上导致严重的性能损失。原来，我只是使用一些预先生成的测试数据，其中是 32位对齐（以RRGGBB00 RRGGBB00 ...的形式用零填充），但我现在要开始使用实际的视频数据。

这引起了我一些显著的性能损失，所以我有两个主要问题：

1. 虽然我可以手动预处理的CPU上的感兴趣像素，然后开始转移，有没有方法可以快速将像素数据传输到GPU，但对齐到32位边界？ （然而，我认为这与预处理具有相同的性能）

2. 是否有另一个库可以用来以不同格式读取视频？例如，我知道即使没有包含Alpha通道，SDL表面也会打包在32位边界中。

我们实现的最终目标将是实时与机器人控制摄像头接口，虽然现在我只是想的东西，可以高效地解码我的测试视频来测试我们的功能检测和运动跟踪具有预定义测试数据的算法。

Answer 1

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

__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; 
}