Hacker News Board Hacker News, browsed like a classic message board
Live data from Hacker News

Compiling LLMs into a MegaKernel: A path to low-latency inference

zhihaojia.medium.com

Post reply on HN
Print view View original

61–70 of 79 posts

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#61

Post by touisteur »

skavi wrote: post #46
refulgentis wrote: post #16

Earlier quoted context omitted.

You've hit the nail on the head. The CPU launch cost of a pre-compiled CUDA graph is tiny. CUDA Graphs are a huge step up from manually launching kernels, but they still treat kernels as monolithic, black-box operations. A megakernel erases the boundaries between those operations. With CUDA Graphs, as in the example in the article, if you have Matmul -> AllReduce, the AllReduce kernel cannot start until the entire Ma…

Ah, that makes a lot of sense. Is this fine grained task scheduling related to CUDA Dynamic Parallelism at all? If not would you have a pointer on where to look? I suppose I could look through the code of this project, but I’d hate to have detangle that from the compiler infrastructure.

Think more of it as 'tasking by hand' where you have one kernel driving the 18k+ cores and you manually (or using device libraries) fine-grained synchronize them, handle memory traffic asynchoronously and pipeline as much as you can.

You might have a look at cooperative groups, also things cuda::pipeline in libcudacxx to handle asynchronous and pipelined memory traffic, and also most of block/warp CUB primitives, and move on up to cuFFTDx, cuBLASDx and now cuSolverDx as the starting toolbox for your fused kernel journey.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#62

Post by nialse »

baq wrote: post #3

Next step - compile straight to verilog so I can buy some LLMs on aliexpress

In five to ten years, when LLMs have stabilized, mapping them straight onto hardware will probably make sense. With today’s processes a hundred billion parameters might fit onto a single silicon wafer using ~1.5 bit precision implemented directly in logic gates. Using higher precision raises the gate count exponentially, so it makes more sense to keep the weights in memory and reuse shared compute blocks for the math for now. We need to get the ultra low precision LLMs working for the future though.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#63

Post by rhdunn »

fc417fc802 wrote: post #52
stirfish wrote: post #42

Earlier quoted context omitted.

I was thinking more like those Bitcoin mining usb Asics that used to be a thing, but instead of becoming ewaste, you can still use them to talk with chatgpt 2 or whatever. I'm picturing an llm appliance.

There is no magic ASIC that can get around needing to do hundreds of watts worth of computations and having on the order of hundreds of gigabytes of very fast memory. Otherwise the major players would be doing that instead of (quite literally) investing in nuclear reactors to power their future data center expansions.

Google have their own ASIC via their TPU. The other major players have leveraged NVIDIA and -- to a lesser extent -- AMD. This is partly due to investment in TPUs/ASICs being complex (need specialist knowledge and fabrication units) and GPU performance being hard to compete with.

Training is the thing that costs the most in terms of power/memory/energy, often requiring months of running multiple (likely 4-8) A100/H100 GPUs on the training data.

Performing inference is cheaper as you can 1) keep the model loaded in VRAM, and 2) run it on a single H100. With the 80GB capacity you would need two to run a 70B model at F16, or one at F8. For 32B models and lower you could run them on a single H100. Then you only need 1 or 2 GPUs to handle the request.

ASICs could optimize things like the ReLU operations, but modern GPUs already have logic and instructions for matrix multiplication and other operations.

I think the sweat spot will be when CPUs have support for high-throughput matrix operations similar to the SIMD operations. That way the system will benefit from being able to use system memory [1] and not have another chip/board consuming power. -- IIUC, things are already moving in that direction for consumer devices.

[1] This will allow access to large amounts of memory without having to chain multiple GPUs. That will make it possible to run the larger models at higher precisions more efficiently and process the large amount of training data efficiently.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#64

Post by gdiamos »

liuliu wrote: post #10
scotty79 wrote: post #4

> Traditional LLM systems often rely on sequences of GPU kernel launches and external communication calls, resulting in underutilized hardware. What? Why? This seems like an obvious optimization if it's possible.

It really is not obvious. These launches are asynchronous, and data movement / computation is overlapped properly through CUDA APIs. Even per-kernel launch cost is reduced with the cudagraph introduction. CUDA programming model relies on each kernel to be computationally expensive to make sense, and these are not true for token generation of LLM. And we are talking about network evaluation at higher than 1000 per sec…

It's surprising to me that the field is willing to invest this much in mega-kernels, but not models that generate multiple tokens in parallel...

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#65

Post by baq »

fc417fc802 wrote: post #52
stirfish wrote: post #42

Earlier quoted context omitted.

I was thinking more like those Bitcoin mining usb Asics that used to be a thing, but instead of becoming ewaste, you can still use them to talk with chatgpt 2 or whatever. I'm picturing an llm appliance.

There is no magic ASIC that can get around needing to do hundreds of watts worth of computations and having on the order of hundreds of gigabytes of very fast memory. Otherwise the major players would be doing that instead of (quite literally) investing in nuclear reactors to power their future data center expansions.

