Companies / NVIDIA CORE AI INFRASTRUCTURE
NVIDIA AI Infrastructure Engineer interview questions NVIDIA hires across every layer this site covers: deep learning library and performance engineers (cuDNN, CUTLASS, TensorRT, TensorRT-LLM, LLM performance), developer technology engineers, Triton Inference Server systems engineers, DGX Cloud performance engineers who benchmark distributed systems to define cluster architecture, and GPU and HPC infrastructure engineers who run NVIDIA's own fleets. The loop is domain-heavy and C++-first: a recruiter call, one or two 45 to 60 minute phone screens that mix C++ fundamentals (virtual dispatch, memory model, move semantics) with graph and array problems, then four to six onsite rounds where the team's domain dominates; candidates consistently report that domain expertise in the team's area is the filter. Reported kernel questions include memory coalescing, shared-memory bank conflicts, warp divergence, parallel reduction, GEMM tiling, and profiling with Nsight. Loops run four to eight weeks. NVIDIA India hires in Pune and Bengaluru.
CHIP AND PLATFORM VENDORS
They build the silicon and the software that drives it, so depth in the team's own domain outranks breadth almost everywhere.
Loop leans on: Microarchitecture, kernels, compilers, interconnect, benchmarking. Compare the other chip and platform vendors →
The NVIDIA AI Infrastructure Engineer interview process Documented How the NVIDIA AI Infrastructure Engineer interview experience actually runs — the rounds, what each stage tests, and the signals candidates report. Last reviewed September 4, 2026.
Role Deep Learning Software Engineer / GPU and HPC Infrastructure Engineer / DGX Cloud Performance Engineer Loop 4 to 8 weeks typical; long loops are a recurring complaint AI tools No first-party statement found; no candidate report of AI being permitted.
1 Recruiter call Team and level.
2 Phone screens One or two rounds of 45 to 60 minutes: resume discussion plus C++ fundamentals (virtual dispatch, memory model, move semantics) and graph or array problems.
3 Take-home (team-dependent) One 2022 report of two 48-hour exercises; most 2023 to 2026 reports do not mention one.
4 Onsite Four to six rounds of about an hour with domain-specific content: systems architecture, performance, team-specific coding (whiteboard CUDA and kernel pseudo-code reported for DevTech and library loops), and a behavioral round; sometimes a hiring-manager or director round.
WHAT THEY'RE EVALUATING
› Expertise in the team's domain, reported as the dominant filter› C++ first, CUDA, Python for tooling› Memory coalescing, bank conflicts, warp divergence, parallel reduction, GEMM tiling, Nsight profiling (reported kernel questions)› TensorRT-LLM, vLLM, SGLang and speculative decoding for LLM performance rolesSpecific kernel questions are mostly aggregator-reported with few dated first-hand debriefs; treat individual items as partial.
Compiled from our research and publicly available information (candidate reports and company interview guides). Interview loops change and are continuously iterated, and they vary by team, level, and region. Treat this as directional preparation, not an official spec, and confirm the exact rounds with your recruiter or hiring point of contact.
Questions modeled on NVIDIA loops 107 questions · 42 unlocked for you
01
Walk me through the CUDA execution model: what are grids, blocks and warps, and what does the hardware actually schedule? ▼ easy ★ Essential New NVIDIA Google CoreWeave 4 replies unlocked
A grid is a request, a block is a residency unit, a warp is what the scheduler issues. Which of those pins to an SM, why 32 matters, and how a GPU hides a 600 ns memory latency with no branch predictor and a cache that is tiny per thread.
02
Describe the GPU memory hierarchy. Where can a byte live on an H100, and what does each level cost? ▼ easy New NVIDIA AMD 4 replies unlocked
Registers, shared memory, L1, L2, HBM, host memory: sizes, bandwidths and latencies for an H100, derived rather than recited, and the habit of asking 'which level am I hitting' before touching a kernel.
03
Why are GPUs so much faster than CPUs for deep learning? Be specific about what the silicon is doing differently. ▼ easy New NVIDIA Google 4 replies unlocked
Not 'more cores.' The real answers are a 10x memory system, a 16x matrix datapath, and a design that spends transistors on lanes and registers instead of on making one thread wait less. With the numbers for a two-socket server against one H100.
04
What is a tensor core, and what does a kernel have to do to actually use one? ▼ easy New NVIDIA 4 replies unlocked
A tensor core multiplies small matrix tiles in one instruction and is 16x faster than the regular lanes, but only for dense matmul at the right precision, with the right shapes and layouts. What the instruction looks like, what it refuses, and how to tell from a profile whether you are on it.
05
Explain arithmetic intensity and the roofline model. Where is the ridge point on an H100, and what does it tell you about a kernel? ▼ medium ★ Essential New NVIDIA Fireworks Together AI 4 replies unlocked
Peak FLOPS divided by bandwidth is one number per chip, and it decides whether any kernel can ever reach peak. How to compute it, how to compute a kernel's intensity from its bytes and FLOPs, and how to read the answer before writing a line of CUDA.
07
bf16, fp16, fp8: what is the difference at the bit level, and where does each one belong in training and serving? ▼ medium New NVIDIA Meta 4 replies unlocked
Every format is a trade between how big a number can be and how finely it is spaced. The bit layouts, the largest and smallest values derived from them, why fp16 needed loss scaling and bf16 did not, and why fp8 comes in two flavors with a scale factor attached.
08
How much faster is an H200 than an H100, really? Which workloads see the gain and which do not? ▼ medium New NVIDIA CoreWeave Lambda 4 replies unlocked
The H200 has the same compute die as the H100 and costs more per hour. The datasheet gives two ratios, 1.43x bandwidth and 1.76x memory, and those two numbers decide exactly which workloads pay back the premium and which ones lose money on it.
09
NVLink versus PCIe: when does the interconnect actually matter for an LLM workload, and when can you get away with PCIe? ▼ medium New NVIDIA Lambda Crusoe 4 replies unlocked
An H100 SXM moves 900 GB/s over NVLink; a PCIe card moves about 64 GB/s per direction over its slot. Whether that 7x to 14x gap costs you anything depends on how many bytes your parallelism strategy sends per step, and the answer is very different for tensor and data parallelism.
10
What is warp divergence, why does it cost you, and how would you find and fix it in a kernel? ▼ medium New NVIDIA 4 replies unlocked
A GPU issues one instruction to 32 threads at once. When those threads want different instructions, the hardware runs each path in turn with the others masked off, and the time is the sum of the paths. The derivation of the cost, the metric that exposes it, and the three fixes that work.
11
A kernel runs at 25% occupancy. Is that a problem? Walk me through what occupancy buys and when you would leave it low on purpose. ▼ medium New NVIDIA Fireworks 4 replies ○ sign in
Occupancy is a means, and the end is enough bytes in flight to cover memory latency. Little's law gives the number of loads an SM needs outstanding; the register file gives the warps you can afford; the kernels that win on H100 usually run at 25% to 50% occupancy with big tiles and no spills.
12
Explain what HBM is and why memory bandwidth, not compute, is the wall for LLM inference. ▼ medium New NVIDIA AMD Micron 4 replies ○ sign in
HBM stacks DRAM dies on top of each other and wires them to the GPU through a silicon interposer with a 1,024-bit bus per stack. That design sets how much bandwidth and capacity a card can have, why the two scale together, and why they have grown more slowly than FLOPS across three generations.
13
What actually changes with Blackwell and the NVL72 rack, and what does it do to how you would serve a large MoE model? ▼ hard New NVIDIA CoreWeave Microsoft 4 replies ○ sign in
Three things change at once: dense compute more than doubles with an fp4 tensor core, HBM3e reaches 8 TB/s, and NVLink stops at 72 GPUs instead of 8. Work through what each does to tensor-parallel degree, per-step weight reads and expert placement for a 671B-parameter mixture of experts.
19
MIG versus MPS: what isolation does each give you when sharing a GPU, and which would you pick for a multi-tenant inference node? ▼ medium New NVIDIA CoreWeave Nebius 4 replies ○ sign in
MIG carves an H100 into up to seven hardware slices with their own SMs, memory and bandwidth; MPS lets processes share one SM pool through a single context. One gives fault and performance isolation at a fixed slice size; the other gives flexibility and a shared failure domain. The numbers that decide it.
20
You are moving a model to fp4 inference on Blackwell. What breaks first, and how would you measure whether the result is acceptable? ▼ hard New NVIDIA Fireworks 4 replies ○ sign in
Four bits give eight magnitudes per sign, and the largest step between them is 50%. Block scaling rescues most of the tensor; a single outlier in a block of sixteen flushes its neighbors to zero. Where the precision goes, which layers fail first, and the measurements that separate acceptable from broken.
21
Design a non-blocking network fabric for 512 H100s. How many switches, how are they wired, and where does the NVLink domain end? ▼ hard New NVIDIA Crusoe Lambda 4 replies ◆ premium
512 GPUs at 400 Gb/s each is 205 Tb/s of injection bandwidth a two-tier fabric must carry with no oversubscribed link. The port arithmetic that gives 16 leaves and 8 spines, the rail wiring that keeps NCCL traffic one hop away, and the line where 900 GB/s of NVLink becomes 50 GB/s of InfiniBand.
22
Explain GPUDirect RDMA. What does it remove from the data path, and when does it change performance? ▼ medium New NVIDIA Crusoe 4 replies ◆ premium
Without it, every byte a GPU sends over the network crosses PCIe twice and passes through host memory. With it, the NIC reads and writes GPU memory directly. The PCIe arithmetic that shows why a 400G NIC cannot be fed the slow way, the NCCL line that tells you which path you are on, and the checkpoint case.
23
An Nsight summary reads DRAM 90%, SM 30%, occupancy 45%, L2 hit rate 20%. What is the bottleneck? ▼ hard New NVIDIA Fireworks Together AI 4 replies ◆ premium
Four percentages from a profiler summary are enough to place a kernel on the roofline, estimate its arithmetic intensity, and rule out half the optimizations people reach for. The reading, the arithmetic that backs it, and the fix list in the order it pays.
25
What can you rely on from the L2 cache on a GPU? When does it save a kernel, and when does it mislead you? ▼ hard New NVIDIA 4 replies ◆ premium
An H100 has 50 MB of L2, split into two partitions, with a bandwidth several times HBM's. That is enough to change the roofline for a kernel whose working set fits and to do nothing at all for the ones that stream. The arithmetic of what fits, the residency controls, and the two ways the L2 lies to a profile.
29
If you could change one thing about GPU architecture for LLMs, what would it be, and what argues against it? ▼ expert New NVIDIA Anthropic OpenAI 4 replies ◆ premium
The workload asks for bytes and the chip delivers FLOPs: bytes per FLOP halved from A100 to H100 and held level on B200. A defensible thesis is more bandwidth and capacity per FLOP, the shoreline and power arithmetic for why it has not happened, and the counterargument that software already routes around the wall.
30
One node in your cluster shows half the expected NVLink bandwidth in nccl-tests. Walk me through isolating it. ▼ expert New NVIDIA CoreWeave Crusoe 4 replies ◆ premium
Half the expected bus bandwidth has four plausible causes needing different fixes: a link trained at reduced width, a fabric that never came up, a topology the job did not expect, or a test that measured latency. The sequence of commands that tells them apart, and what each output looks like on a healthy node.
02
Your fleet is decode-heavy. Is a B300 worth 1.4 times a B200's power for 1.6 times the memory? ▼ medium New NVIDIA Baseten Together AI 4 replies unlocked
Decode is bandwidth-bound, so the FP8 FLOPS number that dominates the marketing does not move it. Where the B300 pays is capacity, and capacity converts into throughput through batch size rather than directly. The arithmetic that decides it, and the case where the B200 wins.
04
After recabling a row to 800G, a third of the links will not come up. Find the cause. ▼ medium New NVIDIA CoreWeave Meta 3 replies unlocked
Links that never come up and links that come up and flap are different faults with different causes. The reach change between 400G and 800G explains most first-time failures, and the counters separate the rest. The order to check, and the one physical measurement that settles it.
07
During acceptance, seven rails hit expected bandwidth and one is 15 percent low. Find it. ▼ hard New Meta CoreWeave NVIDIA 3 replies unlocked
A rail-optimized fabric makes this fault findable because the eight rails are identical by construction, so any asymmetry is a defect. Bisecting by rail, then by switch, then by node, then by port narrows 512 links to one in four steps, and the counters say which of three physical causes it is.
08
New nodes arrive with no software. Walk me through bring-up, and say why the order matters. ▼ medium ★ Essential New NVIDIA CoreWeave Meta 4 replies unlocked
Each layer depends on the ones below it and a mistake at the bottom shows up six layers higher as something else entirely. The order, the dependency almost everyone gets wrong on Blackwell systems, and the validation ladder that keeps a fault attributable to one node instead of five hundred.
10
You are building a fine-tuning service for customer models under 30B. RTX PRO 6000 or H100? ▼ medium New Lambda Labs Modal Baseten 4 replies unlocked
The part with no NVLink can be the correct choice, and the reason is the workload shape rather than the specification. What fits on one card, what MIG partitioning buys a multi-tenant service, and the exact point where the missing scale-up link makes the decision flip.
12
Count the switches, cables and transceivers for a 2,048-GPU rail-optimized fabric. ▼ medium New Crusoe CoreWeave NVIDIA 4 replies ○ sign in
The count follows from three inputs and one topology rule, and getting it wrong by the leaf-to-spine half is the classic error. The arithmetic, the radix check that says two tiers is enough, and the power the transceivers add that nobody put in the budget.
13
After a maintenance window, two-node bandwidth halved. Single-node is unchanged. Diagnose it. ▼ hard New NVIDIA CoreWeave Meta 4 replies ○ sign in
Single-node fine and two-node halved is a specific signature, and it points at one thing before anything else. Why the copy path costs exactly about half, the five causes that produce it, and the one log line that confirms which.
20
A GPU cluster has four separate networks. Name them and say what each one carries. ▼ medium ★ Essential New NVIDIA CoreWeave Crusoe 3 replies ○ sign in
Most architecture diagrams show one network and a cluster has four. Each has a different traffic pattern, a different failure consequence and a different port count, and the one that is under-counted every time is the one you need when everything else is broken.
21
Design the fabric for 10,000 GPUs. Two tiers or three, and what does the choice cost? ▼ hard New Meta Microsoft Crusoe 3 replies ◆ premium
Switch radix decides it, and at 10,000 endpoints the answer sits right at the boundary for current high-radix switches. The arithmetic, what a third tier costs in switches, optics and power, and the two designs that avoid it entirely.
22
For 64 to 72 GPUs, is one GB300 NVL72 rack better than eight HGX B300 nodes? ▼ hard New NVIDIA Microsoft CoreWeave 4 replies ◆ premium
Same GPU count, same silicon, and a completely different machine. The domain size is the whole argument, it is worth an order of magnitude on specific traffic, and it costs a failure domain nine times larger plus a facility that most halls do not have.
25
What does nvidia-fabricmanager do, and what exactly breaks when it is not running? ▼ medium New NVIDIA CoreWeave Meta 3 replies ◆ premium
It configures the NVSwitch fabric so eight GPUs present as one NVLink domain, and when it is absent nothing errors. Jobs run, results are correct, and everything multi-GPU is slower by a factor the team spends weeks not finding. The check that takes one command.
27
Walk me through cabling a rail-optimized cluster. What goes where, and what goes wrong? ▼ medium New Crusoe CoreWeave NVIDIA 4 replies ◆ premium
The rule is one sentence and the execution is thousands of connections where a single transposition creates a fault that is invisible until a collective runs. The mapping, the labelling scheme that makes errors findable, and the verification that has to happen before anyone accepts the cluster.
36
Transceivers are failing at three times the rate you budgeted. Find out why. ▼ hard New Meta CoreWeave NVIDIA 4 replies ◆ premium
Modules do not fail randomly at three times a published rate, so the excess has a cause and it is usually thermal or a batch. The cohort cuts that localize it, the counters that separate a dying module from a dirty connector, and the environmental cause that looks like a hardware defect.
01
Write a CUDA vector add and explain the launch: grid math, the bounds check, and why the copies dominate. ▼ easy New NVIDIA 4 replies unlocked
The kernel is four lines; the interview is about the other forty. How a global index comes out of block and thread ids, why the bounds check exists, what the launch configuration means for a 100M-element array, and the arithmetic that shows the host-to-device copies cost 50 times more than the add.
02
What is memory coalescing, why does a strided access pattern hurt, and how do you see it in a profiler? ▼ easy New NVIDIA Fireworks 4 replies unlocked
A warp issues one load instruction and the memory system turns it into some number of 32-byte sector requests; that number is the whole story. The arithmetic for contiguous, stride-2 and stride-32 access, the row-major matrix where a loop order change gives 8x, and the two Nsight Compute counters that show the waste.
03
Given the addresses each thread in a warp touched, classify the access pattern: coalesced, strided or random. Write the classifier. ▼ medium New NVIDIA 3 replies unlocked
Thirty-two addresses per warp instruction, thousands of instructions: say what the pattern is and how many sectors it cost. The address-delta test for contiguous and strided, the sector count that measures the damage, the code that does both, and the edge cases (misalignment, inactive lanes, mixed widths).
04
Explain shared memory bank conflicts with the bank arithmetic, show a kernel that has them, and fix it with padding. ▼ medium New NVIDIA Together AI 4 replies unlocked
Shared memory has 32 banks, each 4 bytes wide, and a warp's access is as slow as the most-loaded bank. The bank of an address, why a column walk down a 32-wide tile puts all 32 lanes in one bank, the padding by one column that spreads them across all 32, and the profiler counter that confirms the fix.
07
Take a GEMM from naive to 70% of peak: the steps, the speedup at each, and the arithmetic that says why. ▼ hard New NVIDIA Fireworks 4 replies unlocked
The naive kernel reads two bytes per multiply-add; a 128x128 tile with register blocking reads a few hundredths of that. Each step (shared-memory tiles, register blocking, vectorized loads, double buffering, tensor cores) with its intensity arithmetic and speedup, and where a hand kernel stops and cuBLAS begins.
08
Explain occupancy and register pressure: launch bounds, spills, the calculator, and why 50% occupancy can beat 100%. ▼ medium New NVIDIA 4 replies unlocked
Occupancy is how many warps an SM holds against its maximum, and it is a means, not an end. The arithmetic from registers per thread to resident warps, the launch bound that caps registers and the spills that follow, the profiler's occupancy view, and the kernel where halving occupancy doubled speed.
10
What changed between FlashAttention 1, 2 and 3, and why did each change buy what it did? ▼ hard New Fireworks Together AI NVIDIA 4 replies unlocked
Each version fixed a different bottleneck: version 1 fixed memory traffic, version 2 fixed non-matmul work and parallelism, version 3 fixed the fact that softmax and matmul were waiting for each other on Hopper. The arithmetic that shows why a softmax costing one percent of the FLOPs can cost sixty percent of the time.
12
Here is an Nsight Compute report for a slow kernel. Read it, name the bottleneck, and tell me what you would change. ▼ hard New NVIDIA Fireworks 4 replies ○ sign in
Four sections, read in a fixed order, and one number that usually names the bottleneck outright. The report of a kernel at 12 percent of DRAM bandwidth while its memory pipeline reads 82 percent busy, what that gap means, the fix it implies, and the numbers the fixed kernel reports back.
13
Sum 100 million floats on the GPU as fast as the hardware allows. Write the kernel and justify each step. ▼ medium New NVIDIA 4 replies ○ sign in
A sum reads every byte once and does one add per element, so the only question is whether you reach the bandwidth ceiling. The ceiling in milliseconds, four versions from an atomic per element to warp shuffles with sixteen loads in flight, and the bytes-in-flight math behind the last jump.
14
Transpose a large matrix at close to copy bandwidth. Why is the naive version slow, and what does each fix buy? ▼ medium New NVIDIA 3 replies ○ sign in
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.
16
Why do serving engines capture decode steps into CUDA graphs, and what does capture require of the rest of the code? ▼ hard New Fireworks Together AI NVIDIA 4 replies ○ sign in
At batch 1 a decode step is a few milliseconds of GPU work behind several hundred kernel launches, and the CPU cannot issue them fast enough. The arithmetic showing the GPU starve, what replay changes, and the four constraints capture imposes on memory.
18
cuBLAS already gives you a fast GEMM. What does CUTLASS give you that it does not, and when is that worth it? ▼ hard New NVIDIA Fireworks 4 replies ○ sign in
A library call gives you a tuned matmul and nothing around it. Four things a template library gives instead: an epilogue fused into the output stage, one launch for a list of ragged problems, a matmul whose inputs are different types, and a tile schedule that fills the SMs on awkward shapes.
21
How do you overlap host-to-device transfers with compute, and what has to be true before the overlap actually happens? ▼ medium New NVIDIA Lambda 4 replies ◆ premium
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.
22
Why do tensor cores accumulate in fp32 when the inputs are bf16, and why do optimizers keep fp32 master weights? ▼ medium New NVIDIA Meta 4 replies ◆ premium
A float format's resolution is relative, so the damage depends on how large a running total grows against the terms being added. The measured error for a 4,096-term dot product in each direction, the weight update that vanishes entirely, and why bf16 removed loss scaling but not master weights.
23
In an int8 or fp8 GEMM, what exactly gets scaled, where does the dequantization happen, and what breaks if you get it wrong? ▼ hard New NVIDIA Fireworks 4 replies ◆ premium
The matmul runs on tiny formats and the accumulator does not, so every low-precision GEMM is really a scaled integer or float product with a correction applied on the way out. Where the scales come from, the three granularities and what each costs, and why one of them needs a periodic promotion into fp32.
25
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? ▼ medium New NVIDIA OpenAI 4 replies ◆ premium
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.
26
You added three local variables to a working kernel and it got 30 percent slower. Explain what happened and how you would confirm it. ▼ medium New NVIDIA 4 replies ◆ premium
Registers are allocated in fixed steps out of a fixed budget per multiprocessor, so a small increase in live values can cost a whole resident block. The compiler output that shows it in two lines, the occupancy cliff arithmetic, and the two different failure modes that produce the same symptom.
02
Compare data, tensor and pipeline parallelism. What does each one shard, what does each one communicate, and where does each one live? ▼ easy ★ Essential New NVIDIA Meta OpenAI 4 replies unlocked
Three ways to split a training job, one table, and the rule that places each of them: activations on NVLink, gradients on the fabric, stage boundaries in between. With the byte counts that justify the placement.
07
Derive the cost of a ring all-reduce. Why is it bandwidth-optimal, and where does it stop scaling? ▼ medium New NVIDIA Google 4 replies unlocked
Two phases, n minus one steps each, one slice per step: the derivation that gives 2(n−1)/n times the buffer per rank, the argument that no algorithm can beat it on bytes, and the latency term that makes NCCL switch to trees at scale.
08
Do the pipeline bubble math for me. How do micro-batches shrink it, and what does 1F1B change? ▼ medium New NVIDIA Meta 4 replies unlocked
With p stages and m micro-batches the pipeline idles (p−1)/m of its compute time, and the obvious fix, more micro-batches, buys back throughput with activation memory. 1F1B caps that memory at p micro-batches instead of m. The derivation, the numbers for PP16, and the schedules that go further.
09
Why does tensor parallelism stop at 8? Show me the numbers. ▼ medium New NVIDIA Anthropic 4 replies unlocked
Four all-reduces per transformer block, on activations, on the critical path. Inside the NVLink domain they cost 8% of a step; across the NIC they cost more than the step itself. The derivation for a 70B, and the second reason TP stops that has nothing to do with the network.
15
You have 16,384 H100s and a 405B dense model. Choose the parallelism layout and justify the order of the axes. ▼ hard ★ Essential New Meta NVIDIA 4 replies ○ sign in
Tensor parallel 8, pipeline 16, data parallel 128: the layout Meta trained Llama 3 405B on, rebuilt from the memory floor, the compute budget and the bandwidth of each link. Why the axes go in that order, what each one costs per step, and the bubble that the schedule has to hide.
16
Explain mixed-precision training. Why does fp16 need loss scaling and bf16 not, why keep fp32 master weights, and what changes with fp8? ▼ medium New NVIDIA Meta 4 replies ○ sign in
fp16 loses gradients below 6e-8 and overflows above 65,504; bf16 has fp32's range and three digits of precision. The arithmetic that shows why an update of 3e-5 vanishes into a bf16 weight, where each of the 16 bytes per parameter comes from, and what fp8 block scaling adds on top.
23
How do you overlap communication with compute in training, and where does the overlap stop working? ▼ hard New NVIDIA Meta 4 replies ◆ premium
Bucketed all-reduce in backward, FSDP's one-layer-ahead prefetch and split-GEMM tensor parallelism hide most of the bytes, and none of it works below a certain number of tokens per GPU. The ratio that decides it, the SMs the NCCL kernels take, and the tail that is never hidden.
26
Design the infrastructure for RLHF on a 70B: where do rollouts and the learner run, and how do weights move between them every step? ▼ hard New OpenAI Anthropic NVIDIA 4 replies ◆ premium
Generation is a decode workload and the update is a training workload, and they want different software on different GPUs. The per-token arithmetic for each side, the 141 GB broadcast that has to happen every step, the ratio of actor to learner GPUs that follows, and the idle time that on-policy training builds in.
28
We are running GRPO at scale. What does the infrastructure have to do that plain RLHF did not, and where do the GPUs sit idle? ▼ hard New DeepSeek OpenAI NVIDIA 4 replies ◆ premium
Sixteen samples per prompt, a verifier that runs code, no value model, responses to 16k tokens: GRPO moves the cost from the learner to rollouts and rewards. The token arithmetic for one step, the tail that holds a batch for the longest sample, and the three places a fleet idles.
29
Why is fp8 training hard, and how did DeepSeek-V3 make it work? ▼ hard New DeepSeek NVIDIA 4 replies ◆ premium
Three mantissa bits and a range of 448 make fp8 unusable with one scale per tensor. The arithmetic of what an outlier channel destroys, the 128-element block scaling and fp32 promotion that the DeepSeek-V3 report used, the ops that stayed in bf16, and what the 2× peak bought in practice.
11
When does splitting prefill and decode onto separate GPU pools pay for itself, and what does the KV transfer cost? ▼ hard ★ Essential New Anthropic Fireworks NVIDIA 4 replies ○ sign in
Prefill and decode fight over the same GPU and each ruins the other's latency. Putting them on separate pools ends the fight, at the price of shipping every request's cache across the network. The break-even is a number you can derive.
15
vLLM, SGLang or TensorRT-LLM: which engine do you pick for a new deployment, and what would change your mind? ▼ hard New Baseten Together AI NVIDIA 4 replies ○ sign in
Three engines, one hardware roofline, and the difference between them is which part of the roofline each one reaches first on your workload. The decision is a table, and the table has reversal conditions.
28
When does offloading the KV cache to CPU memory or NVMe beat recomputing it? ▼ hard New NVIDIA Perplexity SGLang 4 replies ◆ premium
A paused conversation holds gigabytes of cache that its next turn will need in thirty seconds. Recompute costs prefill; reload costs bandwidth. The comparison is one division per tier, and the answer flips with prompt length.
07
Your model decodes at a tenth of its bandwidth bound at batch one. Explain the gap. ▼ hard New Fireworks AI Together AI NVIDIA 4 replies unlocked
The bound assumes weights stream contiguously and nothing else costs time, and at batch one both assumptions fail badly. Four terms that make up the gap, why sparse models suffer most, and the two fixes that recover most of it.
11
The model ships in FP8. Should you requantize to four bits to fit more of it on fewer GPUs? ▼ medium New Fireworks AI Together AI Baseten 4 replies ○ sign in
Halving the weights halves the GPU count and moves the quality evaluation onto you. What the format actually buys in memory and in speed, why those are different questions, and the evaluation that has to run before it ships.
29
You need a released bf16 model at half the footprint. Walk through quantizing it yourself. ▼ hard New Fireworks AI Together AI Baseten 4 replies ◆ premium
The conversion is the easy part and the evaluation is the project. What calibration data does and why yours should come from production, the per-category evaluation that catches what an average hides, and the parts of the model that should not be quantized at all.
37
The engine logs say CUDA graphs are disabled for your deployment. Does it matter? ▼ hard New Fireworks AI Together AI NVIDIA 4 replies ◆ premium
Graphs remove the per-kernel launch cost, which is a large share of a decode step at low batch and a small one at high batch, so whether it matters is a question about your operating point. The arithmetic, the four reasons an engine turns them off, and the measurement that settles it.
09
A kernel does 4 TFLOP and moves 40 GB in one call. On an H100, is it memory-bound or compute-bound? ▼ medium New NVIDIA Fireworks 4 replies unlocked
Two divisions decide it: the kernel's FLOPs per byte against the card's ridge point. The worked case, the profiler counters that give the inputs, and what to change once you know which wall you hit.
19
How long does an all-reduce of a 70B model's gradients take on eight GPUs? ▼ medium New NVIDIA Meta 4 replies ○ sign in
The ring all-reduce formula, 2(n−1)/n times the buffer over the link bandwidth, applied to 141 GB of bf16 gradients on an NVLink node: 0.27 s by the formula, and the two reasons the measured number is different.
20
How many tokens per second can one B200 decode for a 70B model? ▼ medium New NVIDIA CoreWeave 4 replies ○ sign in
8 TB/s over the bytes per step: 57 tokens per second single-stream in bf16, 113 in fp8, 227 in fp4, and the batch curve on a single card that now holds the whole model. Where the curve bends and what caps it.
27
Estimate the activation memory for training a 70B model at an 8k sequence length ▼ hard New Meta NVIDIA 4 replies ◆ premium
Activations are the term the 16-bytes-per-parameter rule leaves out, and at 8k context they are bigger than the weights: about 34 × sequence × hidden bytes per layer, 2.3 GB, 182 GB per micro-batch across 80 layers. The chain, what FlashAttention removes, and what checkpointing does to the number and the compute.
01
What does NCCL actually do when you call all-reduce, and how does it decide which algorithm to use? ▼ easy ★ Essential New NVIDIA Meta 4 replies unlocked
Before the first byte moves, the library has already discovered the machine's topology, built rings and trees over it, split the work into parallel channels and chosen a wire protocol by message size. What each of those four decisions is for, and the environment variables that let you see and override them.
02
Name the four collectives a training job uses, say what each moves, and match them to the parallelism that needs them. ▼ easy ★ Essential New NVIDIA Google 4 replies unlocked
Four operations cover almost everything a distributed training step sends. What each one does to the data, the bytes each rank moves for a message of size S across N ranks, and which parallelism strategy generates which, worked through for a 70B model so the numbers are concrete rather than symbolic.
03
InfiniBand or RoCE version 2 for a new GPU training cluster. Make the call and say what would reverse it. ▼ medium ★ Essential New NVIDIA Crusoe xAI 4 replies unlocked
Both carry RDMA at the same line rate. One arrives lossless because the transport was designed that way, the other becomes lossless only if a set of switch settings is correct on every port. What that difference costs in operations, what it saves in money and hiring, and the fleet size where the answer flips.
04
What is RDMA, and why can a training cluster not just use TCP at 400 gigabits per second? ▼ easy New NVIDIA Crusoe 4 replies unlocked
A single core moves a few gigabits per second of TCP once copies and interrupts are counted, so filling one 400 gigabit link would take most of a server's cores doing nothing but networking. What RDMA removes, what a queue pair actually is, and where the GPU fits when the data never belongs to the host at all.
05
Explain a rail-optimized fabric. What does it buy over a plain fat tree, and how many switches does one need? ▼ medium ★ Essential New NVIDIA Meta CoreWeave 4 replies unlocked
Wire GPU number three of every node to the same switch and the collective that matters most never leaves it. The rank mapping that makes the layout pay, the switch and cable count for a thousand GPUs worked out, and what happens to a job whose ranks are assigned in the wrong order.
08
What is incast, why does it break an RDMA fabric, and what do ECN and priority flow control each do about it? ▼ hard New xAI NVIDIA Meta 4 replies unlocked
Thirty-two senders aimed at one port fill a switch buffer in microseconds, and what happens next depends on which of two mechanisms fires first. The buffer arithmetic, why the slow mechanism is the safe one and the fast one is the dangerous one, and the counters that tell you which is happening.
09
When does GPUDirect Storage actually help, and when is it just a more complicated read? ▼ medium New NVIDIA WEKA 4 replies unlocked
Reading straight from NVMe into GPU memory skips a copy through host memory, which is worth a lot for one workload and nothing for another. The bandwidth arithmetic for both paths, the two cases where the bounce is the bottleneck, and the far more common case where the CPU has to touch the data anyway.
14
Map a 1,024-GPU job's parallelism onto the hardware. Which dimension goes on NVLink, which on the fabric, and what does a wrong order cost? ▼ hard New NVIDIA Meta 4 replies ○ sign in
Three parallelism dimensions, two link speeds eighteen times apart, and one mapping that makes the job fast. The traffic each dimension generates per step, why the loudest one must stay inside the node, and the arithmetic showing what a wrong rank order costs before anyone notices.
15
Which NCCL environment variables are worth knowing, what does each change, and which should you almost never set? ▼ medium New NVIDIA CoreWeave 4 replies ○ sign in
Four of them diagnose, four configure the machine correctly, and several popular ones override tuning that was already right. Which is which, the symptom each addresses, and why a cluster whose job scripts carry a dozen of these has a configuration problem rather than a tuning one.
17
Why do large training clusters provision roughly 400 gigabits per second per GPU rather than more or less? ▼ medium New NVIDIA Meta xAI 4 replies ○ sign in
The number comes from one requirement: the gradient reduction has to finish inside the backward pass that produces it. Working that requirement backward gives a bandwidth per GPU, and the answer lands near the port speed the industry ships, which is not a coincidence.
24
A mixture-of-experts model does an all-to-all twice per layer. What does that demand of the fabric, and what changes on a rack-scale system? ▼ hard New DeepSeek NVIDIA 4 replies ◆ premium
Two all-to-alls per layer across sixty layers is over a hundred collectives on the critical path of every decode step, each one small and latency-bound. The per-GPU volume worked out, the time on each link type, and why a single large NVLink domain changes the design rather than just improving it.
28
How would you know GPUDirect RDMA is not being used, given that nothing reports an error when it is disabled? ▼ medium New NVIDIA Lambda 4 replies ◆ premium
It fails open: the data still moves, the job still runs, and everything is about half as fast. The one log line that settles it, the four reasons the library declines to use it, and why the bandwidth signature is close enough to exactly half to be diagnostic on its own.
02
MIG, time-slicing and MPS all let several jobs share one GPU. What is the difference, and when would you pick each? ▼ easy ★ Essential New CoreWeave Nebius NVIDIA 4 replies unlocked
Hardware partitions, shared SMs, and context switching are three different promises about isolation and waste. The slice arithmetic for an H100, what each mode gives up, and the one rule about tenants that decides most of it.
06
Two identical 8-GPU jobs get 8 GPUs each. One runs at half the speed of the other. What did the scheduler do, and how do you stop it? ▼ medium New NVIDIA CoreWeave 4 replies unlocked
Same job, same GPU count, half the speed: the slow one was split 4 and 4 across two nodes and its tensor-parallel all-reduces run at fabric speed instead of NVLink speed. The bandwidth arithmetic, the per-parallelism placement rule, and the scheduler policy that enforces it.
09
What does the NVIDIA GPU Operator actually install on a node, and what is Node Feature Discovery doing underneath it? ▼ easy New NVIDIA CoreWeave 4 replies unlocked
A GPU node needs a driver, a container toolkit, a device plugin, a metrics exporter and a validator, in that order, and each needs to know which nodes it belongs on. The operator's component chain, the labels NFD writes to trigger it, and the bring-up arithmetic that rules out doing it by hand.
19
Kubernetes device plugins versus Dynamic Resource Allocation: what changes for GPU scheduling, and what would you adopt in 2026? ▼ medium New NVIDIA Google 4 replies ○ sign in
Device plugins count GPUs and know nothing else; Dynamic Resource Allocation lets a pod ask for devices by attribute and share them by claim. What the two models can and cannot express, what DRA changes for MIG, NVLink domains and multi-node gangs, and a migration stance for a fleet that runs both.
02
What is an XID error, which ones mean the hardware is bad, and which ones mean somebody's kernel has a bug? ▼ easy ★ Essential New NVIDIA CoreWeave Lambda 4 replies unlocked
The driver logs a numbered code when something goes wrong on a GPU, and the number tells you whether to retry the job, drain the node or file a hardware return. The codes worth memorizing in three groups, the action each implies, and the automation that turns a log line into a drained node.
03
With DCGM available on every node, what do you actually collect, what do you alert on, and what do you deliberately ignore? ▼ easy ★ Essential New NVIDIA CoreWeave Microsoft 4 replies unlocked
The exporter offers hundreds of fields and about a dozen of them change decisions. The health fields that predict failure, the performance fields that tell you whether work is happening, the one everybody alerts on that means nothing, and the diagnostic levels with the time each takes.
14
What do NVLink errors look like in telemetry, when is a link degrading rather than broken, and when do you drain the node? ▼ hard New NVIDIA Meta 4 replies ○ sign in
A link that corrupts occasionally does not fail; it retransmits, and each retransmission costs bandwidth that shows up as a slow job with no error anywhere. The counters that reveal it, the arithmetic turning a replay rate into lost bandwidth, and the threshold where watching becomes draining.
20
Walk me through reading a collective flight-recorder dump. What is in it, and how do you find the rank that caused a hang? ▼ hard New Meta NVIDIA 4 replies ○ sign in
Each rank keeps a ring buffer of its recent collectives, and on a timeout every rank dumps it. Comparing sequence numbers across a thousand of those buffers finds the missing participant in one pass, and the three patterns those comparisons produce point at three different causes.
13
Design a KV cache tier across GPU memory, host memory and NVMe so prefixes survive across sessions. Bandwidths, and when it pays. ▼ hard New NVIDIA Perplexity 4 replies ○ sign in
A 70B model's KV cache is 320 KB per token, and a 30k-token conversation is 10 GB that GPU memory cannot keep between turns. The three tiers with their bandwidths, the break-even where reloading a prefix beats recomputing it, the eviction and lookup design, and the traffic shape where the tier is worth its complexity.
28
Design prefill and decode disaggregation for a 70B serving fleet: pool sizes, KV transfer, routing, and when to turn it on. ▼ hard New NVIDIA Fireworks Perplexity 4 replies ◆ premium
Prefill is compute-bound and decode is bandwidth-bound, and one GPU running both does each worse. The pool ratio from the traffic's prompt and output lengths, the KV transfer cost per request and the link it needs, the router that pairs a prefill worker with a decode worker, and the load at which it beats a monolith.
31
Design a training cluster for a one-trillion-parameter MoE. Size it, choose the parallel layout, and map it onto the fabric. ▼ hard ★ Essential New OpenAI Anthropic xAI ◆ premium
The first question is whether 1T is total or active, because storage follows one and compute follows the other. The state budget, the layout that falls out of it, why expert parallelism belongs inside NVLink, and the two failure modes a dense-model plan does not have: router imbalance and all-to-all congestion.
08
You are handed a 40 GB kernel trace. Write the parser that turns it into per-instruction access statistics without running out of memory. ▼ medium ★ Essential New NVIDIA 4 replies unlocked
The classification is a small function; the parser around it is where the problem is. Generators rather than lists, aggregation bounded by the number of distinct instructions rather than by the file, a top-k that never sorts the whole thing, and the malformed lines that stop the job at hour three.
13
Simulate a ring all-reduce on arrays. Verify both the result and the bytes each rank sends. ▼ medium New NVIDIA Google 4 replies ○ sign in
Two phases of N-1 steps each, a chunk index that rotates with the step, and a byte count that should come out to exactly the textbook formula. The implementation, the verification against a direct sum, and the measured traffic matching 2(N-1)/N times the message at every rank count tested.
22
Write a thread-safe counter in C++ that many threads increment. What are the memory-ordering and layout choices? ▼ medium New NVIDIA Cerebras 4 replies ◆ premium
One atomic increment is correct and does not scale, because every thread fights for the same cache line. The ordering that is sufficient and why the default is stronger than needed, the padding that removes contention entirely, and the arithmetic showing when sharding is worth it.
23
Why is std::unordered_map rarely the fastest hash map, and what would you use instead? ▼ medium New Cerebras NVIDIA 4 replies ◆ premium
The standard forces one heap allocation per entry and a pointer dereference per lookup, and that is a specification requirement rather than a library weakness. The measured cost on this machine, the open-addressing layout that removes it, and the two cases where the standard container is still the right pick.
24
Implement a work-stealing deque. Why do the owner and the thieves take from opposite ends? ▼ hard New NVIDIA Anyscale 4 replies ◆ premium
Two ends, two access patterns, and one design choice that removes almost all contention. Why the owner takes the newest task and thieves take the oldest, what that does to cache locality, and the measurement showing why this cannot be demonstrated in Python at all.
28
Write a tiled matrix multiply and explain the cache effect. Then say honestly what your benchmark actually measured. ▼ medium New NVIDIA Google 4 replies ◆ premium
The tiling is ten lines and the theory is a working-set calculation. The interesting part is that the obvious benchmark does not measure what it appears to, because the library underneath is already tiled, and knowing which effect your numbers contain is the skill being tested.
03
Why do you want to work in AI infrastructure, and why now? ▼ easy ★ Essential New OpenAI Anthropic NVIDIA 4 replies unlocked
The answer that fails is a compliment to the company. The answer that lands is a thesis about where the constraint sits, tied to something you have actually done, and it works whether you are coming from distributed systems, kernels, hardware or SRE.
07
Tell me about a performance win you are proud of. How did you measure it? ▼ easy New NVIDIA Fireworks AI Together AI 4 replies unlocked
A speedup number with no method behind it is worth nothing, and experienced interviewers stop listening at the number and start asking about the baseline. What makes a performance claim survive scrutiny, and the four ways a real one turns out to be smaller than it looked.
13
How do you keep current in a field that changes every few months? ▼ easy New NVIDIA Anthropic CoreWeave 4 replies ○ sign in
Naming three newsletters answers nothing. The answer that works has a filter, a source of ground truth that is not a blog post, and one thing you reproduced yourself, because that is the only evidence that you learn rather than accumulate.
18
A vendor says your configuration is the problem. You think it is their firmware. How do you proceed? ▼ medium New CoreWeave Lambda Labs NVIDIA 4 replies ○ sign in
Arguing is the losing move and so is capitulating. What moves a vendor case is a reproducer small enough for their engineer to run, counters that isolate their layer from yours, and a bisection that removes your configuration as a variable entirely.
20
What do you think is the most underrated problem in AI infrastructure right now? ▼ hard New Anthropic OpenAI NVIDIA 4 replies ○ sign in
This is a test of whether you have a position you can defend, not of which problem you pick. What makes a thesis defensible, two worked examples with the arithmetic behind them, and the counterargument you have to be able to state before the interviewer does.
23
You think your tech lead's design is wrong. How do you handle it? ▼ medium New Google NVIDIA Meta 4 replies ◆ premium
Being right is a small part of this. The move that works is converting an opinion into a document with the tradeoff written out, and then committing fully once the decision is made. The two cases where escalating is correct, and how to do it without a fight.
30
Where do you think AI infrastructure is going over the next five years? ▼ expert ★ Essential New Anthropic OpenAI NVIDIA 4 replies ◆ premium
Four claims that are defensible from arithmetic available today, each with the counterargument that could sink it, and what each one implies about the work. Dated to 2026, because a thesis with no date is not a prediction.
More from the tracks NVIDIA's loop tests The highest-signal questions across NVIDIA's core tracks.
8 questions · 1 unlocked for you
06
Is LLM decode memory-bound or compute-bound? Show me the arithmetic that decides it. ▼ medium ★ Essential New OpenAI Anthropic Together AI 4 replies unlocked
At batch 1 a decode step reads every weight byte to do two FLOPs with it. The intensity is the batch size, the ridge is about 295, and the KV cache puts a ceiling on how far batching can push you. The full chain for a 70B model on H100.
14
You have trained on GPUs. What transfers to training on TPUs, and what do you have to relearn? ▼ hard ★ Essential New Google Google DeepMind 4 replies ○ sign in
The roofline, the sharding arithmetic and the parallelism trade-offs transfer unchanged, and the numbers are in the same units. What changes is who writes the kernels, how shapes must behave, and the interconnect topology you shard against. The v6e and v7 numbers worked through against an H100.
15
Explain how a systolic array computes a matrix multiply. Why is it efficient, and why does it want big square tiles? ▼ hard New Google Groq Cerebras 4 replies ○ sign in
An N × N grid of multiply-accumulate cells holds a weight tile still and pumps activations through it, one row per cycle, with no register file or instruction fetch in the inner loop. The cycle count, the fill and drain, and the formula that says why a 128-row tile wastes two thirds of the array.
16
Would you move a cost-sensitive training and serving fleet from H100 to Trainium2? What do you gain, and what do you have to plan for? ▼ hard New AWS 4 replies ○ sign in
A Trn2 instance carries 16 Trainium2 chips with 1.5 TB of HBM and more dense bf16 FLOPS than an 8-GPU H100 node. Whether the cheaper FLOPS reach your workload depends on the Neuron compiler, the kernels you do not have yet, and an MFU number you must measure rather than assume.
17
The 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? ▼ hard New AMD Microsoft Oracle 4 replies ○ sign in
192 GB and 5.3 TB/s per card beat the H100 on both memory numbers at a lower hourly price. Whether that wins in production is a question about achieved bandwidth and ROCm kernel coverage, and there is a short list of numbers to run before you sign the order.
18
Why are Groq and Cerebras so fast at batch 1, and what does that speed cost at scale? ▼ hard New Groq Cerebras 4 replies ○ sign in
Put the weights in on-chip SRAM and the HBM wall disappears: tens of terabytes per second per chip, and a batch-1 step limited by the pipeline rather than the memory. The price is capacity: a 70B needs hundreds of chips per replica, and the cost per token depends on keeping every one of them busy.
24
Why does the batch size change the TFLOPS a GEMM achieves? Draw the curve and explain each region. ▼ medium New Together AI Baseten 4 replies ◆ premium
A linear layer at batch 1 runs at a third of a percent of peak and at batch 4,096 near 70%. The curve between has two regions and a sawtooth: the memory-bound slope where intensity equals M, the plateau where the tensor cores are the limit, and the dips where the tile count does not divide the SM count.
26
What does ECC on a GPU cost you, and why do you keep it on across a fleet? ▼ medium New CoreWeave Lambda Microsoft 4 replies ◆ premium
On an HBM part ECC costs almost nothing; on a GDDR part about 6% of capacity and bandwidth. What it buys is the difference between a corrected bit and a silently flipped exponent that turns 1.0 into infinity across a 512-GPU all-reduce. The mechanisms, the signals, and the expected-loss arithmetic.
Go deeper on the topics NVIDIA's loop tests The tracks that map to a NVIDIA AI Infrastructure Engineer loop, ordered easy to hard.
The concepts NVIDIA's AI Infrastructure Engineer loop assumes you know The vocabulary and mental models behind NVIDIA's questions, from our curriculum. Start with the foundations free; the deeper, interview-defining ideas are part of premium.
⚡ KERNELS & COMPILERSFoundational
CUDA Programming Model CUDA 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 Sign in
Memory Coalescing A 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 🔒 Premium
Shared Memory and Bank Conflicts Shared 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 🔒 Premium
Occupancy and Register Pressure Occupancy is the fraction of an SM's 64 warp slots that are resident, and it is capped by the 65,536 registers and 228 KB of shared memory each block consumes. It decides how much memory latency the hardware can hide for free, but the fastest kernels on a GPU routinely run at 25 percent, so the interview skill is knowing when to raise it and when to stop. 🧩 GPU & ACCELERATOR ARCHITECTUREFoundational
GPU Execution Model A GPU hides memory latency with parallelism instead of caches: thousands of threads in flight, scheduled in warps of 32, pinned to streaming multiprocessors that switch between warps for free whenever one stalls. Every performance conversation in an AI infra loop, from occupancy to why decode is slow, rests on this one mechanism. Foundational
GPU Memory Hierarchy A 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 Sign in
Tensor Cores and Matrix Units Tensor cores are fixed-function units that compute a small matrix multiply-accumulate per instruction, and they are where almost all of a modern GPU's FLOPS live: 989 dense bf16 TFLOPS on an H100 against about 67 from the general-purpose lanes. Only dense, well-shaped matrix multiplication at a supported precision can use them, which is why GEMMs reach peak and nothing else does, and why precision choices are throughput choices. Advanced 🔒 Premium
Memory-Bound vs Compute-Bound Kernels Every kernel is limited by one of two walls: how fast bytes arrive from HBM, or how fast the tensor cores can multiply. Which wall applies is decided by arithmetic intensity against the ridge point, and the two regimes need opposite fixes. Decode, LayerNorm and softmax are memory-bound; prefill GEMMs are compute-bound; the interview question is which one you are looking at and what you would do about it. 🚀 INFERENCE & SERVINGFoundational
Prefill vs Decode An LLM request runs in two phases with opposite hardware profiles: prefill reads the whole prompt in one compute-bound pass and decides time to first token, decode emits one token per forward pass and is bound by memory bandwidth. Every serving decision, from batch size to which GPU to buy to whether to split the two phases across machines, follows from that split. Foundational
The KV Cache The KV cache stores each token's attention keys and values so decode never recomputes them, turning a quadratic cost into a linear one at the price of memory that grows with every token in every concurrent sequence. Its size, 128 KB per token for Llama 3.1 8B and 320 KB for 70B in bf16, is what caps concurrency and context on a given GPU, so it decides batch size, replica count and whether a model fits at all. Core Sign in
Continuous Batching Continuous batching schedules at the granularity of a single decode step instead of a whole request, so a finished sequence's slot is refilled on the next iteration rather than when the longest request in the batch ends. It is the scheduling idea that turned LLM serving from a padded, half-idle GPU into one that stays full, and it decides how the engine's scheduler, memory manager and latency SLOs interact. Advanced 🔒 Premium
PagedAttention PagedAttention stores the KV cache in fixed-size blocks scattered across HBM and maps each sequence's logical positions to physical blocks through a block table, the same trick an operating system uses for virtual memory. It removes the reservation and fragmentation waste of contiguous allocation, lets blocks be shared between sequences, and is why an engine can decide admission by counting free blocks. 💻 CODING FOR INFRAFoundational
The GPU Credit Scheduler Pattern The most widely reported coding problem in AI infrastructure loops is a small scheduler: accounts hold credits, jobs arrive with a cost and a priority, and you must decide which jobs run, in what order, without letting any account overspend, then extend it under follow-ups (refunds, reservations, concurrency limits, fairness). It is not a trick question; it is a test of whether you can model state cleanly, pick the right data structures, keep invariants under mutation, and talk about complexity while typing. This page works the problem from the first line to the fourth follow-up, with the code, the invariants, and the derivations. Core Sign in
Rate-Limiting Algorithms A rate limiter answers one question, 'may this request proceed now?', and the three classic algorithms answer it with different shapes of fairness and memory: the token bucket allows bursts up to a capacity and refills at a rate, the leaky bucket smooths output to a fixed rate, and sliding windows count recent requests exactly or approximately. AI platforms limit in tokens as well as requests, per tenant, across many gateways, which adds two twists: a request's cost is unknown until it finishes, and the counters must be shared. This page derives each algorithm, implements the token bucket correctly, and covers both twists. Advanced 🔒 Premium
Batching Queues and Backpressure Write a request batcher is the coding round's version of the serving engine's scheduler: requests arrive one at a time, the GPU wants them in groups, and the batcher decides when a group is full enough to send without holding anyone too long or accepting more than it can hold. The two knobs are the maximum batch size and the maximum wait, the invariant is a bounded queue, and the follow-ups (priorities, cost-aware batching, cancellation, bounded in-flight batches) are the ideas the real engines carry. This page implements the batcher in asyncio, derives what each knob buys, and walks the follow-ups. Advanced 🔒 Premium
Interval Merging and Utilization Logs Given busy intervals per GPU, when was the whole cluster idle? What was the utilization per hour from a log of start and stop events? Which jobs overlapped? These are the interval problems of the infrastructure coding screen, and they share one tool: sort the endpoints and sweep. The sweep line turns every variant into a single pass with a counter, the sort is the only thing that costs more than linear time, and the edge cases (touching intervals, zero-length events, an unterminated start) are where candidates lose the round. This page works the standard problem and its relatives with code, tests and the complexity derivation. 🔌 NETWORKING & STORAGEFoundational
NCCL and Collective Algorithms NCCL is the library every PyTorch collective lands in, and its choice of ring or tree, channel count and protocol decides whether an all-reduce runs at fabric speed or at a third of it. Knowing what NCCL_DEBUG=INFO prints, and which environment variable changes which decision, is the difference between tuning a cluster and guessing at it. Core Sign in
RDMA, InfiniBand and RoCEv2 Training across nodes moves hundreds of gigabytes per step, and a CPU-driven TCP stack cannot feed a 400 Gb/s link. RDMA lets a NIC write straight into a remote GPU's memory with no kernel and no copies, and it runs over two fabrics: InfiniBand, which is lossless by design, and RoCEv2, which is Ethernet made lossless by configuration. The choice is operational as much as technical, and the numbers that decide it are per-GPU bandwidth, the collective's volume, and who will debug a pause storm at 3 a.m. Advanced 🔒 Premium
Rail-Optimized and Fat-Tree Fabrics A GPU cluster's network is built from two ideas: a fat tree (Clos) that gives every node a path to every other node with a chosen amount of oversubscription, and rail optimization, which wires GPU i of every node to the same leaf switch so the collectives that dominate training stay one hop away. Sizing one is arithmetic on port counts, and the interview question is usually that arithmetic: how many switches, what oversubscription, and where the NVLink domain ends and the fabric begins. Advanced 🔒 Premium
Congestion Control for AI Fabrics Collective traffic is the worst case a network can see: hundreds of senders transmit to the same receiver at the same instant (incast), every flow is large and long-lived, and RDMA cannot tolerate a dropped packet. Congestion control is the set of mechanisms (PFC, ECN with DCQCN, adaptive routing, packet spraying) that keep queues from overflowing without stalling the fabric. On plain Ethernet a busy all-reduce can fall to about 60% of link rate; with a tuned control loop it holds above 90%. Reading the counters that show which one you have is the on-call skill. Where to apply, and official NVIDIA resources Straight from NVIDIA: open roles and the company's own hiring guidance. Prep here, then apply there.
External links to NVIDIA's own pages. Roles and processes change; always confirm on the official site.
NVIDIA INTERVIEW FAQ
What is the NVIDIA AI Infrastructure Engineer interview process? ▲
Deep Learning Software Engineer / GPU and HPC Infrastructure Engineer / DGX Cloud Performance Engineer. Typical loop: 4 to 8 weeks typical; long loops are a recurring complaint. Stages: Recruiter call → Phone screens → Take-home (team-dependent) → Onsite. Key focus: Expertise in the team's domain, reported as the dominant filter. Compiled from public reports; loops change over time, so confirm the exact rounds with your recruiter.
Does NVIDIA hire AI infrastructure engineers? ▼
Yes, in many titles: AI and ML Infrastructure Software Engineer, GPU Clusters (new-grad 2026); DGX Cloud Performance Engineer; Senior GPU and HPC Infrastructure Engineer; Senior HPC AI Cluster Engineer; Deep Learning Software Engineers for LLM performance, TensorRT and TensorRT-LLM; cuDNN and CUTLASS library engineers; Developer Technology engineers; and Triton Inference Server systems engineers. India postings include a DGX Cloud Performance Engineer in Pune and Site Reliability Engineers in Bengaluru.
What does the NVIDIA AI infrastructure interview test? ▼
Reported loop: a recruiter call; one or two phone screens of 45 to 60 minutes mixing C++ trivia and fundamentals with graph and array problems; a team-dependent take-home in some loops (one 2022 report); then four to six onsite rounds of about an hour with domain-specific content (systems architecture, performance, team-specific coding) and a behavioral round, sometimes a hiring-manager or director round. Candidates report that expertise in the team's domain matters more than LeetCode.
What CUDA questions does NVIDIA ask? ▼
Reported by candidates and prep aggregators: explain memory coalescing and its effect on bandwidth; shared-memory bank conflicts and how to avoid them; warp divergence and SIMT execution; implement a parallel reduction and explain synchronization; optimize a GEMM with tiling and coalescing; profile a CUDA application with Nsight Compute and Nsight Systems; kernel fusion versus launch overhead. Treat individual items as partial evidence.
What is the NVIDIA AI infrastructure engineer salary? ▼
Posted 2026 new-grad bands: L2 $120K to $189,750 and L3 $148K to $235,750 for AI and ML Infrastructure Software Engineer, GPU Clusters; L3 $148K to $235,750 and L4 $184K to $287,500 for DGX Cloud Performance Engineer. A Glassdoor copy of a senior HPC AI cluster engineer posting showed $176K to $334K. NVIDIA's ladder runs IC1 to IC6, and equity has driven its total compensation for years.
How long is the NVIDIA loop? ▼
Four to eight weeks typical, with long and drawn-out loops a recurring candidate complaint. No first-party statement on AI use in interviews was found.
Walk into your NVIDIA AI Infrastructure Engineer interview ready Unlock every AI infra interview answer, ordered easy to hard, plus the full concept curriculum, for 6 months. One payment, no auto-renewal. Free questions and concepts in each track, no card needed to start.
Or create a free account to unlock more free answers per topic.
Other AI Infrastructure Engineer interviews to prep Companies whose loops test the same tracks as NVIDIA's.
Independent and not affiliated with NVIDIA. All trademarks belong to their owners.