2

CUDA で簡単な関数を作成しました。画像のサイズを 2 倍にリサイズします。1920*1080 の画像の場合、この関数は完了するまでに最大 20 ミリ秒かかります。その機能を最適化するために、いくつかの異なる方法を試しました。そして、ローカルメモリが主な理由である可能性があることがわかりました。

画像を取得するために3つの異なる方法を試しました。

  • OpenCV の Gpu モジュール
  • OpenCV での GpuMat へのテクスチャ バインド
  • グローバル メモリから GpuMat を直接フェッチする

それらのどれも私を少し改善することはできませんでした。

次に、nvvp を使用して理由を調べます。また、ローカル メモリのオーバーヘッドは、上記の 3 つの条件すべてで最大 95% です。

そこで、コードに目を向けて、nvcc がメモリをどのように使用しているかを調べます。次に、次のような単純な関数を見つけました。

__global__ void performDoubleImage(float* outData, size_t step, const int cols, const int rows)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    if (x >= cols)
        return;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (y >= rows)
        return;
    ((float*)((size_t)outData+y*step))[x] = tex2D(texRef, x/2, y/2);
}

80 バイトのスタック フレームが必要です (ローカル メモリにあります)。

そして、このような別の機能:

__global__ void performFinalDoubleImage(const PtrStepSz<float> in, PtrStepSz<float> out)
{
    out(out.rows-1, out.cols-1) = out(in.rows-1, in.cols-1);
}

88 バイトのスタック フレームも必要です。

問題は、なぜ私の関数がこの単純なタスクで大量のローカル メモリとレジスタを使用するのかということです。また、OpenCV の関数がローカル メモリを使用せずに同じ関数を実行できるのはなぜですか (これは nvvp によるテストで、ローカル メモリの負荷はゼロです)。

私のコードはデバッグモードでコンパイルされています。そして私のカードはGT650(192 SP/SM, 2 SM)です。

4

1 に答える 1