Performance Engineer, Kernels

Sarvam; Bengaluru; Hybrid; INR 4,000,000 - 8,000,000; Full time

Job summary

Sarvam is hiring for a Senior Kernel Engineer in Bengaluru/Chennai hybrid roles. You will own the kernel layer, crafting custom CUDA, DSL-based, and PTX kernels that beat baselines on real workloads.

You’ll influence end-to-end performance in the multi-node serving stack and own explanations for the gains. The role requires shipping kernels in production with measurable p99 improvements and close collaboration with the Performance Engineering team.

Qualifications

• 5+ years in ML systems with 2+ years authoring production CUDA kernels.
• CUDA kernel authoring: thread-block sizing, shared-memory layout, warp primitives, async copies, MMA.
• CUTLASS / CuTe DSL at modify-and-extend level; comfortable with layout algebra.
• PTX at debug-and-modify level; hand-written PTX when needed.
• Nsight Compute and systems fluency; read roofline plots and propose fixes.
• Attention kernels experience with FlashAttention-family, paged, MLA, or related work.
• Multi-architecture awareness from Hopper to Blackwell.

Responsibilities

• Own the kernel layer; author custom CUDA, DSL-based, and PTX kernels to close performance gaps.
• Move production p99 latency and throughput; explain performance improvements.
• Collaborate with serving/runtime teams to integrate kernels into a real workload.

Skills

Kernel authoring

ML systems

Performance tuning

Multi-architecture awareness

GPU optimization

Debugging and profiling

Tools

CUDA

CUTLASS / CuTe DSL

PTX

Nsight Compute

NCCL / NVSHMEM

FlashAttention

Triton

Job description

Part of Sarvam's Performance Engineering team. We are hiring two specialized performance roles - Kernels (this posting) and Inference (companion posting). They are a vertical stack: the kernels team authors the µs-level GPU code, and the inference team integrates it into a running serving stack and owns the system-level numbers.

Location

[Bengaluru / Chennai / Hybrid / On-site]

• Team: Performance Engineering
• Level: Senior

About The Team

Sarvam serves multiple model families - small and large LLMs, Mixture-of-Experts, streaming Indic ASR, and multimodal - across a multi-node, multi-tenant fleet on H100 / H200 / B200. The Performance Engineering team owns the numbers the rest of the company plans against: how fast we serve, how much it costs, and how much we get out of every GPU. This team works at the intersection of the serving runtime, the kernel layer, and the SRE org that keeps the fleet alive.

About The Role

You will own the kernel layer. Where stock libraries - cuBLAS, cuDNN, FlashAttention, out-of-the-box Triton - leave performance on the table, you will author the custom CUDA, DSL-based, and PTX kernels that close the gap.

This is a hard, narrow, high-leverage role. We hire engineers who have shipped kernels that beat published baselines on real workloads, not engineers who have used kernels. When your code lands, production p99 moves, and you own the explanation of why.

What We're Looking For

• 5+ years in ML systems, with 2+ years authoring production CUDA kernels. You have a kernel in production that beat the prior baseline by a measurable margin.
• CUDA at kernel-authoring level: thread-block sizing, shared-memory layout, warp primitives, async copies (cp.async, TMA), and MMA selection.
• CUTLASS / CuTe DSL at a modify-and-extend level, with comfort in the layout algebra.
• PTX at a debug-and-modify level - you have inserted hand-written PTX where the compiler missed.
• Nsight Compute and Systems fluency: you read a roofline plot and propose the fix.
• Attention kernels: you have authored or modified at least one (FlashAttention-family, paged, MLA, sliding-window, or sparse).
• Multi-architecture awareness: what changes from Hopper to Blackwell (TMA, WGMMA, tcgen05).

Strong pluses

• Communication kernels - NCCL / NVSHMEM authoring, custom collectives, expert-parallel dispatch (DeepEP-style), AFD bipartite comms (StepMesh-style), KV transfer (DualPath / Mooncake). Strongly desired; dedicated comms-specialist headcount is expected later.
• Open-source kernel contributions - FlashAttention, CUTLASS examples, vLLM / SGLang kernels, DeepEP, Mooncake, or non-trivial Triton work. For this role, the GitHub filter is the highest-yield signal.
• tcgen05, TMA, CTA-cluster launch and distributed shared memory, async pipelining, and the FP4/microscaling paths
• Grace-side host-path optimization on GH200 / GB200.