skip ↵
open to LLM inference & GPU roles · remote / India

Aman
Karki

I make LLM inference faster and I check the numbers. Two CUDA PRs merged into llama.cpp, and right now I'm working on block-FP8 for Ada GPUs in vLLM. I've also built an inference engine, a vector database and a code analyzer from scratch, each checked against a trusted reference and benchmarked, losses included.

Download résumé See the work
request.trace
tokenizer prefillcompute-bound GEMMs decodebandwidth-bound KV cache FP8 GEMM tiles · 128×128 blocks
>
illustrative, not a benchmarkphase: prefill
0
CUDA PRs merged into llama.cpp
0
backend tests passing on my DUP fix
0tok/s
decode on a T4, my own C++/CUDA engine
0×
over CPU on the same machine
0µs
p50 at 95.4% recall, my vector DB
0×
FP8 path over Marlin on an L4 (measured)
// 01 · NOW

Block-FP8 on Ada GPUs, in vLLM

What I'm working on right now, in public.

● in progress

Ada GPUs (L4, L40S, RTX 4090) leave their FP8 tensor cores unused for block-FP8 models.

In vLLM, block-FP8 checkpoints (DeepSeek-style 128×128 scales) fall back to a weight-only Marlin path on SM89, because the CUTLASS blockwise kernels need Hopper. I measured the cost on an NVIDIA L4, posted the numbers on vllm#58241, and fixed vLLM's broken block-FP8 benchmark on the way (vllm#59261).

