Jobs · Engineering · California

Member of Technical Staff, GPU Kernels

SF Tensor · San Francisco, CA · Yesterday
On-siteEngineering$285k–$315k/yrFull-time

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. 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

  • Write and hand-optimize kernels for real workloads to find the ceiling before the search goes looking for it
  • Profile at the microarchitectural level: SM and CU utilization, warp stalls, memory bank conflicts, register pressure, instruction throughput, occupancy tradeoffs
  • Debug down to the clock behavior, thermal throttling and driver paths where the docs are often missing or just wrong
  • Turn hard-won knowledge into a machine-searchable structure the compiler can explore on its own
  • Work below PTX at the ISA-level, reasoning about SASS and cubins to emit schedules PTX has no way to express
  • Build the performance models, microbenchmarks and tooling that lets us predict kernel behavior
  • 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

  • Track record of hand-writing kernels that match or beat vendor libraries
  • Comfort reading PTX, SASS, GCN/CDNA ISA or equivalent machine-level assembly
  • Fluency with low-level profiling tools: Nsight Compute, Nsight Systems, rocprof, omniperf or their equivalent
  • 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

  • Experience with compiler backends (MLIR, LLVM, codegen, instruction selection, scheduling)
  • Research in superoptimization, program synthesis, formal verification or search-based compilation
  • Work with silicon beyond NVIDIA (AMD MI-series, TPU, Trainium) or mobile edge GPUs (Metal, Mali, Adreno)
  • Experience with distributed AI training at depth
  • Experience with high-speed interconnects (NVLink, NVSwitch, InfiniBand, RoCE)
  • HPC background: large-scale scientific computing, MPI, supercomputing
  • Driver development or 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.

Pay

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

Similar jobs