Helion: A high-level DSL for performant and portable ML kernels (opens in new tab)

I think unlike Gluon/CuTe/ThunderKittens (which distinguish themselves from Triton by being lower level giving you more control, thus being less performance portable and harder to write), Helion distinguishes itself from Triton by being higher level and easier to write.

IMO, this is something that makes sense for PyTorch to release, as "neutral ground" in the industry.

Generally you are unlikely to get Python-level debugging for code that is going to run on GPUs.

> how much of ML development these days touches these “lower level” parts of the stack? I’d expect that by now most of the work would be high level

Every time the high level architectures of models change, there are new lower level optimizations to be done. Even recent releases like GPT-OSS adds new areas for improvements, like MXFP4, that requires the lower level parts to created and optimized.

There are some not so niche communities, like FlashAttention and LinearFlashAttention repos. New code/optimizations get committed on a weekly basis. They find a couple of percents here and there all the time. How useful their kernels actually are in term of producing good results remain to be seen, but their implementations are often much better (in FLOPS) compared to what were proposed in the original papers.

It's just like game optimization, cache-friendliness and memory hierarchy-awareness are huge in attention mechanism. But programming backward pass in these lower-level stacks is definitely not fun, tensor calculus breaks my brain.

a recent wave of interest in bitwise equivalent execution had a lot of kernels this level get pumped out.

new attention mechanisms also often need new kernels to run at any reasonable rate

theres definitely a breed of frontend-only ML dev that dominates the space, but a lot novel exploration needs new kernels

> CUDA that a couple of years ago, and now CUDA is all but irrelevant

What? CUDA won't be irrelevant for years even if all the competitors figure out the holy grail, the ecosystem doesn't suddenly migrate over night. People learning CUDA today will continue to be find jobs and opportunities across the sector for the near future without any worries.

> but at the same time, the ML architecture space has collapsed entirely and everyone is just using transformers.

That's also not true, the ML space is still growing, and lots of things outside of Transformers, but it requires you to actually look and pay attention, not just browse the HN and r/localllama frontpage.

Overall, these do not seem to be the sentiments coming from someone inside the ML space, but rather from an onlookers perspective.

> and now CUDA is all but irrelevant.

Lol this is so wrong it's cringe.

> There's now so many different and opinionated takes on how you should write high performant accelerator cluster code. I love it.

There are literally only 2: SIMT (ie the same as it always was) and tiles (ie Triton). That's it. Helion is just Triton with more auto-tuning (Triton already has auto-tuning).

In what alternative reality is that the case?

Pallas is the Triton equivalent in JAX land. There are some old auto tuning prototypes if you search for Pallas, like this https://github.com/jax-ml/jax-triton/pull/108

They have made it empty only.

>> out = torch.empty([m, n], dtype=x.dtype, device=x.device)

The accumulator has been initialized to zero, since well, they have to add stuff into it.

>> acc = hl.zeros([tile_m, tile_n], dtype=torch.float32)

> idiomatic

No as far as I have seen they generally try to not initialize if its not necessary.

> overhead

There is the memory bandwidth point as you might expect. But additionally when using high level interfaces like pytorch, when you write torch.zeros(512, 512) in pytorch, it launches a whole kernel (tens of micros) just for that line. So that's cpu -> gpu -> back to cpu, and then it does the next line, where it goes to gpu again and uses that memory. So in these cases you make sure to avoid it if its in a hot path. Ideally you want the 2nd kernel to do the initialization itself. When you write cuda c++ yourself this is how you typically do it. Helion being a compiler might be doing this optimization, but runtime based torch can't clearly.

It saves a kernel launch and memory bandwidth for a fill kernel. If you’re going to overwrite the data anyway, why bother?

You mean autotune? I think 10 minutes is pretty normal, torch.compile('max-autotune') can be much slower than that for large models.

Funny, I feel the same way about Triton. Performant Triton looks like CUDA (but with tiles!) except it's ten times harder to debug since it doesn't have the tooling NVIDIA provides.

If I had to run on AMD I'd rather deal with their hipify tooling.

What's the point of Triton compared to Gluon? What's the point of PyTorch compared to Triton?

One of the main values of Triton is that it significantly expanded the scope of folks who can write kernels - I think Helion could expand the scope even more.

It’s not that bad, but I’m not sure why this is relevant?

You’ll need an execution backend.

If you think of Triton as a "baseline", most other DSLs are lower-level than Triton, whereas this is higher-level.