Helion: A high-level DSL for performant and portable ML kernels

(pytorch.org)

149 points | by jarbus 9 days ago ago

51 comments

  • markush_ 3 days ago ago

    Interesting choice from PyTorch to release yet another DSL, on positive side it's one more point in the design space on the other hand it's even more difficult to choose the right technology among Triton, Gluon, CuTe, ThunderKittens and a few others.

    • chillee 3 days ago ago

      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.

  • darknoon 3 days ago ago

    The developers also gave a talk about Helion on GPU Mode: https://www.youtube.com/watch?v=1zKvCLuvUYc

  • bobajeff 3 days ago ago

    It's good to see more effort for making things not device specific but I only see benchmarks for NVIDIA B200 and AMD MI350X. Also what's the experience of using one of these Python DSLs like? Are the tools good enough to make code completion, jump to definition, setting breakpoints, watching variables, copying as expression etc. nice?

  • brap 3 days ago ago

    Asking as someone who is really out of the loop: 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, and the infra would be mostly commoditized.

    • embedding-shape 3 days ago ago

      > 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.

      • westurner 3 days ago ago

        How often do hardware optimizations get created for lower level optimization of LLMs and Tensor physics? How reconfigurable are TPUs? Are there any standardized feature flags for TPUs yet?

        Is TOPS/Whr a good efficiency metric for TPUs and for LLM model hosting operations?

        From https://news.ycombinator.com/item?id=45775181 re: current TPUs in 2025; "AI accelerators" :

        > How does Cerebras WSE-3 with 44GB of 'L2' on-chip SRAM compare to Google's TPUs, Tesla's TPUs, NorthPole, Groq LPU, Tenstorrent's, and AMD's NPU designs?

        • almostgotcaught 3 days ago ago

          this is like 5 different questions all across the landscape - what exactly do you think answers will do for you?

          > How often do hardware optimizations get created for lower level optimization of LLMs and Tensor physics?

          LLMs? all the time? "tensor physics" (whatever that is) never

          > How reconfigurable are TPUs?

          very? as reconfigurable as any other programmable device?

          > Are there any standardized feature flags for TPUs yet?

          have no idea what a feature flag is in this context nor why they would be standardized (there's only one manufacturer/vendor/supplier of TPUs).

          > Is TOPS/Whr a good efficiency metric for TPUs and for LLM model hosting operations?

          i don't see why it wouldn't be? you're just asking is (stuff done)/(energy consumed) a good measure of efficiency to which the answer is yes?

          • westurner 3 days ago ago

            > have no idea what a feature flag is in this context nor why they would be standardized (there's only one manufacturer/vendor/supplier of TPUs).

            X86, ARM, and RISC have all standardized on feature flags which can be reviewed on Linux with /proc/cpuinfo or with dmidecode.

              cat /proc/cpuinfo | grep -E '^processor|Features|^BogoMIPS|^CPU'
            
            There are multiple TPU vendors. I listed multiple AI accelerator TPU products in the comment you are replying to.

            > How reconfigurable are TPUs?

            TIL Google's TPUs are reconfigurable with OCS Optical Circuit Switches that can be switched between for example 3D torus or twisted torus configurations.

            (FWIW also, quantum libraries mostly have Line qubits and Lattice qubits. There is a recent "Layer Coding" paper; to surpass Surface Coding.)

            But classical TPUs;

            I had already started preparing a response to myself to improve that criteria; And then paraphrasing from 2.5pro:

            > Don't rank by TOPS/wHr alone; rank by TOPS/wHr @ [Specific Precision]. Don't rank by Memory Bandwidth alone; rank by Effective Bandwidth @ [Specific Precision].

            Hardware Rank criteria for LLM hosting costs:

            Criterion 1: EGB (Effective Generative Bandwidth) Memory Bandwidth (GB/s) / Precision (Bytes)

            Criterion 2: GE (Generative Efficiency) EGB / Total Board Power (Watts)

            Criterion 3: TTFT Potential Raw TOPS @ Prompt Precision

            LLM hosting metrics: Tokens Per Second (TPS) for throughput, Time to First Token (TTFT) for latency, and Tokens Per Joule for efficiency.

    • anvuong 3 days ago ago

      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.

    • brrrrrm 3 days ago ago

      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

  • dachworker 3 days ago ago

    I'm super excited to give this one a spin. It seems like a neat idea, Triton, but simpler and with automatic autotuning. My head is spinning with options right now. I love how everyone was hyping up CUDA this and CUDA that a couple of years ago, and now CUDA is all but irrelevant. There's now so many different and opinionated takes on how you should write high performant accelerator cluster code. I love it.

    It's also kinda of ironic that right now in 2025, we have all this diversity in tooling, but at the same time, the ML architecture space has collapsed entirely and everyone is just using transformers.

    • embedding-shape 3 days ago ago

      > 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.

    • almostgotcaught 3 days ago ago

      > 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).

      • the__alchemist 3 days ago ago

        Even for non-ML things like chem simulations: CUDA (and cuFFT) are more pleasant to use than Vulkan Compute and vkFFT.

        • ozgrakkurt 3 days ago ago

          I just learned the graphics api of vulkan, can’t imagine anything being less pleasant than vulkan

          • porridgeraisin 3 days ago ago

            Yeah it's quite something. If anyone wants a preview, here's the triangle hello world in vulkan: https://gist.github.com/Overv/7ac07356037592a121225172d7d78f...

            But then again, I've heard that it's this low level because its meant for engine developers.

            • simlevesque 3 days ago ago

              Oh wow that's horrible.

            • anvuong 3 days ago ago

              Really? How low is this level actually? Because I remember my OpenGL class' professor did this in less than 50 lines.

              • porridgeraisin 3 days ago ago

                Imagine writing GlCreateContext yourself, for starters, as has been done in the link I posted.

    • pjmlp 3 days ago ago

      In what alternative reality is that the case?

  • sega_sai 3 days ago ago

    I switched from pytorch to jax just before triton appeared. Does anyone know how jax compares to this autotuning machinery in pytorch ? I know jax does jit, but i don't have a good intuition if jit is better than this type of autotuning.

  • undefined 3 days ago ago
    [deleted]
  • ballpug 3 days ago ago

    Compiling a kernel after assemblage in low-level object oriented languages either uses stable kernel or the cargo fuzzed raw_spinlock code.

    Helion abstracts syntax and design for calculating λ-functions, which converts language in a kernel config.

  • uoaei 3 days ago ago

    Tangential question related to the example kernel: in GPU programming is it idiomatic/standard to initialize the out array as zeros rather than empty? are the performance savings negligible?

    • porridgeraisin 3 days ago ago

      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.

    • saagarjha 3 days ago ago

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

  • mshockwave 3 days ago ago

    Is it normal to spend 10minutes on tuning nowadays? Do we need to spend another 10 minutes upon changing the code?

    • anvuong 3 days ago ago

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

      • Mars008 3 days ago ago

        Add to that it can be done only once by developers before distribution for major hardware. Configs saved. Then on client side selected.

  • bwfan123 3 days ago ago

    I dont get the point of helion as compared to its alternatives like gluon.

    For best performance I would presume one needs low-level access to hardware knobs. And, these kernel primitives are written one-time and reused. So, what is the point of a DSL that dumbs things down as a wrapper around triton.

    • krapht 3 days ago ago

      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.

      • saagarjha 3 days ago ago

        Performant Triton programs are usually simpler and shorter than their CUDA equivalents. This alone makes it easier to write, and I would argue that it helps with debugging too because the model provides a lot more guarantees on how your code executes. That said, some of the tooling is notably poor (such as cuda-gdb support).

        • krapht 3 days ago ago

          Agree on shorter, disagree on simpler. The hard part of understanding GPU code is knowing the reasons why algorithms are the way they are. For example, why we do a split-k decomposition when doing a matrix multiplication, or why are we loading this particular data into shared memory at this particular time, with some overlapping subset into registers.

          Getting rid of the for loop over an array index doesn't make it easier to understand the hard parts. Losing the developer perf and debug tooling is absolutely not worth the tradeoff.

          For me I'd rather deal with Jax or Numba, and if that still wasn't enough, I would jump straight to CUDA.

          It's possible I'm an old fogey with bias, though. It's true that I've spent a lot more time with CUDA than with the new DSLs on the block.

          • saagarjha 3 days ago ago

            I don’t think it is possible to write high performance code without understanding how the hardware works. I just think staring at code that coalesces your loads or swizzles your layouts for the hundredth time is a waste of screen space, though. Just let the compiler do it and when it gets it wrong then you can bust out the explicit code you were going to write in CUDA, anyway.

    • chillee 3 days ago ago

      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.

  • doctorpangloss 3 days ago ago

    Is contributing to Triton so bad? It looks like the blocker is usually LLVM.

    • saagarjha 3 days ago ago

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

  • singularity2001 3 days ago ago

    Anything as long as I don't have to touch propriety cuda and mpx

    • saagarjha 3 days ago ago

      You’ll need an execution backend.

  • jarbus 3 days ago ago

    I posted this 5 days ago, how did this resurface?

  • maknee 3 days ago ago

    How does this compare against other DSLs?

    • chillee 3 days ago ago

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

  • a-dub 3 days ago ago

    numba for gpu kernels... cool!