Member of Technical Staff, GPU Kernels

About the Role

We build the fastest GPU compiler in the world. Most compilers have to preserve correctness at every transform, constraining how far they can search, while we prove correctness at the end instead, allowing us to search a far wider space, with agents, with RL, with anything that works and still guarantee the result. It's why we hold #1 on NVIDIA's own kernel benchmark across hundreds of production kernels.

We're hiring a Member of Technical Staff for GPU Kernel Engineering to establish and push the envelope on what the hardware can actually do before any search ever runs. To this extent, you'll write and hand-optimize kernels that execute during pre-training, post-training and inference, across NVIDIA, AMD, TPU and Trainium. You'll then take those learnings and convert them into structures the compiler can search: instruction sequences, scheduling strategies, cost signals, to ensure that the ceiling you worked for by hand can become the floor for the search finding kernels everywhere else.

You'll have access to better tooling than anywhere else because we own the stack all the way down to the ISA, allowing us to create kernels that others can't even express. For example, our custom LLVM backend emits cubins directly, without a PTX handoff or ptxas doing our scheduling, letting us write kernels that have no expression in CUDA at all. Our team understands the hardware better than anyone else, to this extent we've built a bit-exact software model of Blackwell's tcgen05.

You'll be writing kernels that ship immediately on runs such as pre-training AlphaFold v3 at 3.4× the throughput, post-training robotics models on Trainium or running our custom RL rollout engine on TPU at multi-100B parameter scale.

What You'll Do

• You'll write and hand-optimize kernels for real workloads to find the ceiling before the search goes looking for it
• You'll profile at the microarchitectural level: SM and CU utilization, warp stalls, memory bank conflicts, register pressure, instruction throughput, occupancy tradeoffs
• You'll debug down to the clock behavior, thermal throttling and driver paths where the docs are often missing or just wrong
• You'll turn hard-won knowledge into a machine-searchable structure the compiler can explore on it's own
• You'll work below PTX at the ISA-level, reasoning about SASS and cubins to emit schedules PTX has no way to express
• You'll build the performance models, microbenchmarks and tooling that lets us predict kernel behavior
• You'll work alongside the formal correctness team, so that aggressive, unintuitive kernels ship with a formal proof attached

Expect roughly two-thirds of your time on kernels and one-third on turning what you learned into the compiler search space.

What We're Looking For

• Someone with a track record of hand-writing kernels that match or beat vendor libraries
• Someone comfortable reading PTX, SASS, GCN/CDNA ISA or equivalent machine-level assembly
• Someone fluent with low-level profiling tools: Nsight Compute, Nsight Systems, rocprof, omniperf or their equivalent
• Someone with solid systems programming skills in C++ and CUDA or ROCm/HIP with a working understanding of how high-level ML operations map onto the hardware, including where the framework layers get in the way

Nice to Have

• Someone who's worked on compiler backends (MLIR, LLVM, codegen, instruction selection, scheduling)
• Someone who's done research in the areas of superoptimization, program synthesis, formal verification or search-based compilation
• Someone who's worked with silicon beyond NVIDIA (AMD MI-series, TPU, Trainium) or mobile edge GPUs (Metal, Mali, Aderno)
• Someone who's worked on distributed AI training at depth
• Someone with experience in high-speed interconnects (NVLink, NVSwitch, InfiniBand, RoCE)
• Someone with an HPC background: large-scale scientific computing, MPI, supercomputing
• Someone who's worked in driver development or has a background in EE, computer architecture or hardware design

Why Join Us

This role is for someone who wants to know why things are fast or slow on the actual silicon and is tired of watching that knowledge die undocumented in their own head the moment they move on to the next thing. The compiler takes what you found and runs the same search across every workload we touch, on every vendor's hardware, meaning every optimization compounds across thousands of kernels going forward.

We're a small team operating at frontier scale. We pre-trained foundation models on 4,000 AMD GPUs as a team of three, designed and brought up GB300 NVL72 clusters and designed a TOP500 supercomputer.

We believe that hard problems get solved in person and most of our work happens at our office in San Francisco. We offer relocation assistance and, where possible, we'd like you here as often as possible.

The base salary range for this full-time position is $285,000-$315,000, plus meaningful equity and benefits.