I build and validate the low-level parts of ML systems: accelerator kernels, compiler lowering, and the runtime work that decides what a model actually costs to run. I've measured on NVIDIA, AMD, Google TPU and AWS Trainium hardware, and I publish the losses next to the wins.
Start here: upstream work · kernel and validation projects · raw benchmark data · Hugging Face · LinkedIn
Every number below names its hardware, its baseline and how it was measured. Simulator, CPU-model and compile-only checks are labeled as such, and they aren't presented as hardware speedups. If a design lost, the table says so.
| Project | What I built and measured | Evidence |
|---|---|---|
| Kernel grader: MI300X and Trainium NKI | A histogram kernel graded on an AMD MI300X (ROCm 7.2.4): privatising bins into LDS took a skewed 67M-key run from 86.39 ms with global atomics to 0.10 ms. The same error bound went into an NKI kernel run in the Neuron simulator and on a trn1.2xlarge NeuronCore, four seeds each. The hardware matched a sequential reference bit for bit, the simulator didn't, and no bin exceeded the bound on either path. My own cost model missed by 5.2x, and that's recorded too. | Card is public; files are access-gated |
| spmv-csr-triton-gpu-kernels | Three Triton CSR SpMV kernels (thread-per-row, vector-per-row, nnz-balanced with atomics) on 1M rows and 32M nonzeros, RTX 4090 and MI300X. Compared against each other, not against cuSPARSE or rocSPARSE. vector-per-row won 3 of 4 runs; the usual load-balanced recommendation won once and was 7.7x slower on the even matrix. | Raw timings and versions in evidence/ |
| tpu-kernel-parity-lab | A Pallas masked-softmax kernel inside grouped-query attention, with causal, padding, all-masked-row and Q/K/V gradient parity tests. On one TPU v5e device, XLA beat Pallas in all 10 paired comparisons (1.23x to 2.38x). That's the result, and it's published as one. | 7/7 correctness checks, 20 timed cases, raw data |
| embodied-efficiency | Latency work on a batch-1 VLA flow-matching sampler. Manual CUDA-graph capture ran 5.9x over eager on a T4 (4.82 to 0.82 ms per step) and beat torch.compile. Weight-only int8/int4 didn't buy latency across four experiments, including torchao/Marlin on an L4; it's a footprint lever. Also a runtime safety supervisor evaluated on real DROID actions with injected faults (drift AUC 0.99). |
Results, CI, demo |
| triton-kernel-lab | Fused softmax, elementwise activations and a tiled matmul in Triton, with correctness tests. Benchmark scripts are included, but I haven't published numbers from them. | Code and tests |
Status as of October 2026. "Open" means under review, not accepted.
| Project | My part | Status |
|---|---|---|
| LLVM: NVPTX aligned small memcpy | Author. Lowers 16-byte-aligned 16 to 127 byte copies to v4 loads and stores instead of chained b64 pairs. New lit test; 629 NVPTX lit tests pass. On a T4, an L2-resident 1 MiB copy went from 93.9 to 98.3 GB/s; DRAM-bound sizes were unchanged. Bit-exact on an L4. | Open (intel/llvm mirror) |
| Liger-Kernel: fused linear KL | Author. Runs launches under device_context and pins kernel scalars to fp32, which fixes a cross-device launch error and a torch.compile build failure. Tested on 2x T4; 2 new tests. |
Open |
| vLLM: gfx950 top-k split policy | Author. Device-length-aware work partitioning for top-k with k=2048. Exact match with a CPU model on 35 shapes and compiles for gfx950. GPU timings aren't in yet, so there's no speedup claim. | Open |
| QuACK: PDL waits in grouped GEMM | Wrote the SM100 A-index prefetch and SM90 math warpgroup commits in mcwitt's PR, checked by SASS inspection. Not yet tested for GPU correctness or performance. | Open, pending GPU validation |
| TorchRec: grad_indice_weights | Investigation and a proposed 23-line fix on an issue someone else opened. On a MIG slice of an RTX Pro 6000 Blackwell, the kernel went from 3.86 to 0.054 ms per step and the full step from 19.74 to 15.93 ms, with bit-exact weights and outputs. Local result, not merged. | Proposed in comment |
| FlashAttention: varlen LSE fix | Independent validation on an L4 of another contributor's fix. Reproduced the failure on main and confirmed the fix. I didn't write the fix. | Validation |
| Network UPS Tools: parseconf warning | Author. | Merged |
Other open PRs, smaller in scope: XLA, wgpu/naga, SPIRV-Cross, fio, ComfyUI, speech-to-speech, chai-lab, pygpubench.
- kernel-grader-mi300x-rocm-trainium-nki: MI300X and Trainium runs, simulator vs device, per seed
- tpu-kernel-parity-lab-v5e: every TPU timing sample plus correctness records
- spmv
evidence/: RTX 4090 and MI300X raw timings - embodied-efficiency
RESULTS.md: T4 and L4 runs, including the negative low-bit results
These are my own measurements on rented or free-tier hardware. They're reproducible from the scripts, but nobody has independently reproduced them yet.
Kernel correctness and performance across vendors, and graders that can't be gamed by a kernel that skips the work. I'm building this into ComputeLab, where I'm co-founder and CTO.
Triton · CUDA and PTX/SASS reading · HIP/ROCm · AWS NKI · JAX/Pallas · LLVM (NVPTX) · PyTorch · Python, C++
My kernel and benchmark repos (spmv, tpu-kernel-parity-lab, embodied-efficiency, the kernel grader) are source available under the CodeZorana Source Available License, not open source. Reading, study and reproduction are fine; other use needs a licence. triton-kernel-lab and older projects are MIT.
LinkedIn · Hugging Face · Kaggle · or open an issue on any repo.