3

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

照原样，我正在使用 OpenCV 读取测试视频文件。然而，OpenCV 以压缩字节的形式返回帧RRGGBB RRGGBB ...，或者换句话说，每个像素都与 24 位边界对齐。这使我无法使用合并的内存访问，这会对 GPU 造成严重的性能损失。按原样，我只是使用一些预先生成的32 位对齐的测试数据（在表单中用零填充RRGGBB00 RRGGBB00 ...），但我现在想开始使用实际的视频数据。

这给我带来了一些重大的性能损失，所以我有两个主要问题：

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

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

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

4

1 回答 1

2

我尝试编写一个简单的 CUDA 内核，使用共享内存将 24 位值填充到 32 位。请注意，这不是一个非常整洁的代码（仅适用于 1 个块，依赖于 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;
}
于 2013-04-03T20:11:21.367 回答