Live data from Hacker News

HipKittens: Fast and furious AMD kernels

hazyresearch.stanford.edu

71–80 of 94 posts

Re: HipKittens: Fast and furious AMD kernels

#71
post #5

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

Spending >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?

Re: HipKittens: Fast and furious AMD kernels

#72

Earlier quoted context omitted.

I didn't know you can you buy Google TPUs now?

You can buy older less capable TPUs https://www.seeedstudio.com/Coral-USB-Accelerator-p-2899.htm...

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

Re: HipKittens: Fast and furious AMD kernels

#73

Full 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 NVI…

> Cerebras isn't available anywhere.

That sounds like they're winning.

Re: HipKittens: Fast and furious AMD kernels

#74
post #48
post #17

Earlier quoted context omitted.

The MBAs have always been in charge to an extent. But the real issue is we don't want to invest in beating Nvidia on quality. Otherwise we wouldn't be doing stock buybacks and instead use the money on poaching engineers. The mindset is that we maintain a comfortable second place by creating a shittier but cheaper product. That is how AMD has operated since 1959 as a second source to Fairchild Semiconductor and Intel.…

Came into this thread hoping for good news about GPUs and instead there's some surprisingly thoughtful management discussion! > What they should do is bump TC by 60-70% and simultaneously lay off 50% of the engineers. Tell me you're an engineer without telling me you're an engineer. The problem is they don't know which half and they can't know. It's an issue of legibility and transparency - put yourself into the shoe…

> How can you possibly know who's good and who's not? Think of the information they have at hand - what the chain of command tells them. What if the chain of command itself is the problem? Think about how you yourself could protect a bad employee if you were a manager. You could! How can they possibly find the truth?

Senior managers should look at what people are actually doing. It doesn't take that much time. If tickets and PRs/MRs/changes are searchable by author, reviewer, and the files they touch (if they aren't, that's your problem right there) then it takes a few minutes to figure out who did the critical work, and who doesn't do much of anything.

In big tech, I've had senior managers (1-3 levels up) that do this, and ones that don't. The ones that do it are great managers. Under this type of manager, people are usually focused on making things actually work and making projects successful. The ones that don't do it can be good managers, but usually aren't. Under these types is where politics festers and dominates, because why wouldn't it? If you don't let the actual work guide your understanding, you're left with presentations and opinions of others.

When I do this (a few times a year), it takes 10 minutes for the easy cases, 1 hour for the hard cases, and once you do a few of these kinds of investigations in the same work area, you start to understand what the collaborators are doing before even looking at them specifically. So you're talking a few weeks of work for 100s of people. A few weeks a few times a year is not too much to ask someone to spend on their primary responsibility as a senior manager.

Past some point in scale, this does become impractical, I don't expect the CEO of a 10k person company to be doing this. But at that scale, the metrics are different anyways.

Re: HipKittens: Fast and furious AMD kernels

#75
post #2

One 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 a…

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

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

Re: HipKittens: Fast and furious AMD kernels

#76
post #68
post #40

Earlier quoted context omitted.

Does consumer hardware (non-MI) need proprietary kernel drivers for running rocm + pytorch?

No. But you might need a specific version of rocm built for your gpu. These are built on https://github.com/ROCm/TheRock Right now AI support on AMD is officially only on specific models. But they are working hard to turn this around to have broader support. And making progress.

Vulkan compute is also getting some good press as a local llm platform (at least on the linux side), will be interesting to see which crosses the line to "can ship production quality apps on this" first.

Re: HipKittens: Fast and furious AMD kernels

#77

Earlier quoted context omitted.

How far is Tinygrad from being able to represent/search the kind of optimisations listed in the article? i.e.: 1. data layouts to avoid local memory bank conflicts 2. read patterns from global memory to optimize L2 cache reuse 3. warp specialisation How complex is it to add these into tinygrad?

1 and 2 are supported, 1 you need to specify, 2 will be found with BEAM. We are working on reimplementing HipKittens in tinygrad, all the stuff is there to do it. See the amd_uop_matmul example. tinygrad doesn't support 3 yet, it's not needed on any AMD GPUs, and not needed on NVIDIA consumer. It wouldn't be hard to add, but it's important to figure out how it best fits with the existing abstractions. I think everyth…

Good luck with the AMD contract! I imagine HipKittens came at just the right time.

Re: HipKittens: Fast and furious AMD kernels

#78
> 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!

Re: HipKittens: Fast and furious AMD kernels

#79
Anyone know whether there are things built on 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.

Re: HipKittens: Fast and furious AMD kernels

#80
post #53

Earlier quoted context omitted.

I 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 rem…

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

I got one of those running whisper yesterday, hopeful the bigger llms will run shortly. You'd need rocm 7 which seems to be much better than 6.4 was.
Post reply on HN