AI Infra Interviews logo
← CUDA, Triton & Kernel Engineering / 14
mediumNewNVIDIA

Transpose a large matrix at close to copy bandwidth. Why is the naive version slow, and what does each fix buy?

A transpose moves every byte exactly once, so a device copy is the honest ceiling. Why the naive kernel pays eight times the write traffic, why staging through shared memory does nothing until you pad, and the counter that tells the two problems apart.

Updated Sep 2026 · Grounded in real AI infrastructure interview loops and written to a senior-engineer editorial bar, with every number worked and every diagram hand-built.

A transpose moves every byte exactly once, so a device copy is the honest ceiling. Why the naive kernel pays eight times the write traffic, why staging through shared memory does nothing until you pad, and the counter that tells the two problems apart.

more free answers with an account · no card

The concepts behind this question

Ranked by how closely each one overlaps this question's topic, so the first card is the thing to read if the answer above moved too fast.

Advanced
⚡ Kernels & Compilers🔒 Premium
Shared Memory and Bank ConflictsShared memory is the programmer-managed SRAM inside each SM, split into 32 four-byte banks that serve one word each per cycle. When several lanes of a warp hit the same bank at different addresses the access serializes, and a 32-way conflict makes a shared-memory-bound loop run over ten times slower. Padding, XOR swizzles, cp.async and TMA are the tools that decide whether a tiled kernel gets the bandwidth it staged data for.
Foundational
🧩 GPU & Accelerator Architecture
GPU Memory HierarchyA GPU has four places a byte can live, and they differ by a thousandfold in bandwidth: registers, shared memory on the SM, a chip-wide L2, and HBM off-chip. Almost every kernel optimization is a decision about which level a value is read from and how many times. Knowing the sizes and bandwidths for an H100 cold is what lets you say why a kernel is slow before you profile it.
Core
⚡ Kernels & CompilersSign in
Memory CoalescingA warp's 32 threads issue one memory request together, and the hardware serves it in 32-byte sectors. Coalescing is arranging addresses so those sectors are full of bytes the warp will use. It decides whether a bandwidth-bound kernel moves at the HBM rate or at an eighth of it, and it is the pattern NVIDIA's trace-classification interview question tests.
Advanced
⚡ Kernels & Compilers🔒 Premium
Tiled Matrix MultiplicationA matrix multiply has enough reuse to be compute-bound, but only if the kernel captures that reuse in shared memory and registers instead of re-reading HBM. Tiling is how: a block owns an output tile, streams K-slices of A and B through shared memory, and each thread accumulates a small register tile. It is the live-coding exercise that separates people who know the roofline from people who have climbed it.
Explore the full curriculum →
UP NEXT ON YOUR JOURNEY
Next in this trackWhat does torch.compile actually do to your model, and when does it fail to help?Next in this trackWhy do serving engines capture decode steps into CUDA graphs, and what does capture require of the rest of the code?Next in this trackSketch a paged attention kernel. What changes from FlashAttention once the KV cache is not contiguous?
FEDITOR'S NOTE

Scored on naming the copy kernel as the ceiling, on the sector arithmetic for the strided writes, on knowing that the shared-memory fix trades a global problem for a bank-conflict problem, and on the profiler evidence that separates them.

DISCUSSION · 0

No comments yet — be the first to share your approach.