How do you overlap host-to-device transfers with compute, and what has to be true before the overlap actually happens?

Chunk the work, put copies and kernels on different streams, and the total stops being copy time plus compute time. The pipeline arithmetic that says how many chunks are enough, the three conditions without which the calls run in sequence anyway, and the two levels of asynchrony that are often confused.

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.

20 answers per topic instead of 10, plus saved progress and bookmarks · no cardor unlock all 283 remaining answers · ₹2,000 / $25

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.

Foundational

⚡ Kernels & Compilers

CUDA Programming ModelCUDA splits a program into a host that allocates, copies and enqueues work, and a device that runs thousands of identical threads organized as a grid of blocks. Getting the split right, and knowing that a launch returns before the kernel runs, decides whether your first live-coding kernel produces a correct number or a silent zero.

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.

Advanced

🔌 Networking & Storage🔒 Premium

Data Loading Pipelines for TrainingThe dataloader is the only part of a training job that runs on the CPU, the disk and the network at once, and it is the part most often found starving the GPUs. A pipeline that keeps 1,024 accelerators fed has to read sharded files sequentially, decode and tokenize in parallel workers, prefetch several batches ahead, pin memory for the PCIe copy, and do it deterministically enough to resume mid-epoch. The symptom of failure is a GPU at 30% utilization with nothing wrong on the GPU.

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.

Explore the full curriculum →

UP NEXT ON YOUR JOURNEY

Next in this trackWhy do tensor cores accumulate in fp32 when the inputs are bf16, and why do optimizers keep fp32 master weights?Next in this trackIn an int8 or fp8 GEMM, what exactly gets scaled, where does the dequantization happen, and what breaks if you get it wrong?Next in this trackA mixture-of-experts layer is 256 small matmuls. Why is a loop over cuBLAS calls slow, and what does a grouped GEMM do instead?

FEDITOR'S NOTE

Scored on the pipeline arithmetic, on naming pinned memory and non-default streams as preconditions rather than optimizations, and on separating host-to-device overlap from the device-side asynchronous copy inside a kernel.

DISCUSSION · 0

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