1

私は、sift アルゴリズムを cuda で実行する学校のプロジェクトに取り組んでいます。私はある時点で、画像のすべてのピクセル(X)のマグニチュード値を、その近傍(A、B、C、D)の値に基づいて計算しています:

   A
B  X  C
   D

入力配列から必要な値を簡単に取得できたため、グローバルメモリを使用してなんとか作成できました。

しかし、今は最初に入力配列を共有メモリに入れることでそれを作りたいと思っていますが、スレッドが正しいピクセルを共有メモリに入れる方法に本当に苦労しています。画像の境界のパディングを考慮する必要があります。

パディングが含まれるように、そこに入れたい画像の部分よりも多くの共有メモリが必要であることはわかっていますが、スレッドブロックに共有メモリスペースよりも多いスレッドまたは少ないスレッドを含める必要があるかどうか、および何を指定するかはわかりません読む。誰かがこれをどのように考えるかについての一般的なアイデアを私に与えることができれば、そこからそれを取ることができます...

ありがとう！

4

1 に答える 1

1

グレースケール画像を通過し、sobel フィルターを適用するコードを提供しました: (Sobel は、neighbor(A,B,C,D) 関数に似たフィルターです)

#define QUANTUM_TYPE short
__global__ void sobel_gpu(QUANTUM_TYPE *img_out, QUANTUM_TYPE *img_in, int WIDTH, int HEIGHT){
    int x,y;
    x=blockDim.x*blockIdx.x+threadIdx.x;
    y=blockDim.y*blockIdx.y+threadIdx.y;
    QUANTUM_TYPE LUp,LCnt,LDw,RUp,RCnt,RDw;
    int pixel;

    if(x<WIDTH && y<HEIGHT){
        LUp = (x-1>=0 && y-1>=0)? img_in[(x-1)+(y-1)*WIDTH]:0;
        LCnt= (x-1>=0)? img_in[(x-1)+y*WIDTH]:0;
        LDw = (x-1>=0 && y+1<HEIGHT)? img_in[(x-1)+(y+1)*WIDTH]:0;
        RUp = (x+1<WIDTH && y-1>=0)? img_in[(x+1)+(y-1)*WIDTH]:0;
        RCnt= (x+1<WIDTH)? img_in[(x+1)+y*WIDTH]:0;
        RDw = (x+1<WIDTH && y+1<HEIGHT)? img_in[(x+1)+(y+1)*WIDTH]:0;
        pixel = -1*LUp  + 1*RUp +
                -2*LCnt + 2*RCnt +
                -1*LDw  + 1*RDw;
        pixel=(pixel<0)?0:pixel;
        pixel=(pixel>MAXRGB)?MAXRGB:pixel;
        img_out[x+y*WIDTH]=pixel;
    }
}

コードはグローバル メモリに対して機能し、境界を安全に処理します。私の完全なコードは、BMP 画像を読み取り、それにフィルターを適用して、結果の BMP をディスクに保存します。こちらから入手できます(Linux と Windows の両方で、CPU と GPU の実装が統合されています)。

ちょっとした作業で共有メモリスタイルに変えることができます。まず、各ブロックに与えるタスクの量を決定する必要があります。次に、タスクを複数の共有メモリ シンク/ドレインに分割します。CUDA SDK の行列乗算の例は、完璧なアイデアを提供します。

于 2012-10-09T08:25:48.357 に答える