100 DAYS OF GPU Course preview
Day 12Module 2
in-technical-review

Why transposing a matrix is slow

This matrix transpose assigns one element, one read, and one write to each thread. Every thread stays active, and the kernel has no arithmetic, branch, or barrier. It moves as many bytes as a straight copy of the same matrix.

out[col * width + row] = in[row * width + col];

It will not run at copy speed because its output addresses are strided.

NVIDIA's own programming guide uses this kernel to make the point and says which half is at fault: "the writing of c is not coalesced, because consecutive values of threadIdx.x ... are writing elements to c that are ld (leading dimension) elements apart from each other." (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html section 2.3.4.1, checked 2026-08-30.)

This page compares the transpose with a copy that uses the same data and launch. It then shows why moving the strided access from the write to the read does not remove the cost.

One warp reads a row and writes a column

A warp is 32 threads issuing one instruction together, and threadIdx.x is the fastest-varying index inside a block. So the 32 lanes of a warp hold 32 consecutive values of col on one row of the matrix.

Take the read first. in[row * width + col], with row fixed across the warp and col running over 32 consecutive values, covers 128 contiguous bytes of global memory. That is four 32-byte sectors, so four transactions.

Every byte fetched gets used. Day 11 measured this case.

Now the write. out[col * width + row] puts the same 32 lanes width elements apart, because col is now the leading index. On a 4096-wide matrix, adjacent lanes are 16,384 bytes apart.

Every lane uses its own sector. The warp therefore asks for 32 sectors to provide 128 bytes.

The transpose reads four sectors and writes 32, for 36 sectors per warp. A copy of the same elements uses four sectors for each side, for a total of eight. The transpose therefore requests 4.5 times as many sectors before cache and DRAM effects.

This is a model, not a measurement. Day 11 found the predicted sector counts, but bandwidth kept falling after the transaction count stopped rising. Use the 36-to-8 ratio to predict the trend, not the measured result.

One warp touches 8 sectors for a copy but 36 for either transpose: 4 on the contiguous side and 32 on the side whose addresses are 16 KiB apart.

The fix everyone tries first

The obvious repair is to move the problem. Swap the two subscripts and the strided side becomes the read instead of the write:

out[row * width + col] = in[col * width + row];

The output and element count stay the same, but each warp now writes 128 contiguous bytes. This moves the strided access to the read.

The new form reads 32 sectors and writes four, still 36 in total. A transpose maps a contiguous run in one matrix to addresses spaced width apart in the other. When one thread owns one element, either the read or the write must be strided.

Move the reorder into the block. Load a square tile with coalesced reads, transpose it in shared memory, then use coalesced writes. Day 13 adds the tile and barrier.

Measuring it against something honest

The full program is in code/day12-transpose/transpose.cu. It follows three benchmark rules.

The program measures its own copy limit. Day 11's coalesced copy reached 232.9 GB/s, but it used a 1D grid, 256 threads per block, and a 256 MiB buffer. Comparing that result with this transpose would also include the block shape, grid shape, buffer size, and clock state.

The copy row here uses the same launch as the transpose. This isolates the address pattern.

Every kernel moves the same bytes from the same launch. All three read and write 16,777,216 floats using one grid and block shape. Only the swapped subscript differs. Day 11 shows how changing the byte count can reverse a benchmark result.

Correctness is checked exactly before timing. These kernels move floats without computing on them, so the check needs no tolerance. A tolerance could hide an index error.

The program first fills the output buffer with 0xFF bytes, a NaN pattern. Any element that a kernel misses will fail instead of retaining an earlier result.

Here are the three kernels. The ceiling:

__global__ void copyRows(const float* __restrict__ in, float* __restrict__ out,
                         size_t width) {
    const size_t col =
        blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const size_t row =
        blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
    out[row * width + col] = in[row * width + col];
}

The transpose, with the strided write:

__global__ void transposeNaive(const float* __restrict__ in,
                               float* __restrict__ out, size_t width) {
    const size_t col =
        blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const size_t row =
        blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
    out[col * width + row] = in[row * width + col];
}

And the same transpose with the strided side moved to the read:

__global__ void transposeStridedRead(const float* __restrict__ in,
                                     float* __restrict__ out, size_t width) {
    const size_t col =
        blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const size_t row =
        blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
    out[row * width + col] = in[col * width + row];
}

