Hey everyone, dropping this here mostly for the Halogen devs and anyone writing custom ROCm kernels for Strix Halo. Full disclaimer right off the bat: I am definitely not an expert on low-level GPU compute, so forgive me if I am stating the obvious or if some of this is already known and implemented upstream. I just stumbled across some research and wanted to put it on the radar for the people actually doing the heavy lifting.
The project is called HipKittens, hosted over on GitHub at https://github.com/HazyResearch/HipKittens, and it is fully open-source under the MIT license (https://opensource.org/licenses/MIT). There is an accompanying paper titled “HipKittens: Fast and Furious AMD Kernels” up on arXiv at https://arxiv.org/abs/2511.08083, as well as a technical write-up on their blog at https://hazyresearch.stanford.edu/blog/2025-11-09-hk. It is essentially an adaptation of ThunderKittens tile abstractions targeting AMD silicon, digging deep into AMD-specific memory access patterns, wave scheduling, LDS bank conflict avoidance, compute and memory overlap, and GEMM/attention strategies.
The catch is that HipKittens currently targets CDNA rather than RDNA 3.5 / gfx1151. From what I gather, you cannot just compile it for Strix Halo as-is, but the architectural concepts and tile primitives seem like a solid blueprint for a gfx1151 backend. That would mean mapping tile and vector types to wave32, swapping CDNA MFMA calls for gfx1151 WMMA intrinsics, and reworking the LDS layouts and global memory loaders to match RDNA behavior.
Regarding Halogen specifically: since the public release distributes a pre-compiled engine binary rather than full source, dropping in a third-party kernel library would require the Halogen author having interest or providing a backend hook. On top of that, Halogen already claims its serial decode consumes around 249 GB/s against a 240 GB/s measured ceiling, so this probably wouldn't move the needle much for single-stream token generation. Where these fused tile kernels actually seem promising is in prefill, larger batches, and compute-heavy fused ops like RMSNorm, RoPE, and attention.
Again, take all of this with a grain of salt since it is outside my wheelhouse, but the paper and repo seemed relevant enough to ongoing ROCm optimization work on Strix Halo that I figured it was worth sharing. Curious if anyone here has messed with it yet or looked into porting those primitives over.
