13 ms·
HipKittens: Fast and furious AMD kernels
Related post: https://hazyresearch.stanford.edu/blog/2025-11-09-amd-brr https://hazyresearch.stanford.edu/blog/2025-11-09-amd-brr
- bratao 11mo agoOne thing I don't understand about Nvidia’s valuation is that right now a small number of algorithms have 'won,' such as Transformers. The data is very important. Compared to the past where customized code was much more common, such as modeling code and HPC, the ecosystem was very important and it was almost impossible to implement all CUDA and related code. Competitors now only need to optimize for a narrow set of algorithms. If a vendor can run vLLM and Transformers efficiently, a massive market becomes available. Consequently, companies like AMD or Huawei should be able to catch up easily. What, then, is Nvidia’s moat? Is InfiniBand enough?"
- wmf 11mo agoInfiniband is being replaced with UEC (and it isn't needed for inference). For inference there is no moat and smart players are buying/renting AMD or Google TPUs.
- mandelken 11mo agoI didn't know you can you buy Google TPUs now?
- knowitnone3 11mo agoYou can buy older less capable TPUs https://www.seeedstudio.com/Coral-USB-Accelerator-p-2899.html https://www.seeedstudio.com/Coral-USB-Accelerator-p-2899.htm...
- fooblaster 11mo agothese are not remotely like anything Google uses in the datacenter, even a decade ago.
- i80and 11mo agoThe Coral TPUs are closer if anything to what's in Pixel phones. In particular they're limited to iirc 8-bit integer types, which puts them in a very different category of applications compared to the kind of TPUs being talked about here.
- mattlondon 11mo agoYou can pay to use them https://cloud.google.com/tpu https://cloud.google.com/tpu
- amypetrik8 11mo agoWhat you really want to buy is a Ming Mecca chip. Original model came out around 2003, but they've been iterating. These things are bigger than AMD or nvidia silicon, actually even much larger than a gigantic Cerebras wafer, typically 500-900 million USD in price. As you could guess, Ming Mecca is not broadly publicized, historically used for NSA crypto cracking although now adapted to AI and used for data crunching from gathered messages. More recently all those gathered messages have been used for training strategic /tactical intelligence developments to oversee and deploy resources optimally via a cluster of, at least last I heard, 18 Ming Mecca v7 chips
- patagurbon 11mo agoDo you have evidence for this? I don’t think Nvidia is switching to Ultra Ethernet, just adding it to the product line-up
- wmf 11mo agoSorry, I don't mean Nvidia is adopting UEC (they probably hate it). I should have said UEC can substitute for Infiniband.
- LtdJorge 11mo agoThe vast amount of CUDA libraries for anything you can think of. I think there’s where they have the biggest leverage.
- bryanlarsen 11mo agoTo rephrase the OP's point: transformers et al are worth trillions. All the other CUDA uses are worth tens or hundreds of billions. They've totally got that locked up, but researchers is a smaller market than video games.
- observationist 11mo agoAI is going to be so ubiquitous, something principled and open is going to supersede cuda at some point, as HTML5 did for Flash. CUDA isn't like an x86 vs ARM situation where they can use hardware dominance for decades, it's a higher level language, and being compatible with a wide range of systems benefits NVIDIA and their competitors. They're riding out their relative superiority for now, but we're going to see a standards and interoperability correction sometime soon, imo. NVIDIA will drive it, and it will gain them a few more years of dominance, but afaik nothing in their hardware IP means CUDA compatibility sacrifices performance or efficiency. They're also going to want to compete in the Chinese market, so being flexible about interoperability with their systems gains them a bit of market access that might otherwise be lost. There's a ton of pressure on the market to decouple nvidia's proprietary software from literally everything important to AI, and they will either gracefully transition and control it, or it will reach a breaking point and someone else will do it for (and to) them. I'm sure they've got finance nerds and quants informing and minmaxing their strategy, so they probably know to the quarter when they'll pivot and launch their FOSS, industry leading standards narrative (or whatever the strategy is.)
- toasterlovin 11mo ago> as HTML5 did for Flash Uh, Flash died because Apple refused to support it on mobile Safari. Perhaps Flash would have died anyway, but that is the proximate cause. And Apple's competitors were falling over themselves to market Flash support as a competitive advantage vs. iPhone.
- o11c 11mo agoThe thing the "just optimize AI" crowd misses is that this isn't like optimizing a programming language implementation, where even the worst implementation is likely only 100x slower than a good implementation. AI is millions of times slower than optimal algorithms for most things.
- mountainriver 11mo agoTransformers aren’t really one thing, the way they are implemented is wildly different. If it wasn’t then vllm and TRL would be easy
- deleted 11mo ago[deleted]
- ekropotin 11mo agoIt’s all about deeply entrenched ecosystem NVIDIA had been building around CUDA for decades. It’d super hard to replicate this hardware-software platform. Plus strategic partnerships with cloud providers. And InfinityBand, yes
- ivape 11mo agoI don’t think NVDA will have anything like a real moat, and more like whatever the difference was between iOS and Android. The gist of it is, the big bang of AI has happened and that universe is rapidly expanding, just like it once did for smart phones. There is the Apple of AI which is NVDA, and then there is Android (AMD). Moats are irrelevant here because the universe has just started rapidly expanding for them. Apple didn’t really “win” out against Android, and it would be a very wrong way of measuring what actually happened. Yet, Apple could have been seen as more premium during various points of that timeline. The truth of the matter was, it was never a swimming race at any point in that smartphone timeline. It was simply a flood that you could convince yourself was an orderly race. I believe the same is happening now, and it’s in Nvidias interest to maintain the narrative that there is a race and they are winning it. Believing something like this during the smartphone era would have been foolish.
- ehnto 11mo agoThey also don't actually have a moat in the sense that they have patented technology keeping others out of the game. The other chip makers are coming for their lunch eventually.
- jillesvangurp 11mo agoYou are right to question their moat. My view on this is that there's a lot of pressure from essentially all other trillion dollar companies (MS, Google, Amazon, Apple, etc.) to not get locked into a NVidia only ecosystem. Each of those do their own chips. They also use Nvidia but not exclusively. An Android or IOS phone has no nvidia capable chips whatsoever. Neither do most laptops. Apple's M series CPUs don't support it at all typically. And with the exception of some gaming or workstation class laptops, most windows/linux laptops come with either AMD or Intel GPUs. Or lately Qualcomm ARM based architectures with custom GPUs. Nvidias valuation and moat are centered around data center class GPUs used for training. I don't think they effectively have that space to themselves for much longer. Google is already using their own TPUs at scale for both training and inference. They still use some Nvidia stuff but they seem to be able to keep that off the critical path for anything that needs to run at "Google scale". OpenAI just ordered a bunch of AMD hardware. A lot of AI engineers use Apple laptops that rely on the M series hardware. In short, the Cuda moat is shrinking. It's still relevant of course and there are a lot of tooling and frameworks that depend on it. That's why everybody still uses it. But not exclusively. And there's a lot of extremely well funded and active development to cut loose from it. AMD of course wants in. So does Intel. And so does everybody else. This HipKittens thing looks like it makes some big steps towards a more neutral software ecosystem.
- vagab0nd 11mo agoIf your competitor has a 5-year lead, and is working as hard as you are, or harder, then you are not gonna catch up any time soon. Also yes networking.
- dwheeler 11mo agoThat's only true if future improvements are easy to create as past ones, that customers care as much about those improvements, and there are no other differentiators. For example, many companies do well by selling a less capable but more affordable and available product.
- ACCount37 11mo agoBy far the easiest way to implement that "small number of algorithms" is with universal number-grinding hardware. Which also protects you against any architectural developments. Hardware takes a damn long time to make.
- wewewedxfgdf 11mo agoYou'd think AMD would swing in on something like this and fund it with the money needed to succeed. I have no knowledge of it but my guess is no, AMD never misses an opportunity to miss an opportunity - when it comes to GPUs and AI.
- LtdJorge 11mo agoFirst rule of AMD stock is nobody understands AMD stock. I guess it’s also the same for AMD’s software endeavors.
- elteto 11mo agoFrom the performance comparison table, basically AMD could be NVIDIA right now, but they aren’t because… software? That’s a complete institutional and leadership failure. Ironically, building chips is the actual _hard_ part. The software and the compilers are not trivial but the iteration speed is almost infinite by comparison. It goes to show that some companies just don’t “get” software. Not even AMD!
- wmobit 11mo agoI'd go so far as to say it's the exact opposite. It's faster and easier to change the hardware than the software.
- elteto 11mo agoCounterproof: attempt to modify your graphics card. Then attempt to modify a piece of code. Which one was easier?
- _lyxd 11mo agoYou're saying it like hardware and software are disjoint. You design hardware with software in mind (and vice versa); you need to if you want performance rivaling nvidia. This codesign, seeing their products are not only usable but actually tailored to maximize resource utilization in real workloads (not driven by w/e benchmarks), is where AMD seems to lack. Why oversimplify the premise and frame your take as some 'proof'. Just use the term counter-argument/example
- LtdJorge 11mo agoAhh, composable-kernel. The highest offender in the list of software that have produced unrecoverable OOMs in my Gentoo system (it’s actually Clang while compiling CK, which uses upwards of 2.5GB per thread).
- slavik81 11mo agoI was recently reviewing a CK package for Debian. My test build crashed due to OOM using -j32 on a 64GB workstation, so I tried with -j1 to be safe. That completed successfully after 190 hours! I think I may need to reduce the number of architectures it's built for to successfully compile it on the official Debian buildd infrastructure, but my (unverified) understanding is that most of its reverse dependencies only need the header-only parts of the library anyway. I'm told they're working on improving the build times via a few different methods.
- LtdJorge 11mo agoSame, -j32 with 64GB on a 3950x. I use 50% of ZRAM, but it’s still not enough most of the times, so I had to make a config called less-threads that only uses 24, with ZRAM enabled. I also use OOMD, but I have to work on separating my systemd units better, OOMD has killed my greetd session before, and with that my entire tree of userland processes :D
- nalllar 11mo agoSpending >10 minutes doing template instantiation for a single kernel for a single ISA is impressive! `device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance`, what are you doing to our poor friend clang?
- LtdJorge 11mo agoAnd they say Rust is slow!
- georgehotz 11mo agoFull disclosure, we have a contract with AMD to get Llama 405B training on MI350X on MLPerf. Things are turning around for AMD. If you have an AMD card, go to pytorch.org, click Linux+ROCm and install PyTorch. 3 years ago, this was hopeless. Today, most mainline things work. I ran nanochat on MI300X and it just worked. I think that's true about MI350X now too. The MI350X machine is stable. They are clearly behind NVIDIA, nobody doubts that. And a lot of investment into software will be required to catch up, ecosystem, compiler, and driver. But 2 years ago they seemed hopeless, now they don't. Things take time. HipKittens is a great codebase to study to see where AMD's LLVM backend is still lacking; compare it to the CUDA Kittens. For training, it's NVIDIA and Google in first. AMD in second. And nobody in third. Intel and Tenstorrent are not remotely close. Huawei examples segfaulted. Groq gave up selling chips. Cerebras isn't available anywhere. Trainium had a 5 day wait time to get one instance and I lost interest.
- latchkey 11mo agoAs CEO of an AMD NeoCloud for the past 2 years, it is so nice to hear all this and also see the turn around. It is what I bet my business on from the start and I can concur with what George is saying 100%. The out of box experience can be a bit rough around the edges on bleeding edge stuff, but it isn't anything near as bad as it used to be. For example, a month ago nanochat wasn't working well and now it is. The important thing is that people now care enough to make it work. At the end of the day, AI does need viable options. Having a monopoly on all AI hardware and software might be a good thing for share holders, but isn't a good thing for what is looking like a fundamental technology, akin to the internet.
- ivape 11mo agoThat’s interesting, I was specifically looking for AMD hardware being offered by neoclouds, they seem to be rare. I like your bet though. The difference between NVDA and AMD has never really existed on a hardware level for decades. AMD has always been on par, and software is software, it will catch up. AMD will be a stock many people will miss because the opportunity has presented itself at the height of AI bubble talk, and this will leave many in the dust. Doubling and tripling of their market cap is pretty much a forgone conclusion.
- villgax 11mo agoTotally ignored B300 for some reason
- semessier 11mo agowithout having implemented inference, just by looking at it from a math perspective this is base linear algebra/BLAS. I am very much wondering what a lean inference optimized API with covering 80% of all use cases across dtypes and sparsity would look like. Probably a far cry from what's in CUDA and probably all that's needed for practical inference.
- 999900000999 11mo agoWith these new developments, are there any implications for getting LLMs running well on consumer AMD chips ? For example, the following laptop which I'm thinking of picking up, has both a strong AMD CPU/IGPU and a RTX 5080. Could we see the AMD side competing with the RTX? I know a dedicated gpu will always be faster though. >HP OMEN MAX 16-ak0003nr 16" Gaming Laptop Computer - Shadow Black Aluminum AMD Ryzen AI 9 HX 375 (2.0GHz) Processor; NVIDIA GeForce RTX 5080 16GB GDDR7; 32GB DDR5-5600 RAM; 1TB Solid State Drive
- ehnto 11mo agoI run Qwen3 Coder 30b through Ollama on an RTX7900XTX. It works great, I suspect some load gets passed to the 32gb system memory and Ryzen 7 CPU. It's not quite as fast as like Sonnet 4 from an API, but it's really not that bad. It's really great for quick questions so I don't have to google stuff, and it's probably Sonnet4 level of competency at achieving coding tasks. No API served model has been fast enough to remove the urge to do something else while waiting for bigger tasks, so the UX is more or less the same in that regard. Opencode + ollama + Qwen3 Coder has been a very reasonable alternative to ClaudeCode with Sonnet4. That is amazing for something running locally. It is possible that if you actually need AI to be doing all your coding, that you're going to feel differently about the setup. But as a small assistant it's great.
- electroglyph 11mo agonot the best model to use as a showcase, it's blistering fast on anything that isn't a toaster
- ehnto 11mo agoGreat! That's what I am pointing out, it's a 30b param model that fits into an AMD card and runs great. That's what we want.
- christkv 11mo agoThat's great I have been eyeing a Strix Halo and was wondering how well smaller models are doing. This is great news from the perspective of running local agents.
- jiehong 11mo ago> what is raw assembly? can't understand it? that's the point! Raw assembly vs cooked assembly? Also, I think this attitude wasn’t the most common on CPUs, and people used to write assembly by hand just fine (and sometimes some still do). I think we shouldn’t be afraid of assembly like that. Compilers could write that assembly in the end, just like the do for CPUs!
- yunnpp 11mo agoYeah, comments like these really make you question the authors' background in optimization. Never mind that AMD actually publishes ISA specs for all of their graphics IPs -- it is not their point that you don't understand it -- what's holding GPU programming back is often that the underlying assembly primitives are not exposed in the high level languages. I also do wonder what 'raw assembly' is supposed to be. Is it like sushi? Perhaps it is left as future work in the paper for the authors to answer.
- JonChesterfield 11mo agoAnyone know whether there are things built on https://github.com/HazyResearch/ThunderKittens https://github.com/HazyResearch/ThunderKittens? I think this is a port of that to HIP, where generally ports of cuda things to hip are of vague professional interest, but much more so if the library is used by other things.
- nextworddev 11mo agoLong $amd?