CUDA原地转置错误

Question

CUDA原地转置错误

cudatransposein-place

4

我正在实现一个用于转置图像的CUDA程序。我创建了两个内核。第一个内核用于离线转置，对于任何图像大小都能完美工作。

然后我创建了一个用于原地转置正方形图像的内核。但是输出结果是不正确的。图像的下三角被转置了，但上三角保持不变。结果图像在对角线上呈现出阶梯状的图案，每一步的大小都等于我内核所使用的二维块的大小。

离线内核:

如果源(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];
}

Wrapper Function:

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;
}

示例输入和错误输出：

enter image description here

我知道这个问题与原地转置逻辑有关。这是因为我的非原地转置核心在不同的源和目标上工作得非常完美，但如果我将单个指针用作源和目标并传递给它，则也会给出相同的错误结果。

我做错了什么？请帮助我修正原地转置核心。

- sgarizvi

你能否添加图片，包括问题出现前、纠正后和错误纠正后的图片。这有助于问题的可视化。此外，如果您可以包含不在适当位置的内核的代码，那就更好了。 - 1-----1

@ks6g10... 已添加两个。嗯...“after-correct”是不是很明显呢？ - sgarizvi

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

你的原地内核正在覆盖图像中将要被另一个线程用于转置操作的数据。因此，对于方形图像，在覆盖之前应该缓冲目标数据，然后将目标数据放在其适当的转置位置上。由于使用这种方法每个线程有效地进行了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;
        }
}