Note. The block is 32 by 32, which is 1024 threads and not the course default of 256. The tile shape forces it: 32 in x is one warp per block row, which is what makes one side of each kernel perfectly coalesced. The matrix is 4096 wide, so the grid tiles it exactly and no kernel here needs a bounds check. Six static_asserts say so at compile time, which is where a fact that never consults the device belongs.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5)
Matrix: 4096 x 4096 floats, 64 MiB per buffer
Block: 32 x 32 threads. Mean of 10 runs after 3 warm-ups.

kernel                              time (ms)       GB/s   of copy
---------------------------------   ---------  ---------  --------
copy, both sides coalesced              0.698      192.2      1.00
transpose, strided write                2.781       48.3      0.25
transpose, strided read                 1.875       71.6      0.37

All 3 rows moved exactly 128 MiB, from the same grid and the same
block, with the same index arithmetic. Only which subscript is
swapped differs, so the gap is the address pattern and nothing else.

The transpose reaches one quarter to one third of copy bandwidth on this card. The strided write is slower than the strided read, at 48.3 against 71.6 GB/s. A read miss can use cache, while a write miss must reach global memory.

All three kernels move exactly 128 MiB from the same grid, block, and index arithmetic. Only the swapped subscript changes. The address order alone cuts bandwidth by up to three quarters.

Day 13 fixes it by staging a tile through shared memory so both sides can be coalesced.

Run it yourself

A free Colab T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o transpose transpose.cu

There is no Compiler Explorer embed on this page. The program allocates two 64 MiB device buffers, checks 16,777,216 elements on the host three times, and runs 42 launches. That workload sits too close to Compiler Explorer's 20-second run limit.

Shrinking the matrix would also shrink the write stride under test. If you have no GPU, read /setup/learn-cuda-without-a-gpu.

Exercise

Run the program, then answer in two sentences: why is the naive transpose's fraction of the copy row far larger than one over thirty-two, and why does moving the strided side to the read not change it?

Time: 25 to 40 minutes. Submit: your three-row table, the naive transpose's fraction of the copy row, and the two sentences.

Check: before any timing, the harness compares all three outputs against the expected values exactly and prints the first bad index with its row and column. It then checks your two ratios against the range recorded for your GPU and reports them either way. If the copy row is not the fastest of the three, the measurement is broken and not your card.

Hint 1

Count what one warp asks the memory system for on each side of each kernel, then count what it uses. Do the copy first, so you have something to divide by.

Hint 2

A sector the warp half uses is not necessarily wasted. Ask who wants the rest of it, and how soon. Then ask what is different about day 11's stride-32 row.

Solution

Per warp the copy costs four sectors read and four written; the transpose costs four read and 32 written. That puts the transpose at 8/36 of the copy, 22 percent, already far above one over thirty-two, because only one of its two sides is strided.

The measured fraction usually beats even that, and the reason is reuse. A warp writing 32 elements down a column uses four bytes of each sector it touches, but the warps beside it write the neighbouring columns, so those sectors fill in before eviction. Day 11's stride-32 row has no such neighbour.

Swapping the subscripts leaves the count at 36, because a transpose maps a contiguous run to a run width apart in either direction.

Moving the strided access does not remove it. The sector count follows the transpose, so day 13 changes how the kernel performs that transpose.

Pitfalls

You compare your transpose against a bandwidth number from another program. A copy from a different block shape, grid shape or buffer size gives a ratio that contains all of those differences. Measure the copy in the same program and the same launch, which is what this page's copy row is for.

You swap the subscripts and report a fix. Both variants cost 36 sectors per warp. If your two transpose rows are far apart, that is a finding about read and write paths and worth chasing, but neither one is a coalesced transpose.

You transpose in place. With in and out sharing a buffer, one block may overwrite an element before another block reads it. __syncthreads() orders threads within one block only. Day 14 explains barrier scope.

You pick a matrix size that hides the problem. At 32 columns the write stride is 128 bytes and one warp's writes still fall inside a handful of cache lines. The guide draws the line in the same place: "When ld is larger than 32 (which occurs whenever the matrix sizes are larger than 32), this is equivalent to the pathological case shown in Figure 13." Benchmark at a size you would use.

You conclude that coalescing is unfixable here. It is fixable, it just is not fixable by choosing a better subscript. Every fast transpose stages the reordering through a shared-memory tile, and that is one lesson away.

Go deeper

Next

Day 13 fixes this kernel by loading a tile with both sides coalesced, reordering inside the block, and measuring the result against the copy row you just produced. Day 15 then explains why the obvious tile is slower than it should be.