Matrix Transpose: Pad Away Shared-Memory Bank Conflicts

This transpose already coalesces both global-memory paths through a shared tile. The remaining bottleneck is inside shared memory: reading a column from tile[32][32] makes a warp touch 32 distinct words with the same bank index. Add one padding column so successive rows rotate through successive banks without changing the logical 32 by 32 tile.

Languages: CUDA C++, Python, Rust

Lesson

A tiled transpose fixes global coalescing by loading rows and later reading the tile in the opposite orientation. With a physical row stride of 32 floats, a warp reading a column accesses word indexes separated by exactly 32. Those are distinct addresses in the same bank, so the shared-memory request is serialized. Declaring tile[32][33] keeps the same 32 by 32 logical data but makes the physical row stride relatively prime to 32; successive rows rotate across all banks. This is a layout fix, not an algorithm change.

Padding changes the physical shared-memory stride, not the logical matrix shape. Compare both bank serialization and runtime; fewer conflicts alone do not quantify speedup.

What you will learn

Identify the shared-memory column read that remains after global accesses are coalesced.
Derive why a 32-word row stride maps one column to one bank.
Use a one-column pad to rotate column accesses across all 32 banks.
Separate logical tile dimensions from physical shared-memory layout.

Starter code

    // Global reads and writes are coalesced, but the transposed shared-memory
    // read below walks a column of a 32-word-stride tile.
    __shared__ float tile[32][32];

    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    if (x < N && y < N) {
        tile[threadIdx.y][threadIdx.x] = A[y * N + x];
    }
    __syncthreads();

    int outX = blockIdx.y * 32 + threadIdx.x;
    int outY = blockIdx.x * 32 + threadIdx.y;
    if (outX < N && outY < N) {
        C[outY * N + outX] = tile[threadIdx.x][threadIdx.y];
    }