CUDA shared memory and tiling, measured on a transpose

Someone working through GPU Puzzles, on the puzzle whose whole point is the tile:

"I find that I can pass the two tests but I haven't use the shared memory. I feel that my solution is not correct but I don't know why the shared memory is needed here."

https://github.com/srush/GPU-Puzzles/issues/18 (checked 2026-08-29)

Their solution is correct, but it does not use shared memory to improve speed. A checker that only compares output values cannot show that difference.

This page times the transpose from day 12, where out[c][r] = in[r][c] makes either the read or write strided. You will move that strided access into shared memory, measure the gain, and test the 48 KiB default limit per block.

Where the stride has to go

A warp issues one load for all 32 lanes, and the hardware groups those addresses into 32-byte sectors. Thirty-two consecutive float addresses use four transactions with no unused bytes. Day 11 measured this coalesced case.

A transpose cannot coalesce both global-memory accesses in a one-thread-per-element kernel. Here the matrix is 2003 by 3001, and the naive kernel reads in[row * cols + col] with 32 consecutive col values. That is 128 contiguous bytes in four sectors.

The store uses out[col * rows + row], so its 32 addresses are 2003 floats, or 8,012 bytes, apart. It needs 32 sectors to provide 4 bytes to each lane. Swapping the subscripts moves the stride from the store to the load but does not remove it.

Shared memory is storage on the SM, with one allocation per thread block. It has no cache-line or sector traffic. Move the strided access there.

Use a 32 by 32 tile. One block reads an input tile by rows into shared memory, then reads that tile by columns and writes output rows. Both global-memory accesses are coalesced.

The strided step now occurs in shared memory. The tile uses 4,096 bytes. A Tesla T4 has 64 KiB per SM and a 48 KiB default limit per block (FACT-SHEET.md section 3), so the kernel uses one twelfth of the default per-block limit.

The tile saves no bytes at all

Tiling often lets a kernel load data once and use it many times instead of returning to global memory. Matrix multiplication does this because each value in a tile contributes to many output values. Day 16 covers that use of the memory hierarchy.

A transpose reuses nothing. Both the naive and tiled kernels read and write every element exactly once. The tiled version moves the same bytes across the same bus.

The tile changes the address pattern, not the byte count. This explains the quote at the top: both kernels return the right values, so an output-only check cannot tell them apart. A performance exercise must check both correctness and speed.

Shared memory is not a cache. Hardware does not fill or evict it, and its initial contents are undefined. Your block writes and reads it, then the allocation ends with the block.

Three rules, and where the barrier goes

The full program is in code/day13-shared-memory/shared_memory.cu. It has four kernels: a row-major copy, a naive transpose, and tiled transposes with static and dynamic shared memory.

All four launch the same grid and block and move the same bytes. Four elements per thread, 32 by 8 threads per block, one 32 by 32 tile of the matrix per block. The byte total is computed once from the two size constants rather than per kernel, so no row can win by doing less work. Day 11 exists because a published benchmark got that wrong.

Guard the load, not the barrier. __syncthreads() makes each thread's tile writes visible before other threads read them. The bounds checks wrap the load and store, while every non-exited thread reaches the barrier.

Days 5, 7, and 8 used if (i < n) { ... } instead of an early return. That form remains safe when you add this barrier.

Time with events, after warming up every kernel. Day 9 covers why, and the program uses that page's timeKernel unchanged.

__global__ void transposeTiledStatic(const float* __restrict__ in,
                                     float* __restrict__ out, size_t rows,
                                     size_t cols) {
    __shared__ float tile[kTileDim][kTileDim];

    size_t col = blockIdx.x * static_cast<size_t>(kTileDim) + threadIdx.x;
    size_t rowBase = blockIdx.y * static_cast<size_t>(kTileDim) + threadIdx.y;
    for (int j = 0; j < kTileDim; j += kBlockRows) {
        const size_t row = rowBase + static_cast<size_t>(j);
        if (row < rows && col < cols) {
            tile[threadIdx.y + j][threadIdx.x] = in[row * cols + col];
        }
    }

    __syncthreads();

    col = blockIdx.y * static_cast<size_t>(kTileDim) + threadIdx.x;
    rowBase = blockIdx.x * static_cast<size_t>(kTileDim) + threadIdx.y;
    for (int j = 0; j < kTileDim; j += kBlockRows) {
        const size_t row = rowBase + static_cast<size_t>(j);
        if (row < cols && col < rows) {
            out[row * rows + col] = tile[threadIdx.x][threadIdx.y + j];
        }
    }
}

The first loop writes tile[threadIdx.y + j][threadIdx.x], so each warp writes a tile row. The second reads tile[threadIdx.x][threadIdx.y + j], so it reads a column.

