In this post, we shall be looking at how to optimize transpose kernel for real world applications where the matrix dimensions are seldom powers of 2.

Let us start with the transpose kernel available from <a href="http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/">here</a>. It has been optimized to avoid bank conflicts as well but only works with matrices with dimensions that are multiples of 32.
<pre>
template<typename T>
__global__ 
void transpose32(T * out, const T * in, unsigned dim0, unsigned dim1)
{
    __shared__ T shrdMem[TILE_DIM][TILE_DIM+1];

    unsigned lx = threadIdx.x;
    unsigned ly = threadIdx.y;

    unsigned gx = lx + blockDim.x * blockIdx.x;
    unsigned gy = ly + TILE_DIM   * blockIdx.y;

#pragma unroll
    for (unsigned repeat = 0; repeat < TILE_DIM; repeat += blockDim.y) {
        unsigned gy_ = gy+repeat;
        shrdMem[ly + repeat][lx] = in[gy_ * dim0 + gx];
    }
    __syncthreads();

    gx = lx + blockDim.x * blockIdx.y;
    gy = ly + TILE_DIM   * blockIdx.x;

#pragma unroll
    for (unsigned repeat = 0; repeat < TILE_DIM; repeat += blockDim.y) {
        unsigned gy_ = gy+repeat;
        out[gy_ * dim0 + gx] = shrdMem[lx][ly + repeat];
    }
}
</pre>
I have done a minor modification to the kernel by templating the matrix element data type such that i can use the same kernel with other types such as double, int, unsigned etc.