Select the top k logits from a 128,000-token vocabulary on the GPU. What shape does the kernel take and why not just sort?

A row of logits is half a megabyte, so the kernel is not bandwidth-bound and the cost is in how many passes you make over it. Why a full sort does far more work than the question asks, the two shapes that fit small and large k, and the four-pass radix select with a reference that matches a sort exactly.

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.

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.

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.

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.

Explore the full curriculum →

UP NEXT ON YOUR JOURNEY

Next in this trackYou added three local variables to a working kernel and it got 30 percent slower. Explain what happened and how you would confirm it.Next in this trackA training step runs at 20 percent model FLOPs utilization. Profile it and find where the missing time goes.Next in this trackPort a hand-written CUDA kernel to MI300X. What translates mechanically, and what silently computes the wrong answer?

FEDITOR'S NOTE

Scored on rejecting a full sort with a reason, on choosing the structure from the size of k, on the per-block-then-merge shape for small k, and on knowing the radix select is a fixed number of passes regardless of vocabulary size.

DISCUSSION · 0

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