12 ms·
Surprisingly fast AI-generated kernels we didn't mean to publish yet
- terhechte 1y ago[flagged]
- yahoozoo 1y agoVery cool. They used o3 and Gemini 2.5 Pro but unfortunately they don’t mention which one produced the better kernels.
- reliabilityguy 1y agoIs my understanding correct that they assumed a fixed size of the input? If so, why is it surprising that generic implementations in PyTorch are worse?
- GaggiX 1y agoPytorch uses different kernels depending on the input size. There is a reason why it's so massive to download.
- reliabilityguy 1y agoSure, some degree of customization is expected. However, I doubt that PyTorch implements every input size separately.
- saagarjha 1y agoGenerally tiling to a handful of sizes gets you pretty close to "specialized to this size" performance.
- reliabilityguy 1y agoSure, but the size can be anything.
- saagarjha 1y agoYes, you pick the size that’s closest
- Workaccount2 1y agoVery fascinating result, and it seems they wrote this blog post out of pure excitement to share their findings, and maybe to have someone throw cold water on it before publishing, ha. Who knows if this is the actual fabled path of "self improvement", but results like this are what we expect to find on such a path.
- suddenlybananas 1y ago> Who knows if this is the actual fabled path of "self improvement" Seems doubtful as this works only on an extremely well-defined evaluation function.
- EMIRELADERO 1y agoThat may be true, but this is the first example I've seen where the concept is successfully implemented in a noticeable way. It's just like image generation: the first iteration is the worst it will ever be.
- observationist 1y agoEach time you define another task well enough for the system to work, you generalize the system just a little bit - repeat enough times and you can start to expand, develop taxonomies of functions, precisely define function spaces and metrics for improvement. This might not be a bootstrap for recursive self improvement generally, but it could definitely inform the theory or design of a system that does bootstrap rsi.
- suddenlybananas 1y agoThat's an entirely different idea that may or may not work. This is not evidence of that.
- observationist 1y agoThe structure of their research - the process, the specific task, and the data they generate - will help inform how other research gets performed. Instead of GPU kernels, maybe the next task is something like neuron modules, looking for structures that improve on attention blocks, or things like that - each time you run through an experiment like this, you're creating foundational data upon which other experiments can be run and improved. Once you've done enough of them, you can generalize. It could be that the end result is the knowledge of strict boundaries of LLM capabilities, that they can only operate in specific domains, or only improve to a certain extent, and some currently unspecified defect limits the level of improvement. The underlying idea of specifying a domain and task conditions, then letting an LLM run thousands of experiments, is a great search technique. The hope is that there is no implicit defect and that the methodology will extend and generalize - it's not too complex a notion to think that you could have an LLM create a broad range of individual tasks, with a meta-goal of identifying better and more general recursive improvement processes and algorithms.
- thorum 1y agoMy takeaway - from this article, from Google’s AlphaEvolve [1], and the recent announcement about o3 finding a zero day in the Linux kernel [2] - is that Gemini Pro 2.5 and o3 in particular have reached a new level of capability where these ideas that were tried unsuccessfully with other models, suddenly just work. [1] https://deepmind.google/discover/blog/alphaevolve-a-gemini-powered-coding-agent-for-designing-advanced-algorithms/ https://deepmind.google/discover/blog/alphaevolve-a-gemini-p... [2] https://sean.heelan.io/2025/05/22/how-i-used-o3-to-find-cve-2025-37899-a-remote-zeroday-vulnerability-in-the-linux-kernels-smb-implementation/ https://sean.heelan.io/2025/05/22/how-i-used-o3-to-find-cve-...
- jiggawatts 1y agoGemini Pro 2.5 is the first AI that I can productively use for anything other than human language translation, but it's just barely crossed that threshold. Sometimes I get success hit rates below 20%. When 3.0 comes out, that... that's going to start getting a little scary.
- jacob019 1y agoWhat domain?
- jiggawatts 1y agoSRE / DevOps / coding mostly in the Azure and .NET ecosystems. The problems I have to solve tend to be the horrible ones that nobody has answers to, anywhere on the Internet, so unsurprisingly the AIs aren't good at it either. The trick has been to use the AIs for what they are good that, which used to be "nothing" for me at least, but now I can use them productively for certain "spot" tasks. Random examples: - Cross-language and cross-platform benchmarking of a bunch of different database clients to see how they stack up. I gave the AI a working example in one language and got it to whip up a series of equivalents with other DB drivers and languages. Sure, it's trivial, but it's way faster than doing it myself! - Crash dump analysis using WinDbg. I read somwhere that "vibe debugging" of kernel dumps totally works, so when I had an actual crash I gave it a go for laughs. With AI help I managed to extract the name of the specific file that had NTFS corruption and was crashing the server. Deleted the file, restored it from backups, and the server was good to go again! - If you ever watch the top mechanical engineers on YouTube, they all make their own tools instead of just buying them. Jigs, extenders, unusual sizes, etc... IT work is the same. As a recent example, I got Gemini to make me a code-AST rewriter for a specific issue I wanted to clean up in bulk across a huge code base. Using the Roslyn compiler SDK is a bit fiddly, but it spat out a working tool for me in under an hour. (This is not something you can solve with a script full of regex, it needed a proper parser to handle commented-out blocks and the like.)
- brrrrrm 1y agowhat's going to be interesting is to see the large space of fused kernels being tackled by AI generated code. that might include gemm + relu + gemm + a norm of some kind - which would be annoyingly exhaustive to 1. sweep with a tuner and 2. handwrite as a human
- AtlasBarfed 1y agoUh, what is a "kernel" in the sense of AI? Because it sure looks like this isn't an OS kernel.
- philipkglass 1y agoThis is GPU terminology: https://cvw.cac.cornell.edu/gpu-architecture/gpu-characteristics/kernel_sm https://cvw.cac.cornell.edu/gpu-architecture/gpu-characteris... A function that is meant to be executed in parallel on an attached GPU is called a kernel. In CUDA, a kernel is usually identified by the presence of the __global__ specifier in front of an otherwise normal-looking C++ function declaration.
- ekelsen 1y ago"FP32 is less common in modern ML workloads and often less optimized on recent hardware compared to FP16 or BF16, which may partly explain why it’s easier to achieve performance gains over PyTorch with FP32 kernels." People haven't spent time optimizing the fp32 versions of these kernels in years. This will be much more interesting if they can improve the kernels where developer effort has gone and that are actually used.
- suddenlybananas 1y agoI wonder if it's using known improvements from the fp16/bf16 kernels that are transferable to fp32?
- moralestapia 1y ago>People haven't spent time optimizing the fp32 versions of these kernels in years. Wow, so, you're basically saying the AI created new algos in a domain with no pre-existing solutions? Awesome!
- Aurornis 1y agoNo one said the AI created new algorithms nor that there weren’t pre-existing solutions. The implication was that the FP32 versions of these kernels have lagged behind the more popular versions. There was opportunity to translate the advancements from other kernels into these. Someone would need to look closely to see exactly what was done, but it’s premature to suggest anything like “new algos” or “no pre-existing solutions” This is a great use case for LLMs, though. I often do something similar where I make improvements to something I use most frequently and ask an LLM to translate that pattern to other similar parts of the code.
- moralestapia 1y ago>The implication was that the FP32 versions of these kernels have lagged behind the more popular versions. Help me understand this 'cause I'm a bit slow these days ... Does that mean optimized FP32 versions of these kernels were already there or not?
- adityamwagh 1y agoSometimes I think of LLMs as kind of a hive mind. It’s trained on thought processes of so many humans. I think that’s why it’s able to do these kinds of things given the fact that it has so much information and context compressed in weights.
- MangoToupe 1y agoThe market itself is also kind of a hive-mind metaphor. Worth thinking about.
- suddenlybananas 1y agoMaybe we could replace it with a central planning now that we can distill information.
- MangoToupe 1y agoWhoops you just did a communism
- gpm 1y agoA "vertical integration" in the capitalist world ;)
- MangoToupe 1y agoThis got a legitimate chortle out of me
- yieldcrv 1y agoa non-human standing committee following the directives of a trust could work
- MangoToupe 1y agoWhat like you want to govern by divining patterns of snake coils or bird guts?
- constantcrying 1y ago>and test for correctness by checking the numerical equality of the two outputs over many random inputs. This is fundamentally different to how any human would approach this problem. And also different to how some recent advances in this area were made, where AI actually came up with superior and correct algorithms. This approach also seems quite unfortunate and makes many of theses results somewhat doubtful.
- gotoeleven 1y agoHow else would you do the verification?
- constantcrying 1y agoSee: https://www.nature.com/articles/s41586-022-05172-4 https://www.nature.com/articles/s41586-022-05172-4 IIRC there was another paper recently, with similar methodology about computing xAx. These papers produce algorithms which aren't empirically correct, but provably correct. They do this by operating on a graph data structure, which describes the algorithm and then verifying the algebraic equality to the correct result. There is a substantial difference here. And I think utilizing algorithms which only are empirically correct can be dangerous.
- ekelsen 1y ago"the reference code is in the default FP32, and given a tolerance threshold (1e-02)" that's a huge tolerance and allows them to use fp16 operations to replace the "fp32" kernel.
- unignorant 1y agoyeah, it seems likely the underlying task here (one reasoning step away) was: replace as many fp32 operations as possible in this kernel with fp16. i'm not sure exactly how challenging a port like that is, but intuitively seems a bit less impressive maybe this intuition is wrong but would be great for the work to address it explicitly if so!
- AlotOfReading 1y agoOnly seems to have done that in a couple places, like the MatMul. The softmax kernel (https://github.com/ScalingIntelligence/good-kernels/blob/main/Softmax/src.py https://github.com/ScalingIntelligence/good-kernels/blob/mai...) seem to be entirely bog-standard, and the layernorm kernels are only slightly more interesting.
- ekelsen 1y agoI looked at the softmax kernel and the cast that it does from a float* to a float4* is extremely brittle -- it's trivial to break by offsetting the input slightly. Very likely a kernel for a standard library could not employ such a trick that relies on alignment of input pointers. Certainly not without a fallback.
- constantcrying 1y agoThis means the results are useless. Did they even check the relative error at all? Replacing float32 operations with float16 is also pointless. There is nothing to be gained by doing this, as it removes the actual accuracy advantage of float32s, which would the single most important reason to use that version of the algorithm.
- 1y ago
- JSR_FDED 1y agoCould this be used to create kernels for OpenCL, ROCm, etc?
- vessenes 1y agoBy far the most interesting part (after the 400% speed up in some cases) is the methodology: rather than hill climb on operations, they forced a language reasoning step between iterations to encourage diversity of search. This seems to have worked. Very very interesting.
- lucidrains 1y agooh wow, I was looking for use of islands or map-elites that I missed this.. thought it was the blandest mimetic evolution possible
- vessenes 1y agoJust anecdotally I feel like hill climbing on operations is just so slow; I’m not saying it doesn’t work, but it always feels one step away from brute force search. I really like the idea of just throwing stuff at the LLM and giving it access to old strong variants in context.
- deleted 1y ago[deleted]
- deleted 1y ago[deleted]
- david-gpu 1y agoDisclaimer: This used to be my bread and butter, but I'm really rusty after five years of not working on this sort of stuff. That said, after quickly skimming the example AI-generated kernel I am not seeing anything novel there. While working at nVidia I did see a handful of techniques that, frankly, blew my mind. Thus, I wonder what makes this AI-generated kernel faster than the standard pyTorch kernel, which I presume is simply delegating all the heavy lifting onto cuDNN. My guess, and it's just a guess, is that they are comparing the fastest AI-generated kernel they produced for a very particular set of parameters against whatever kernel cuDNN is picking for that same scenario, and perhaps the subsystem inside cuDNN that picks which kernel to execute out of the very large database it manages chose a suboptimal candidate. Researchers tend to completely ignore this issue and assume that cuDNN is always able to choose the very best kernel in every possible scenario, something that is just not realistic. Maybe there is something else going on, but these sort of "we have beaten this heavily optimized proprietary library" always seem to miss this very important point. Kind regards to any NVidia insiders who may read this. You guys are the brightest people I've ever met.
- zahlman 1y ago> Thus, I wonder what makes this AI-generated kernel faster than the standard pyTorch kernel All of this stuff is way outside my wheelhouse, but maybe "the standard pyTorch kernel" is just a low bar? (https://news.ycombinator.com/item?id=44144346 https://news.ycombinator.com/item?id=44144346)
- david-gpu 1y agoThanks for the link. I did have a quick look at the github repo yesterday, trying to find where this "default pytorch implementation" lives, but it wasn't immediately apparent due to the layers of abstraction in the code. There's little doubt in my mind that the authors of the blog post went after the largest performance delta between their best AI-generated code and the default pytorch implementation, without looking into real-world nuances like "picking a good kernel out of a zillion options" or "production-quality output accuracy". It's cool that this sort of thing is being researched, but the results need to be taken with a grain of salt. Maybe I'm too jaded.
- MangoToupe 1y ago> Our results are benchmarked on an Nvidia L40S At the very least they could have used consumer hardware. I don't even know how to parse that model it's so consumer-alien.
- Maxious 1y agoIt's the AD102 chipset ie. RTX 4090 with perfect binning (like the never released 4090 Ti would have had) and 48GB of VRAM soldered on https://www.techpowerup.com/gpu-specs/?architecture=Ada%20Lovelace&gpu=AD102&sort=generation https://www.techpowerup.com/gpu-specs/?architecture=Ada%20Lo...
- klingenm 1y agoThis sounds more like using AI (llm) as one small step, where the randomness in the output is used to implement a Genetic Algorithm, than being "AI-generated" (admittedly technically correct). (Edit, typo)
- FL33TW00D 1y agoTried a replication here. The LayerNorm kernel is not numerically stable so cannot be counted as valid. They only test with zero mean and unit std, so the catastrophic cancellation doesn't show up until after. EDIT: looks like they've since generated another one that is numerically stable! great work
- userbinator 1y agoAm I the only one who was enticed into this article by thinking they had AI generate an OS kernel?
- dgfitz 1y agoNope, I was as well.
- miki123211 1y agoI think how the authors of this post think about "AI agents" is really interesting. Most people think of agents like they think of human employees. They set up a limited number of agents to run in parallel (often just one), with each agent running in a loop and doing one task at a time. They're still in a world where you have a fixed (on the timescale of hours or days) number of employees, each employee can only do one thing at a time, and transferring tasks between employees is slow and costly. LLMs don't really work like that. You effectively have an infinite number of agents that you can conjure out of thin air at any time. There's no cost advantage to performing LLM requests in series rather than in parallel. If you realize this, the pattern of each agent fanning out and forking itself into as many sub-agents as are needed to fulfill the task becomes obvious. This is exactly what the authors have done. I think a better way to think of agents is as "tasks" or "jobs", like those you might find in Celery or sidekik, and apply the learnings from those.
- neom 1y agoFor fun last month I decided to see if i could build a fully functional business of agents. It's 175 python files (employees) build up of internal roles within those files (tasks). So what I have is 175 employees who are able to pass output around each other, understand the work, complete the work, understand where to send the output. The whole system has the ability to do around 275 base processes (same as a business at > 100MM arr) I started on a Friday afternoon and slept a little bit and finished on Monday afternoon. After I had it running I sent it to a VC friend to show them and they sent back the deck of a startup that is in stealth with $25MM doing it the exact same way. With 1 month and a designer and an engineer, I could have it mvp functional for anyone to use ($40k?). Times are changing. Here is kinda how it looks: https://s.h4x.club/9ZuO4XQR https://s.h4x.club/9ZuO4XQR / https://s.h4x.club/jkuB8ZED https://s.h4x.club/jkuB8ZED (I've evolved it a little since this, and if you're an engineer and look at my files and think, this guy is a moron: I know!:))
- literalAardvark 1y agoEngineers who would judge someone's frontier MVP like that are not even worth worrying about. This is epic work. Would love to see more of it but I guess you're gonna take it the startup route since you have connections. Best of luck.
- bgwalter 1y ago> They are performing close to or in some cases even beating the standard expert-optimized production kernels shipped in PyTorch. The PyTorch code base is NOT written by performance experts in any way. This is the wrong base line. Nothing about that code base is clean or hand optimized . The "AI" generation methodology seems to give many instructions and even descends into instruction trees, manually throwing away results etc. So it requires, as usual, extreme guidance.
- poltomo 1y agoBeating pytorch and tensorflow kernels has been easy to do with ml compilers since ~2018. You typically train and evaluate your model in one of these frameworks then hand off the computation graph to a compiler like Apache TVM or your hardware vendor’s proprietary one. They should test their kernels against those kernels. ML guided heuristic search over compute schedules is as old as 2013 (Halide for image processing)
- Mathnerd314 1y ago> we didn't mean to publish yet I was thinking this was about leaking the kernels or something, but no, they are "publishing" them in the sense of putting out the blog post - they just mean they are skipping the peer review process and not doing a formal paper.
- t-vi 1y agoNote that PyTorch's kernels are somewhat generic in shape. It has always been relatively easy to get speedups by specializing the shape, e.g. Apache TVM had that (back before it was "Apache" even).