Two things:

• This is a visualization of the following statement from an NVIDIA blog on efficient matrix transpose kernel.
  • "All kernels in this study launch blocks of 32×8 threads (TILE_DIM=32, BLOCK_ROWS=8 in the code), and each thread block transposes (or copies) a tile of size 32×32."
• The trick to avoid bank conflicts during transposing a tile and storing it in SRAM.

The tile in HBM and the threads in a block are shown here.

#define TILE_DIM    32
#define BLOCK_ROWS   8

// 32 x 8 = 256 threads per block, one block per 32 x 32 tile
dim3 dimGrid(width / TILE_DIM, height / TILE_DIM, 1);
dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);
transpose<<<dimGrid, dimBlock>>>(odata, idata);

__global__ void transposeNoBankConflicts(float *odata, const float *idata)
{
  __shared__ float tile[TILE_DIM][TILE_DIM+1];
    
  int x = blockIdx.x * TILE_DIM + threadIdx.x;
  int y = blockIdx.y * TILE_DIM + threadIdx.y;
  int width = gridDim.x * TILE_DIM;

  for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
     tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];

  __syncthreads();

  // ...
  // code to transpose and write back to HBM
  // ...
}

Transposing when writing back from SRAM to HBM

When writing back to HBM, values from columns of SRAM are read, then transposed and written back to HBM row-wise. Note that reading values from the same column can lead to worst-case bank conflicts. This is avoided by adding an extra column to SRAM.

__global__ void transposeNoBankConflicts(float *odata, const float *idata)
{
  __shared__ float tile[TILE_DIM][TILE_DIM+1];
    
  // ...
  // code to read from HBM to SRAM
  // ...

  x = blockIdx.y * TILE_DIM + threadIdx.x;  // transpose block offset
  y = blockIdx.x * TILE_DIM + threadIdx.y;

  for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
     odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
}

This figure shows the first few rows of the SRAM layout where each word shows the bank it belongs to. Note how the banks of each column are unique now.

Adapted from An Efficient Matrix Transpose in CUDA C/C++ (Mark Harris, NVIDIA).