After the barrier, col and rowBase use the other blockIdx. The block that reads tile (x, y) writes tile (y, x). Reusing the input indices would copy the matrix instead of transposing it.

Note. The tile is [32][32] with no padding, so the column read in the second loop puts all 32 lanes of a warp on one of shared memory's 32 banks. That 32-way bank conflict is deliberate: this kernel should beat the naive one and should still miss the copy ceiling, and day 15 closes the gap with one extra column.

Static, dynamic, and the wall at 48 KiB

There are two ways to get a tile, and they differ only in when its size is fixed.

__shared__ float tile[32][32];   // size fixed when you compile
extern __shared__ float tile[];  // size fixed when you launch

Pass the dynamic size in bytes as the third part of the execution configuration: kernel<<<grid, block, bytes>>>(...). Each block gets one unsized array, so code must compute two-dimensional indices by hand. Use the static form when the tile size is fixed.

Dynamic shared memory has a default per-block limit. On a T4, the default is 48 KiB and the opt-in limit is 64 KiB. A larger request needs per-kernel opt-in and otherwise fails at launch, not compile time.

The program requests 4 KiB more than the reported default, prints the error, opts in, and launches again.

    touchDynamicShared<<<1, kThreadsPerBlock, requestBytes>>>(d_probe, floats);
    const cudaError_t beforeOptIn = cudaGetLastError();
    std::printf("%zu bytes of dynamic shared memory, no opt-in: %s\n",
                requestBytes, cudaGetErrorString(beforeOptIn));

    // cudaFuncSetAttribute's third parameter is an int, so the narrowing is
    // written out rather than left to the call. requestBytes is bounded by
    // optInBytes a few lines up, so it fits.
    CUDA_CHECK(cudaFuncSetAttribute(touchDynamicShared,
                                    cudaFuncAttributeMaxDynamicSharedMemorySize,
                                    static_cast<int>(requestBytes)));

    touchDynamicShared<<<1, kThreadsPerBlock, requestBytes>>>(d_probe, floats);
    const cudaError_t afterOptIn = cudaGetLastError();
    std::printf("%zu bytes after cudaFuncSetAttribute:         %s\n",
                requestBytes, cudaGetErrorString(afterOptIn));

Hardware. The two numbers are sharedMemPerBlock and sharedMemPerBlockOptin in cudaDeviceProp, and the program prints both for your card. Neither is shared memory per SM: the docs give per-SM and per-block figures that differ by 1 KB on several architectures (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , Table 31, checked 2026-08-29). A card with no headroom above its default makes the probe say so and skip.

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: 2003 x 3001 floats, 22.9 MiB per buffer
Shared memory per block: 49152 bytes default, 65536 bytes opt-in
Shared memory per SM:    65536 bytes
Tile: 32 x 32 floats, 4096 bytes per block

kernel                     time (ms)        GB/s    vs copy
------------------------  ----------  ----------  ---------
copyRowMajor                   0.228       211.2      1.00x
transposeNaive                 0.737        65.2      3.24x
transposeTiledStatic           0.406       118.6      1.78x
transposeTiledDynamic          0.407       118.2      1.79x

Every row above moved exactly 45.9 MiB. All four launch the same
grid and the same block and read and write every element once, so
no row can win by doing less work.

53248 bytes of dynamic shared memory, no opt-in: invalid argument
53248 bytes after cudaFuncSetAttribute:         no error
all four kernels match the CPU reference at 6011003 elements

The CUDA 13.0 re-run on the same T4 preserved every correctness and launch result. Absolute bandwidth was about 10% lower for the naive and tiled transposes, but the useful comparison did not move: tiling was 1.83x faster than the naive transpose, static and dynamic shared memory were tied at 106.9 GB/s, and the 53,248-byte request still failed before the opt-in and succeeded after it.

Tiling improves bandwidth by 1.8 times. The naive transpose runs at 65.2 GB/s and the tiled one at 118.6 GB/s. The gap to a straight copy falls from 3.2x to 1.8x, while both tiled kernels move the same 45.9 MiB as the copy.

Static and dynamic shared memory measure the same, 118.6 against 118.2 GB/s. Choose on whether the size is known at compile time, not on speed.

A request above the default needs opt-in. Asking for 53,248 bytes of dynamic shared memory fails at launch with invalid argument, because this device defaults to 49,152 bytes per block. The same request succeeds after cudaFuncSetAttribute raises the limit toward the reported 65,536-byte maximum.

The error text does not name shared memory. Day 74 uses the same opt-in for a pipeline.

The tiled version still runs below copy speed. It reads the tile down a column, and every lane of that read lands in the same bank. Day 15 measures what that costs.

