1

シーケンシャルスムージングアルゴリズムがあります

void triangularSmooth(unsigned char *grayImage, unsigned char *smoothImage, const int width, const int height, const float *filter, NSTimer &timer, dim3 grid_size, dim3 block_size) {
for ( int y = 0; y < height; y++ ) {
    for ( int x = 0; x < width; x++ ) {
        unsigned int filterItem = 0;
        float filterSum = 0.0f;
        float smoothPix = 0.0f;

        for ( int fy = y - 2; fy < y + 3; fy++ ) {
            for ( int fx = x - 2; fx < x + 3; fx++ ) {
                if ( ((fy < 0) || (fy >= height)) || ((fx < 0) || (fx >= width)) ) {
                    filterItem++;
                    continue;
                }

                smoothPix += grayImage[(fy * width) + fx] * filter[filterItem];
                filterSum += filter[filterItem];
                filterItem++;
            }
        }

        smoothPix /= filterSum;
        smoothImage[(y * width) + x] = static_cast< unsigned char >(smoothPix);
    }
}
}

私はCUDAで実装しており、grayImageのピクセルを保持するために共有変数を使用したいと考えています。ただし、その前はそのまま実行しようとしています。この目的のために、私はカーネルコードを持っています：

__global__ void smooth(unsigned char *grayImage, unsigned char *smoothImage, const int width, const int height, const float *filter)
{

        int x = blockIdx.x*blockDim.x + threadIdx.x;
        int y = blockIdx.y*blockDim.y + threadIdx.y;

        unsigned int filterItem = 0;
        float filterSum = 0.0f;
        float smoothPix = 0.0f;

        for ( int fy = y - 2; fy < y + 3; fy++ ) {
            for ( int fx = x - 2; fx < x + 3; fx++ ) {
                if ( ((fy < 0) || (fy >= height)) || ((fx < 0) || (fx >= width)) ) {
                    filterItem++;
                    continue;
                }

                smoothPix += grayImage[(fy * width) + fx] * filter[filterItem];
                filterSum += filter[filterItem];
                filterItem++;
            }
        }
        smoothPix /= filterSum;
        smoothImage[(y * width) + x] = static_cast< unsigned char >(smoothPix);
}

そして、次のように呼び出します。

const float filter[] = {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 2.0f, 2.0f, 2.0f, 1.0f, 1.0f, 2.0f, 3.0f, 2.0f, 1.0f, 1.0f, 2.0f, 2.0f, 2.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f};
dim3 gridSize((width*height)/1024,(width*height)/1024,1);
dim3 blockSize(256,256,1);
smooth <<< gridSize, blockSize >>> (grayImage, smoothImage, width, height, filter);
cudaDeviceSynchronize();

問題は、ピクセルのように見える滑らかな画像がすべて間違っている（混同されている）ことです。これはグリッドとブロックの寸法からですか？私は他の可能な次元の多くを試しました。正しい方法は何でしょうか？

私はGTX480、バージョン-2.x、スレッドブロックのグリッドの最大次元-3、スレッドブロックのグリッドの最大x-、y-、またはz-次元-65535、ブロックあたりのスレッドの最大数-を使用しています1024

4

2 に答える 2

1

画像フィルタリングに関連するこの回答を見ると、次のように画像のブロックとグリッドを作成することをお勧めします。

dim3 blockSize(16,16,1);
dim3 gridSize((width + blockSize.x - 1)/blockSize.x,(height + blockSize.y - 1)/blockSize.y,1);

よくあるもう1つの間違いは、カーネルに渡すフィルター配列がホストに割り当てられていることです。デバイス上に同じサイズの配列を作成し、係数をホストからデバイスにコピーします。そのデバイス配列をカーネルに渡します。

また、各スレッドで何度も合計を計算するのではなく、ホスト側でフィルター係数の合計を計算し、それを引数としてカーネルに渡すことを強くお勧めします。

境界条件により、範囲外のメモリアクセスが発生する可能性があります。カーネルで境界条件を明示的に処理します。または、境界条件が自動的に処理されるように、入力画像にCUDAテクスチャを使用するのが簡単な方法です。

于 2013-02-07T08:23:00.900 に答える
1

まず、寸法が完全に無効です。この場合、以下が機能するはずです。

dim3 blockSize(16, 16, 1);
dim3 gridSize((width + blockSize.x - 1)/ blockSize.x, (height + blockSize.y - 1) / blockSize.y, 1);
smooth <<< grid_size, block_size >>> (grayImage, smoothImage, width, height);

修正後、cuda-memcheckを使用すると、次のような結果が得られました。

========= Invalid __global__ read of size 4
=========     at 0x00000120 in cudaFilter
=========     by thread (4,1,0) in block (1,0,0)
=========     Address 0x05100190 is out of bounds

これは、カーネルコード内の値が範囲外であることを示しています（ほとんどの場合、配列インデックス）。さまざまな変数をチェックすると、filter[]が空であることがわかりました。

最後に、filter []をカーネルに渡す場合は、次のようなものを使用してCPUからGPUにコピーする必要があります。

cudaMemcpy(filterGpu, filter, 25 * sizeof(float), cudaMemcpyHostToDevice);

あるいは、フィルターが他の場所で必要ない場合（ここでの場合のように）、代わりにカーネル内で宣言することができます。

于 2013-02-07T09:14:13.460 に答える