3

現在、GPU (CUDA/C++) にモーション トラッキング アルゴリズムを実装しており、これまでのところ非常に高速化されています。ただし、おそらく予想できるように、主なボトルネックは、CPU から GPU へのフレーム (画像) データの実際の転送です。

そのまま、OpenCV を使用してテスト ビデオ ファイルを読み込みます。ただし、OpenCV は、フレームを の形式RRGGBB RRGGBB ...でパックされたバイトとして返します。つまり、各ピクセルは 24 ビット境界に揃えられます。これにより、結合されたメモリ アクセスを使用できなくなり、GPU のパフォーマンスが大幅に低下します。現状では、 32 ビットで整列された事前生成されたテスト データを使用しているだけですが(フォームにゼロが埋め込まれていますRRGGBB00 RRGGBB00 ...)、実際のビデオ データを今すぐ使用したいと思います。

これにより、パフォーマンスが大幅に低下するため、主な質問が 2 つあります。

  1. CPU で対象のピクセルを手動で前処理して転送を開始することはできますが、ピクセル データを GPU にすばやく転送し、代わりに 32 ビット境界に合わせる方法はありますか? (ただし、これは前処理と同じパフォーマンス ヒットがあると思います)

  2. ビデオを別の形式で読み込むために使用できる別のライブラリはありますか? たとえば、アルファ チャネルが含まれていなくても、SDL サーフェスが 32 ビット境界でパックされていることはわかっています。

実装の最終目標は、ロボット制御用のカメラとリアルタイムでインターフェースすることですが、今のところは、テスト ビデオを効率的にデコードして、事前定義されたテスト データを使用して機能検出とモーション トラッキング アルゴリズムをテストできるものが必要です。 .

4

1 に答える 1

2

共有メモリを使用して 24 ビット値を 32 ビットにパディングする単純な CUDA カーネルを作成してみました。これは非常にきちんとしたコードではないことに注意してください (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 に答える