Run it yourself

The program is one file with no shared headers, so it fits in Compiler Explorer. Target sm_75 or lower. A higher target may compile but fail on the runner's GPU.

The matrix has 6,011,003 floats, and the program launches four kernels thirteen times each. If a shared runner times out, reduce kRows and kCols. The static_assert at the top checks that neither size divides by the tile.

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

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

Exercise

Write solve() for the transpose so that both the global read and the global write are coalesced, then report your tiled time as a fraction of the harness's copy baseline and say in one sentence what the remaining gap is.

Time: 30 to 45 minutes. Submit: your shared_memory.cu, plus the ratio and the sentence.

Check: the harness runs the 2D case ladder, including (31, 33) and (611, 613), and reports the smallest failing case rather than the first. A failure prints the case, the first ten wrong elements as output row and column, and one line naming what they have in common, which here is usually "cells one tile did not write".

It then measures its own coalesced copy over your buffers, on your card, seconds before it times your kernel, and prints your ratio against it. Correct but no faster than the naive kernel it ships with is a distinct outcome and it says so.

Hint 1

One of the two global accesses has to be strided and you cannot make both contiguous, so stop trying to. Where could the strided step happen instead, and what would have to be true of that place for the stride not to cost anything?

Hint 2

Work out which thread reads input element (r, c) into the tile, and which thread writes it back out. They are not the same thread. That is the whole reason there is a barrier between the two loops, and why the block indices on the second loop come from the other axis.

Solution

A 32 by 32 tile in __shared__, a __syncthreads(), and the output block indices swapped, which is the kernel above. The diff against the starter is one shared array, one barrier and two recomputed index lines.

On the shape: your tiled row should sit much closer to the copy baseline than the naive row does, and should not reach it. What is left over is the column read out of an unpadded tile, which puts every lane of the warp on one bank, and day 15 closes it with one extra column.

Shared memory lets you change the order of global-memory accesses. The byte count stays the same. Check the 32 addresses used by each warp, and move a needed stride into shared memory.

Pitfalls

The tile holds values your block never wrote. Shared memory is not zeroed and promises nothing about its contents. A learner on GPU Puzzles reported that "the value got carried over from Block 0, even after calling cuda.syncthreads()" (https://github.com/srush/GPU-Puzzles/issues/20 , checked 2026-08-29). Write every cell you read, and where a guard leaves cells unwritten, do not read them.

You guard with an early return above the barrier. It does not hang, which is why it survives review. The Programming Guide says __syncthreads*() "wait until all non-exited threads in the thread block simultaneously reach the same __syncthreads*() intrinsic call in the program or exit" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html , checked 2026-08-29), and a thread that returned has exited.

You get a tile with holes in it, read back as the previous block's leftovers. Guard the load, leave the barrier alone. Day 14 covers the semantics, including the case that does hang: a barrier inside a condition that is not uniform across the block.

You read the tile with the indices you wrote it with. tile[threadIdx.y][threadIdx.x] on both sides is a copy with extra steps, caught at the first off-diagonal element. Write the index out rather than hiding it in a macro: an unparenthesized index macro is a documented way to lose an afternoon, which is why this course keeps two macros and neither takes an index.

A kernel that asks for more than the default shared memory never runs. The launch is refused, cudaGetLastError() returns a code whose string names no shared memory, and a program missing the two checks after a launch prints an empty output buffer instead. The fix is cudaFuncSetAttribute with cudaFuncAttributeMaxDynamicSharedMemorySize, per kernel, before the launch, and it reaches only as far as your card's sharedMemPerBlockOptin. Day 74 uses it in a pipeline.

More shared memory per block is not better. The SM has a fixed budget, so a block claiming half of it means two blocks fit per SM and occupancy drops. Ask for what the tile needs. Day 45 measures the trade, and cp.async is what you reach for when you want both.

Go deeper

• CUDA Programming Guide, "Shared Memory", including the transpose and the dynamic form: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)
• CUDA Programming Guide, Compute Capabilities appendix, Table 31, for shared memory per SM and per thread block on every architecture: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-29)
• cuda-samples, cpp/6_Performance/transpose, NVIDIA's own ladder over this exact kernel: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/6_Performance/transpose (checked 2026-08-29)
• Programming Massively Parallel Processors, 4th edition, chapter 5, on the memory hierarchy and tiling: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Day 14 takes the one line this page told you to leave alone and asks what it actually guarantees, then breaks a tile kernel and proves the race with compute-sanitizer. Day 15 measures the gap this page's tiled kernel leaves between itself and the copy row, and closes most of it by making the tile 32 by 33. Both sit in module 2, and day 16 is where a tile finally buys you reuse instead of only order.