Я реализую программу 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 ядра.