CUDA 就地转置错误

Question

我正在实施一个用于转置图像的 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;
}

示例输入和错误输出：

我知道这个问题与 in-place 转置的逻辑有关。这是因为我的不合适的转置内核可以完美地用于不同的源和目标，如果我将源和目标的单个指针传递给它也会给出相同的错误结果。

我做错了什么？帮助我纠正就地内核。

Answer 1

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