--- title: How To Build Your Own LLM Runtime From Scratch | daily.dev description: A deep technical walkthrough of building a custom CUDA LLM inference runtime from scratch for Qwen2.5-Coder-7B on an NVIDIA H100. The author covers the full... canonical: https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok twitter:card: summary_large_image twitter:site: @dailydotdev og:type: website og:site_name: daily.dev og:title: How To Build Your Own LLM Runtime From Scratch | daily.dev og:description: A deep technical walkthrough of building a custom CUDA LLM inference runtime from scratch for Qwen2.5-Coder-7B on an NVIDIA H100. The author covers the full... og:url: https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok og:image: https://api.daily.dev/og/posts/MEzU7RJOk.png og:image:alt: How To Build Your Own LLM Runtime From Scratch og:image:width: 1200 og:image:height: 630 og:locale: en --- > ## Documentation Index > Fetch the complete documentation index at: https://daily.dev/llms.txt > Use this file to discover all available pages before exploring further. # How To Build Your Own LLM Runtime From Scratch **[Towards Data Science](https://daily.dev/sources/tds)** · 21 min read · 14 upvotes · 1 comments ## Summary A deep technical walkthrough of building a custom CUDA LLM inference runtime from scratch for Qwen2.5-Coder-7B on an NVIDIA H100. The author covers the full decode stack: custom INT4 weight packing (.nanoqwen format), paged KV cache, warp-specialized attention kernels, and CUDA graph capture. Three key bugs are documented in detail: (1) placing __syncthreads() inside a warp-id branch causing silent softmax corruption at KV page tails, (2) eager kernel launches causing 119 ms/token latency that CUDA graph capture reduced to ~17 ms/token (7x speedup), and (3) INT8 activation quantization losing to FP16+prmt.b32 INT4 unpack on Hopper for GEMV shapes. The runtime achieves ~60 tok/s decode vs llama.cpp's ~200 tok/s, with the gap framed as a learning exercise rather than a production competitor. ## Full article daily.dev links to this article rather than hosting it. Read it at the original source: ## Questions this post answers ### Why did wrapping CUDA decode in a captured CUDA graph make such a big performance difference on H100? Capturing decode into a CUDA graph eliminated per-kernel launch overhead, dropping steady-state latency from about 119 ms/token with eager launches to about 17 ms/token, a 7x speedup with identical kernels. Eager decode required roughly 280 separate cudaLaunchKernel calls per token across 28 layers, each forcing a costly CPU-driver round trip that left the GPU idle 90% of the time. _For anyone chasing decode latency, daily.dev surfaces the CUDA graph tricks that make or break inference throughput._ ### Why does putting __syncthreads() inside a warp-id branch cause silent corruption instead of a crash? __syncthreads() is a block-wide barrier, so if it sits inside a branch like if (warp_id == 0), only that warp reaches it while the other warps skip past and read shared memory before data has fully landed. In a warp-specialized paged-attention kernel this let consumer warps read stale SMEM at partial KV tail pages, producing confident but wrong softmax outputs rather than an obvious crash. _Debugging warp-specialization races is easier when daily.dev keeps CUDA synchronization gotchas within reach._ ### Why did FP16-activation INT4 unpacking with prmt.b32 outperform INT8-activation dp4a on Hopper GEMV shapes? For gate_proj, up_proj, and down_proj GEMV shapes on Hopper, the prmt.b32-based INT4 unpack with FP16 activations beat the __dp4a INT8-activation path, even though INT8 activations use half the memory bandwidth of FP16. The extra kernel needed to quantize FP16 activations to INT8 before each GEMV became the new bottleneck, so the INT8 path was implemented, benchmarked, and then removed. _Choosing quantization tactics for GPU kernels gets easier when daily.dev tracks real hardware-specific benchmark tradeoffs like this._ ## Community discussion Top comments from developers on daily.dev. **@akashskypatel** · 0 upvotes > What platform do people use to rent GPU compute that doesn't break the bank/is free for learning? ## Similar posts on daily.dev - [The Infrastructure Behind Making Local LLM Agents Actually Useful](https://daily.dev/posts/the-infrastructure-behind-making-local-llm-agents-actually-useful-ff4qnxir1) · Towards Data Science · 1 upvotes · 0 comments - [How I massively improved my AI inference performance without buying new hardware](https://daily.dev/posts/how-i-massively-improved-my-ai-inference-performance-without-buying-new-hardware-jgyj1odfm) · Red Hat Developer · 0 upvotes · 0 comments - [The LLM Inference Optimization Stack: From Quantization to Speculative Decoding Part 1](https://daily.dev/posts/the-llm-inference-optimization-stack-from-quantization-to-speculative-decoding-part-1-ryexzxqsr) · DigitalOcean Community · 0 upvotes · 0 comments - [Your old GPU can still run big LLMs – you just need the right tweaks](https://daily.dev/posts/your-old-gpu-can-still-run-big-llms-you-just-need-the-right-tweaks-3sa1xwzxu) · XDA Developers · 0 upvotes · 0 comments - [5 Techniques to Optimize LLMs in Production](https://daily.dev/posts/5-techniques-to-optimize-llms-in-production-lmz4pypvm) · Daily Dose of Data Science \| Avi Chawla \| Substack · 0 upvotes · 0 comments --- Tags: [#data-science](https://daily.dev/tags/data-science), [#gpu](https://daily.dev/tags/gpu), [#ai-inference](https://daily.dev/tags/ai-inference), [#cuda](https://daily.dev/tags/cuda) [View this post on daily.dev](https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok) ```json {"@context":"https://schema.org","@graph":[{"@type":"Organization","@id":"https://daily.dev/#organization","name":"daily.dev","url":"https://daily.dev","logo":{"@type":"ImageObject","url":"https://daily.dev/apple-touch-icon.png","width":180,"height":180},"sameAs":["https://twitter.com/dailydotdev","https://github.com/dailydotdev","https://www.linkedin.com/company/daily-dev-ltd"]},{"@type":"WebSite","@id":"https://daily.dev/#website","url":"https://daily.dev","name":"daily.dev","publisher":{"@id":"https://daily.dev/#organization"},"potentialAction":{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https://daily.dev/search?q={search_term_string}"},"query-input":"required name=search_term_string"}}]} {"@context":"https://schema.org","@type":"TechArticle","headline":"How To Build Your Own LLM Runtime From Scratch","url":"https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok","mainEntityOfPage":{"@type":"WebPage","@id":"https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok"},"datePublished":"2026-07-22T15:36:53.125Z","dateModified":"2026-09-14T08:41:53.611Z","description":"A deep technical walkthrough of building a custom CUDA LLM inference runtime from scratch for Qwen2.5-Coder-7B on an NVIDIA H100. The author covers the full...","image":"https://media.daily.dev/image/upload/f_auto,q_auto/v1/posts/a5d9c9e4758c09baedd91589fa180a3c?_a=AQAEuop","thumbnailUrl":"https://media.daily.dev/image/upload/f_auto,q_auto/v1/posts/a5d9c9e4758c09baedd91589fa180a3c?_a=AQAEuop","isAccessibleForFree":true,"articleSection":"Towards Data Science","inLanguage":"en","publisher":{"@type":"Organization","name":"daily.dev","url":"https://daily.dev","logo":{"@type":"ImageObject","url":"https://daily.dev/apple-touch-icon.png","width":180,"height":180}},"author":{"@type":"Organization","name":"Towards Data Science","logo":"https://media.daily.dev/image/upload/t_logo,f_auto/v1/logos/tds","url":"https://daily.dev/sources/tds"},"commentCount":1,"discussionUrl":"https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok","interactionStatistic":[{"@type":"InteractionCounter","interactionType":{"@type":"LikeAction"},"userInteractionCount":14},{"@type":"InteractionCounter","interactionType":{"@type":"CommentAction"},"userInteractionCount":1}],"keywords":"data-science,gpu,ai-inference,cuda","timeRequired":"PT21M"} {"@context":"https://schema.org","@type":"BreadcrumbList","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https://daily.dev"},{"@type":"ListItem","position":2,"name":"Towards Data Science","item":"https://daily.dev/sources/tds"},{"@type":"ListItem","position":3,"name":"How To Build Your Own LLM Runtime From Scratch"}]} {"@context":"https://schema.org","@type":"WebPage","@id":"https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok","comment":[{"@type":"Comment","text":"What platform do people use to rent GPU compute that doesn’t break the bank/is free for learning?","datePublished":"2026-07-22T17:23:49.728Z","url":"https://daily.dev/posts/MEzU7RJOk#c-3zJuRKKdh","author":{"@type":"Person","name":"Akash","url":"https://daily.dev/akashskypatel","image":"https://avatars.githubusercontent.com/u/8129618?v=4"}}]} {"@context":"https://schema.org","@type":"FAQPage","@id":"https://daily.dev/posts/how-to-build-your-own-llm-runtime-from-scratch-mezu7rjok#faq","mainEntity":[{"@type":"Question","name":"Why did wrapping CUDA decode in a captured CUDA graph make such a big performance difference on H100?","acceptedAnswer":{"@type":"Answer","text":"Capturing decode into a CUDA graph eliminated per-kernel launch overhead, dropping steady-state latency from about 119 ms/token with eager launches to about 17 ms/token, a 7x speedup with identical kernels. Eager decode required roughly 280 separate cudaLaunchKernel calls per token across 28 layers, each forcing a costly CPU-driver round trip that left the GPU idle 90% of the time. For anyone chasing decode latency, daily.dev surfaces the CUDA graph tricks that make or break inference throughput."}},{"@type":"Question","name":"Why does putting __syncthreads() inside a warp-id branch cause silent corruption instead of a crash?","acceptedAnswer":{"@type":"Answer","text":"__syncthreads() is a block-wide barrier, so if it sits inside a branch like if (warp_id == 0), only that warp reaches it while the other warps skip past and read shared memory before data has fully landed. In a warp-specialized paged-attention kernel this let consumer warps read stale SMEM at partial KV tail pages, producing confident but wrong softmax outputs rather than an obvious crash. Debugging warp-specialization races is easier when daily.dev keeps CUDA synchronization gotchas within reach."}},{"@type":"Question","name":"Why did FP16-activation INT4 unpacking with prmt.b32 outperform INT8-activation dp4a on Hopper GEMV shapes?","acceptedAnswer":{"@type":"Answer","text":"For gate_proj, up_proj, and down_proj GEMV shapes on Hopper, the prmt.b32-based INT4 unpack with FP16 activations beat the __dp4a INT8-activation path, even though INT8 activations use half the memory bandwidth of FP16. The extra kernel needed to quantize FP16 activations to INT8 before each GEMV became the new bottleneck, so the INT8 path was implemented, benchmarked, and then removed. Choosing quantization tactics for GPU kernels gets easier when daily.dev tracks real hardware-specific benchmark tradeoffs like this."}}]} ```