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 on 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>
A minor modification has been done 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.

One immediate thing we can do to make this kernel applicable to matrices that have dimensions which are not powers of 2 is to add index checking. Following are the changes required to do so.
<pre>
template<typename T>
__global__
void transposeNoSC(T * out, const T * in, unsigned dim0, unsigned dim1)
{
    /*
     * same code as transpose32 kernel
     */
#pragma unroll
    for (unsigned repeat = 0; repeat < TILE_DIM; repeat += blockDim.y) {
        unsigned gy_ = gy+repeat;
        if (gx<dim0 && gy_<dim1)
            shrdMem[ly + repeat][lx] = in[gy_ * dim0 + gx];
    }
    __syncthreads();
    /*
     * same code as transpose32 kernel
     */
#pragma unroll
    for (unsigned repeat = 0; repeat < TILE_DIM; repeat += blockDim.y) {
        unsigned gy_ = gy+repeat;
        if (gx<dim1 && gy_<dim0)
            out[gy_ * dim0 + gx] = shrdMem[lx][ly + repeat];
    }
}
</pre>