Port a hand-written CUDA kernel to MI300X. What translates mechanically, and what silently computes the wrong answer?

The source translation is a script and takes an afternoon. The assumptions underneath it are the problem: a wavefront is 64 lanes rather than 32, so a warp-level reduction written for 32 ports cleanly and reduces half the data. What breaks silently, what breaks loudly, and the numbers that change every tuning decision.

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

🧩 GPU & Accelerator Architecture🔒 Premium

AMD Instinct and ROCmAMD's Instinct line competes on memory: 192 GB on the MI300X, 256 on the MI325X, 288 on the MI355X, with dense fp8 peaks at or above NVIDIA's at each generation. The catch is software: ROCm and HIP run most PyTorch and the main serving engines, but the attention, MoE and quantization kernels arrive on CUDA first and measured throughput has trailed the spec sheet. Interviewers want the numbers, the wavefront-of-64 difference, and a clear-eyed statement of when the memory argument wins.

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.

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

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.

Explore the full curriculum →

UP NEXT ON YOUR JOURNEY

Next in this trackDesign a kernel benchmarking setup that will not lie to you. What does it control for, and how do fake speedups get published?Next in this trackYour fused attention kernel matches the reference at 512 tokens and drifts at 8,000. Find the bug.Closest to this topicThe MI300X has 192 GB per card. When does AMD actually win against an H100 for LLM serving, and what would you check before betting on it?

FEDITOR'S NOTE

Scored on naming wavefront 64 as the silent correctness hazard, on the local-memory and matrix-instruction differences that force retuning, and on knowing that the machine balance differs so the tuned configuration does not transfer.

DISCUSSION · 0

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