Jane Street CUDA, Triton & GPU Performance interview questions

Q: How should I prepare for the Jane Street CUDA, Triton & GPU Performance round?

Work the cuda, triton & gpu performance questions on this page in order, study the underlying concepts linked below, then review Jane Street's full interview process for how this round fits the loop. For Jane Street specifically, interviewers weight this: Low-level GPU knowledge: PTX, SASS, warps, cooperative groups, Tensor Cores, memory hierarchy.

Q: Does Jane Street hire inference engineers?

Yes. Jane Street job postings seen on 16 October 2026 include Machine Learning Performance Engineer. The same team posts under different titles from one req to the next, so judge a posting by its listed duties.

Q: What is the Jane Street inference interview process?

Jane Street publishes partial guidance about its interviews: Application review; Technical interview over Zoom; Final rounds in person. It is company-wide guidance, not specific to inference roles, so confirm the inference loop with your recruiter.

CUDA, Triton & GPU Performance is a core part of the Jane Street Inference Engineer loop. The memory hierarchy in numbers, coalescing and bank conflicts, occupancy, Hopper TMA and WGMMA, Blackwell tensor memory, decode GEMV versus prefill GEMM, fused MoE and dequant kernels, CUDA graphs, and reading Nsight like an interviewer. Questions tagged to Jane Street come first. After them, the top-voted questions from the CUDA, Triton & GPU Performance track. Each answer is written the way a senior serving engineer would argue it at a whiteboard, numbers included.

WHAT JANE STREET LOOKS FOR HERE · Low-level GPU knowledge: PTX, SASS, warps, cooperative groups, Tensor Cores, memory hierarchy. See the full Jane Street interview process →

CUDA, Triton & GPU Performance questions tagged to Jane Street

2 questions · 1 unlocked for you

Open full answer →

More CUDA, Triton & GPU Performance questions for Jane Street's loop

Sorted by reader votes. These are not tagged to Jane Street; they cover the cuda, triton & gpu performance ground its Inference Engineer loop is known to test.

13 questions · 4 unlocked for you

Open full answer →

Concepts behind Jane Street's CUDA, Triton & GPU Performance round

What these answers take for granted. Foundation concepts are free to read; the advanced ones are Premium.

Foundational

The GPU Memory Hierarchy in NumbersAn H100 SXM has 33 MiB of registers, up to 29 MiB of shared memory, a 50 MB L2 and 80 GB of HBM3 at 3.35 TB/s. Knowing those capacities decides which inference tensors can be reused on chip, why decode streams every weight from HBM each step, and when a benchmark is secretly measuring L2.

Advanced

🔒 Premium

Bytes in Flight: Little's Law for GPU MemoryLittle's law says sustained memory bandwidth equals the bytes a kernel keeps outstanding divided by memory latency. On an H100 SXM that is about 2 MB across the GPU, or 15 KB per SM, and it decides whether a bandwidth-bound decode kernel can reach peak at all, before any question of coalescing or caching.

Advanced

🔒 Premium

CUDA Graphs and Launch Overhead in DecodeA decode step is hundreds of small kernels, and once tensor parallelism or a small model shrinks the GPU time per step, the CPU cost of launching them becomes the step time. CUDA graphs record the whole step once and replay it with one launch, which is why every serving engine captures decode graphs per batch-size bucket.

Explore the full curriculum →

JANE STREET CUDA, TRITON & GPU PERFORMANCE FAQ

What CUDA, Triton & GPU Performance questions does Jane Street ask in interviews?▲

Jane Street's Inference Engineer loop draws cuda, triton & gpu performance questions such as "A kernel uses 128 registers per thread and 48 KB of shared memory per block. How many warps are active per SM on an H100?", "You are about to report one number for your new kernel in a design review. How do you measure it so the number survives scrutiny?", "Your decode GEMV kernel reaches only about half of HBM bandwidth on an H100. Find the reasons and get it above 80%.". The memory hierarchy in numbers, coalescing and bank conflicts, occupancy, Hopper TMA and WGMMA, Blackwell tensor memory, decode GEMV versus prefill GEMM, fused MoE and dequant kernels, CUDA graphs, and reading Nsight like an interviewer. Every question in the topic is listed below from easiest to hardest, each with a worked answer.

How should I prepare for the Jane Street CUDA, Triton & GPU Performance round?▼

Does Jane Street hire inference engineers?▼

What is the Jane Street inference interview process?▼

Other Jane Street interview rounds

The other tracks Jane Street's Inference Engineer loop tests.

Jane Street attention kernels→Jane Street inference physics→

Prep the whole Jane Street Inference Engineer loop

CUDA, Triton & GPU Performance is one round. Unlock every answer across Jane Street's full loop with Premium. Free questions in every track let you try the depth first.

Every worked answer and derivation in all 12 topics, every concept page, and the full practice-test bank.

6 months · One payment · No auto-renewal

Unlock all answers · ₹2,000 / $25 Full Jane Street interview guide

Independent and not affiliated with Jane Street. All trademarks belong to their owners.

Jane Street CUDA, Triton & GPU Performance interview questions

CUDA, Triton & GPU Performance questions tagged to Jane Street

A kernel uses 128 registers per thread and 48 KB of shared memory per block. How many warps are active per SM on an H100?

You are about to report one number for your new kernel in a design review. How do you measure it so the number survives scrutiny?

More CUDA, Triton & GPU Performance questions for Jane Street's loop

Your decode GEMV kernel reaches only about half of HBM bandwidth on an H100. Find the reasons and get it above 80%.

In a decode GEMV over a row-major weight matrix, which thread-to-element mapping coalesces, and why does the transposed one not?

Your shared-memory transpose tile has bank conflicts. Do the bank arithmetic, apply the padding fix, and say when padding is the wrong fix.

Why do tensor core GEMMs use tiles like 128 by 128 by 64, and what limits the tile size on Hopper?

A GEMM launches 140 output tiles on a 132-SM H100. What does that cost, and how do you fix it?

Write a fused RMSNorm kernel in Triton and count its memory traffic against the unfused PyTorch version.

Fused MoE kernels sort tokens by expert before the GEMM. Why, and what does a grouped GEMM do that a loop over experts cannot?

In a W4A16 GEMM, where does the INT4 unpacking happen, and how does the kernel keep that work hidden under the memory loads?

Here is an Nsight Systems timeline of one decode step. Find the launch gaps, synchronizations and CPU stalls, and price each one.

Hopper replaced Ampere's cp.async and mma.sync with TMA and WGMMA. What changed in each, and how does it change the way you design a kernel?

Blackwell added tensor memory and 2-CTA MMA. What problem does each one solve, and what changes in a GEMM kernel written for it?

Could you keep the KV cache of hot sequences resident in the H100's 50 MB L2 during decode, and would it be worth it?

A kernel runs 20% slower on an H200 than on an H100, even though the H200 has 43% more memory bandwidth. How can that happen?

Concepts behind Jane Street's CUDA, Triton & GPU Performance round

Other Jane Street interview rounds

Prep the whole Jane Street Inference Engineer loop