03 About
Ibteshamul Haque
CUDA kernels and LLM inference systems — and the measurement discipline to tell when something actually got faster.
Currently
Finishing COLIDE — custom CUDA kernels for an IoT intrusion-detection model, now in manuscript prep — and building Slipstream, an LLM serving engine organised around a memory-aware scheduler.
Profile
I work on inference systems close to the metal: writing CUDA kernels by hand when the framework leaves something on the table, building the serving machinery around them, and — increasingly — on the much less glamorous question of whether the resulting numbers mean what they appear to mean.
That last part is not a side interest. Most of what I have learned in the past year came from a measurement being wrong rather than a kernel being slow. A kernel I had explained away as floating-point noise turned out to be a data race. A benchmark whose within-session variance looked tight moved up to 27% between sessions. A published recall figure I was working against turned out to measure grounding rather than discovery, because the prompt handed the model the answer key. In every case the fix was cheap and the wrong version had been sitting there, plausible, for a while.
So the through-line across the work below is: make it fast, then find out whether it actually got faster, then write down what the number does not cover.
What I work on
- Custom CUDA kernels for small, control-flow-heavy models — the shape that production compilers handle worst, where launch overhead dominates arithmetic and a bidirectional LSTM breaks graph capture outright.
- LLM inference systems — KV-cache-constrained scheduling, quantized decode paths, and the block-scaled low-precision GEMMs underneath them.
- On-device and air-gapped deployment, where calling an API is not an available answer and the model has to fit next to everything else already on the box.
- Benchmark methodology — multi-session ranges over hero numbers, both sides of a ratio backed by real distributions, and every published figure tied to an artifact that can re-derive it.
Currently
COLIDE — custom CUDA kernels for a CNN-BiLSTM intrusion-detection model, with an on-device LLM that explains alerts without ever entering the detection path. The detection and kernel gates are closed and the work is in manuscript preparation.
Slipstream — an LLM inference engine built from scratch, where the research sits in the scheduler rather than in any single kernel. Currently on development hardware; the targets are written against A100-class runs that have not happened yet.
qgemm-mx — recovering block-scaled FP4 throughput on Hopper, which has the
memory saving and no native FP4 multiply. Early: the measurement harness and correctness floor exist,
the wgmma kernel does not.
RISC-V parameter extraction — an audit of what a published extraction benchmark actually measures, done through an LFX mentorship with the Parameter SIG, which produced sixteen merged fixes upstream.
How I work
I would rather disclose a limitation than tell a clean story. Every project I run keeps an explicit list of what is not claimed, and in more than one case that list is the most useful document in the repository — it is what stops a convenient sentence getting written six months later by someone who has forgotten the caveat, including me.
Concretely, that means: negative results stay in the table; withdrawn numbers get tombstoned rather than deleted; partial and complete measurements never share a ratio; and where I can, a script fails the build if a retired figure reappears.
Elsewhere
Code is on GitHub. Email is on the contact page and I read it.