4

画像を転置するための CUDA プログラムを実装しています。2 つのカーネルを作成しました。最初のカーネルは場違いな転置を行い、あらゆる画像サイズに対して完全に機能します。

次に、正方形の画像をその場で転置するためのカーネルを作成しました。ただし、出力は正しくありません。イメージの下の三角形は転置されますが、上の三角形は同じままです。結果の画像には、対角線に階段のようなパターンがあり、階段の各ステップのサイズは、カーネルに使用した 2D ブロック サイズと同じです。

場違いなカーネル:

src と dst が異なる場合、どの画像サイズでも完全に機能します。

template<typename T, int blockSize>
__global__ void kernel_transpose(T* src, T* dst, int width, int height, int srcPitch, int dstPitch)
{
    __shared__ T block[blockSize][blockSize];

    int col = blockIdx.x * blockSize + threadIdx.x;
    int row = blockIdx.y * blockSize + threadIdx.y;

    if((col < width) && (row < height))
    {
        int tid_in = row * srcPitch + col;
        block[threadIdx.y][threadIdx.x] = src[tid_in];
    }

    __syncthreads();

    col = blockIdx.y * blockSize + threadIdx.x;
    row = blockIdx.x * blockSize + threadIdx.y;

    if((col < height) && (row < width))
    {
        int tid_out = row * dstPitch + col;
        dst[tid_out] = block[threadIdx.x][threadIdx.y];
    }
}

インプレース カーネル:

template<typename T, int blockSize>
__global__ void kernel_transpose_inplace(T* srcDst, int width, int pitch)
{
    __shared__ T block[blockSize][blockSize];

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

    int tid_in = row * pitch + col;
    int tid_out = col * pitch + row;

    if((row < width) && (col < width))
        block[threadIdx.x][threadIdx.y] = srcDst[tid_in];

    __threadfence();

    if((row < width) && (col < width))
        srcDst[tid_out] = block[threadIdx.x][threadIdx.y];
}

ラッパー機能:

int transpose_8u_c1(unsigned char* pSrcDst, int width,int pitch)
{
    //pSrcDst is allocated using cudaMallocPitch

    dim3 block(16,16);
    dim3 grid;
    grid.x = (width + block.x - 1)/block.x;
    grid.y = (width + block.y - 1)/block.y;

    kernel_transpose_inplace<unsigned char,16><<<grid,block>>>(pSrcDst,width,pitch);

    assert(cudaSuccess == cudaDeviceSynchronize());

    return 1;
}

サンプル入力と間違った出力:

ここに画像の説明を入力 ここに画像の説明を入力

この問題は、インプレース転置のロジックと関係があることを知っています。これは、異なるソースと宛先に対して完全に機能している場違いな転置カーネルが、ソースと宛先に単一のポインターを渡すと、同じ間違った結果をもたらすためです。

私は何を間違っていますか？インプレース カーネルの修正を手伝ってください。

4

1 に答える 1

3

インプレース カーネルがイメージ内のデータを上書きしています。このデータは、転置操作に使用する別のスレッドによって後で取得されます。したがって、正方形の画像の場合、宛先データを上書きする前にバッファリングしてから、宛先データを適切な転置場所に配置する必要があります。この方法を使用すると、スレッドごとに実質的に 2 つのコピーを実行しているため、使用するスレッドの数は半分で済みます。このようなものが動作するはずです：

template<typename T, int blockSize>
__global__ void kernel_transpose_inplace(T* srcDst, int width, int pitch)
{

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

    int tid_in = row * pitch + col;
    int tid_out = col * pitch + row;

    if((row < width) && (col < width) && (row<col)) {

        T temp = srcDst[tid_out];

        srcDst[tid_out] = srcDst[tid_in];
        srcDst[tid_in] = temp;
        }
}
于 2013-01-05T19:50:32.907 に答える