If you focus on just the matmuls, no CUDA, no architectures, no infinibands, everything-on-a-chip - put input tokens in input registers, get output tokens from output registers from a model that's baked into gates - you should be able to save some power. Not sure if 10x or 2x or 100x, but certainly there are gains to be had.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#66

Post by rhdunn »

otabdeveloper4 wrote: post #60
fc417fc802 wrote: post #52

Earlier quoted context omitted.

There is no magic ASIC that can get around needing to do hundreds of watts worth of computations and having on the order of hundreds of gigabytes of very fast memory. Otherwise the major players would be doing that instead of (quite literally) investing in nuclear reactors to power their future data center expansions.

24 gigabytes is more than enough to run a local LLM for a small household or business. This is "gaming PC" territory, not "space heater". I mean people already have PS5's and whatnot in their homes. The hundreds of gigabytes thing exists because the big cloud LLM providers went down the increasing parameter count path. That way is a dead end and we've reached negative returns already. Prompt engineering + finetunes i…

It depends on 1) what model you are running; and 2) how many models you are running.

You can just about run a 32B (at Q4/Q5 quantization) on 24GB. Running anything higher (such as the increasingly common 70B models, or higher if you want to run something like Llama 4 or DeepSeek) means splitting the model between RAM and RAM. -- But yes, anything 24B or lower you can run comfortably, including enough capacity for the context.

If you have other models -- such as text-to-speech, speech recognition, etc. -- then those are going to take up VRAM for both the model and during processing/generation. That affects the size of LLM you can run.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#67

Post by fho »

Somewhat relevant anecdote: we had a small CUDA competition (10-ish years ago). Some embrassingly parallel CV algorithm.

I tried to be smart and cache intermediate results that were shared by multiple kernels.

When the results were published I was stumped to see that others were orders of magnitude faster then me.

Turns out they didn't bother with caching at all. The overhead of recalculating everything a thousand times was tiny compared to the overhead of doing roundtrips through RAM.

I assume it's the same thing here. By compiling into MegaKernels, layer boundaries are squashed. There likely will be _more_ calculations and less shared intermediate results. But overall it's still a win due to less memory roundtrips.

There has to be a sweet spot, especially for convolution networks. No idea if the MegaKernel takes this into account.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#68

Post by andy12_ »

How is this possible? I mean, I thought that sometimes you had no choice but to separate computation into several kernels. But here they literally allow cuda threads to dinamically perform tasks assigned by scheduler threads? I only have a little experience writing cuda kernels, so I have my mind blown.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#69

Post by fc417fc802 »

rhdunn wrote: post #63
fc417fc802 wrote: post #52

Earlier quoted context omitted.

There is no magic ASIC that can get around needing to do hundreds of watts worth of computations and having on the order of hundreds of gigabytes of very fast memory. Otherwise the major players would be doing that instead of (quite literally) investing in nuclear reactors to power their future data center expansions.

Google have their own ASIC via their TPU. The other major players have leveraged NVIDIA and -- to a lesser extent -- AMD. This is partly due to investment in TPUs/ASICs being complex (need specialist knowledge and fabrication units) and GPU performance being hard to compete with. Training is the thing that costs the most in terms of power/memory/energy, often requiring months of running multiple (likely 4-8) A100/H10…

> ASICs could optimize things like the ReLU operations, but modern GPUs already have logic and instructions for matrix multiplication and other operations.

Right but at that point you're describing an H100 plus an additional ASIC plus presumably a CPU and some RAM. Or a variant of an H100 with some specialized ML functions baked in. Both of those just sound like a regular workstation to me.

Inference is certainly cheaper but getting it running quickly requires raw horsepower (thus wattage, thus heat dissipation).

Regarding CPUs there's a severe memory bandwidth issue. I haven't kept track of the extreme high end hardware but it's difficult to compete with GPUs on raw throughput.

Re: Compiling LLMs into a MegaKernel: A path to low-latency inference

#70

Post by fc417fc802 »

otabdeveloper4 wrote: post #60
fc417fc802 wrote: post #52

Earlier quoted context omitted.

There is no magic ASIC that can get around needing to do hundreds of watts worth of computations and having on the order of hundreds of gigabytes of very fast memory. Otherwise the major players would be doing that instead of (quite literally) investing in nuclear reactors to power their future data center expansions.

24 gigabytes is more than enough to run a local LLM for a small household or business. This is "gaming PC" territory, not "space heater". I mean people already have PS5's and whatnot in their homes. The hundreds of gigabytes thing exists because the big cloud LLM providers went down the increasing parameter count path. That way is a dead end and we've reached negative returns already. Prompt engineering + finetunes i…

Only if you'll settle for less than state of the art. The best models still tend to be some of the largest ones.

Anything that overflows VRAM is going to slow down the response time drastically.

"Space heater" is determined by computational horsepower rather than available RAM.

How big a context window do you want? Last I checked that was very expensive in terms of RAM and having a large one was highly desirable.

Post reply on HN