Now: an SM89 blockwise GEMM in CUTLASS 2.x (Ada's FP8 tensor cores, per-token activation scales), open for review in vllm#61224. Profiling found register spills in the mainloop; scaling once per 128-wide K block removed them. End to end on an L4 (Qwen3-8B-FP8): ~1.4× throughput over Marlin with the same gsm8k accuracy; batch-1 decode is still ~4% slower, and that's what I tune next. Test coverage for SM89 is up for review in vllm#60171, and an int8 small-M fix for a user-reported slowdown in vllm#62128.

prefill / batchesFP8 tensor cores win: 1.4–2.8× over Marlin at M ≥ 256
decode (M ≤ 64)Marlin is still a little ahead at batch 1 (~4% end to end). That's the hard part.
speedup vs today's Marlin pathL4 · M=1024
Marlin (today)
1.00×
gate_up_proj
2.81×
down_proj
1.91×
o_proj
1.71×
qkv_proj
1.48×

FP8 tensor-core path (CUTLASS per-tensor, the upper bound for a blockwise kernel) vs Marlin. Qwen3-8B layer shapes, activation quant included. Measured on an NVIDIA L4, vLLM main @ 741edeebe.

// 02 · OPEN SOURCE

Code in the engines people actually run

Small, well-tested changes, with the test command and the result in every PR.

vllm-project/vllm #61224open

CUTLASS blockwise FP8 GEMM for Ada (SM89) with per-token scales

Block-FP8 models on Ada ran through Marlin, a weight-only fallback. This kernel uses the FP8 tensor cores: spill-free mainloop, tiles picked by M. End to end on an L4 (Qwen3-8B-FP8): ~1.4× throughput, same accuracy; batch 1 still ~4% slower.

Oct 2026 · 6 files view PR →
vllm-project/vllm #62128open

int8 W8A8 GEMM on Ada: fix the small-M slowdown

A user reported int8 GEMMs getting ~1.5× slower at M=17 on an RTX 4090. The dispatch switched to bigger tiles and left most SMs idle. 16-row tiles for small N: ~13% less GEMM time over M 17–64 on the 4090, identical results, 400 tests passing. A reviewer's run on a smaller 60-SM Ada card showed slowdowns, so I reworked the rule to count GPU waves first: now faster on all three Ada GPUs tested (4090, L4, 4500 Ada) with no slower cells.

Oct 2026 · 1 file view PR →
vllm-project/vllm #60171open

Run the block-FP8 kernel tests on Ada (SM89)

The whole test file skipped below SM90, so the L4 CI job never ran it. Lowered the skip, gated the Hopper-only tests, and fixed a 1-ULP rounding tie. On an L4: 221 passed, 0 failed.

Oct 2026 · 1 file view PR →
vllm-project/vllm #59261open

Fix the block-FP8 GEMM benchmark for the current kernel API

It crashed on main after an API change, and its "triton" and "cutlass" lines never forced the kernel in their name. Now each provider asserts it runs the kernel it names. Run on an L4 and an H100 across 11 DeepSeek-V3 shapes.

Sep 2026 · 1 file view PR →
vllm-project/vllm #58241research

Block-FP8 on SM89: L4 baseline and a kernel plan

Measured Marlin against FP8 tensor-core paths on four Qwen3-8B layers from M=1 to 4096, then proposed a CUTLASS 2.x blockwise kernel that must not regress decode.

Sep 2026 · 40 measurementsview issue →
ggml-org/llama.cpp #28897merged

CUDA: i16/i32 support for GGML_OP_DUP

Both types were quietly running on the CPU: the i32 path existed but a capability check blocked it, and i16 had none. Full backend suite on two T4s: 16,097/16,097. Approved by Georgi Gerganov, merged by am17an.

Sep 2026 · 0 regressionsview PR →
ggml-org/llama.cpp #27573merged

CUDA: 1D pooling (POOL_1D)

A new CUDA kernel for average and max pooling that matches the CPU reference exactly at padded edges. Tested on all 216 kernel/stride/padding combinations. Merged by Georgi Gerganov.

Aug 2026 · 216/216 casesview PR →
// 03 · BUILT FROM SCRATCH

Systems I built end to end

No framework doing the real work. Each one is checked against a reference and benchmarked, and the losses are published too.

FLAGSHIP · LLM INFERENCE ENGINE · C++20 / CUDA

verbum.cpp

Runs Qwen3-0.6B end to end with no PyTorch and no existing runtime: safetensors loader, BPE tokenizer, GQA attention, RoPE, SwiGLU, KV cache and sampling. Diffing every layer against HuggingFace caught five bugs that never crashed, they just gave slightly wrong answers.

3e-5max logit diff vs HuggingFace
28 tok/sdecode on a T4 (38× CPU)
~25%of the bandwidth roofline, room left
3.99×smaller weights, per-row INT8
$ ./test_logits models/qwen3-0.6b position=0 max_diff=8.1e-06 OK position=5 max_diff=4.15 FAIL -> RoPE convention mismatch, only visible past position 0 switch to split-half rotation position=5 max_diff=3.0e-05 OK > all 3 prompts match HuggingFace
__global__ void matmul_tiled( const float* A, const float* B, float* C) { // shared-memory tiling, coalesced loads __shared__ float tile[TILE][TILE]; // one of: matmul, rmsnorm, rope, swiglu, // attention decode }
$ ./generate -m qwen3-0.6b --cuda --greedy \ -p "The capital of France is" loaded 3006.53 MB, tied_embeddings=true The capital of France is Paris. The capital of Italy is Rome... 28.01 tok/s on NVIDIA T4
EMBEDDED VECTOR DATABASE · C++20 / PYTHON

Lattice

Closer to SQLite than to a service: hand-written HNSW, a WAL storage engine with crash recovery, mmap segments and scalar quantization. Benchmarked against Qdrant and Chroma, including where it loses.

583 µsp50 at 95.4% recall, SIFT10K
1.3 msp50 on SIFT1M
9.4×faster than brute force at 50K
PyPIpip install pylattice-db
$ lattice bench --dataset sift 10,000 vectors, 128 dims p50 latency 583 µs recall@10 95.4% SIFT1M p50 1.3 ms > HNSW 9.4× faster than brute force at 50K published: pylattice-db on PyPI
class HNSWIndex { public: void insert(const Vector& v) { // probabilistic level assignment int level = assign_level(); greedy_search_and_link(v, level); } };
ARCHITECTURE ANALYTICS · C++20 / PYTHON · VS CODE

RAAG

A Tree-sitter parser on a std::jthread pool reads a repository in parallel, and a Python engine builds the dependency graph and computes coupling, instability and cohesion. GraphRAG keeps AI refactoring inside the code a change can reach, and a CI gate blocks PRs that push a module past its instability limit.

3.69×speedup on 8 cores
1,290files/s, zero parse failures
913dependency edges analyzed
CI gate+ VS Code extension
$ gh pr checks 42 61 files, 126 dependencies, 0 cycles ✗ RAAG Architectural Analysis 4 error(s): instability past threshold (0.40) pipeline.py I=0.73 this blocks the merge
class ThreadPool { public: explicit ThreadPool(size_t n) { // cooperative cancellation via jthread for (size_t i = 0; i < n; ++i) workers_.emplace_back( [this](std::stop_token st) { run(st); }); } };
// 04 · WRITING

Notes from the work

What I built, what broke, and what I'd change. On Hashnode and Medium; four are published in Level Up Coding.

vLLM · GPU waves · Hashnode · Level Up Coding

My GPU fix was 13% faster on one card and 81% slower on another

A reviewer with a smaller Ada GPU, the one idea that explained it (waves), and the rule that fixed it on three Ada GPUs.

Hashnode → · Level Up Coding →
vLLM · part 2 · Level Up Coding

Blockwise FP8 on Ada GPUs: my kernel was correct and 4x too slow

Register spills in the main loop, the scale-once fix, and ~1.42× end to end on an L4.

read →
vLLM · Level Up Coding

Blockwise FP8 on Ada GPUs: why your L4 falls back to Marlin, measured

What 128×128 block scaling is, why it needs its own kernel, and what the fallback costs on a real L4.

read →
llama.cpp · Level Up Coding

Two CUDA PRs into llama.cpp, and what they taught me about a 100K-star codebase

A pooling kernel, a silent CPU fallback, and how ggml decides where ops run.

read →
verbum.cpp

What actually happens inside a transformer forward pass

Five bugs that never crash. They just quietly give wrong answers.

read →
verbum.cpp

INT8 quantization, the second time around

Catching myself about to publish the flashy number instead of the honest one.

read →
verbum.cpp × Lattice

Two from-scratch systems, and the day they talked

The demo where the inference engine and the vector database finally meet.

read →
Lattice

Benchmarking against Qdrant and Chroma

Real recall and latency numbers, including where it loses.

read →
Lattice

Building an HNSW index from scratch

The core algorithm behind an embedded vector database.

read →
Lattice

What I learned building a storage engine from scratch

The design decisions that held up, and the ones I'd change.

read →
RAAG

Building a parallel C++ source parser

jthread, stop_token, and the deadlock I didn't see coming.

read →
RAAG

An AI refactoring tool that can't see more code than the dependency graph allows

Scoping AI refactoring to a change's blast radius.

read →
// 05 · ABOUT

How I got here

Computer science, a year of music, then all-in on inference.

Dec 2025 – now

Software Engineer (Independent)

Full-time on open source and my own systems work in C++20, CUDA and Python. Contributing to llama.cpp since Aug 2026 and to vLLM since Sep 2026. Everything is checked against a reference, benchmarked and written up.

Dec 2024 – Nov 2025

Career break: music

A year of songwriting, composition and vocals. Picked through a three-stage nationwide audition for a label release (Rangdaari, 128K+ views), and wrote, sang and edited my debut original, Ye Meri Thi Zindagi. Listen →

2019 – 2024

B.Tech, Computer Science and Engineering

University of Petroleum and Energy Studies, Cyber Security and Forensics specialization.

Live from GitHub

public repos…
contributions (last year)…
followers…
latest activity…
last 14 days
fetched live from api.github.com
// 06 · CONTACT

Working on inference, GPUs or C++ systems? Let's talk.

I'm looking for roles in LLM inference, GPU/CUDA performance and C++ systems, remote or at an India office. Try the terminal, or just email me.

aman@inference:~
connected.
type help to see commands.
aman@inference:~$
OBJECTIVE
routing…
TOKENS
0
LEVEL
1
BEST
0
STATION REACHED
arrows / WASD / swipe · avoid red OOM · P pause · C camera · ESC quitskip → résumé

ROUTE THE REQUEST

You're a token stream on a GPU die. Visit every station, collect the scale keys, and reach the résumé portal.

  1. Follow the compass to the station glowing green. Each one tells part of my story
  2. Red towers are OOM. Hit one, a wall or yourself and the request dies
  3. Every station makes you faster. Press C to switch camera
  4. Finale: grab 3 FP8 scale keys, then enter the portal to download the résumé
3

OUT OF MEMORY

The request hit an OOM block.

just the résumé

REQUEST SERVED

Every section visited. Streaming résumé.pdf…