another day, another IMA is too dangerous to debug
(3D modelling is the only safe permitted use)
Thank you for everyone who commented on other public MXFP8 kernels that I didn't know about. At least for the toy 8192^3 benchmark shape here, K3-generated Gluon is the best 😳
I haven't checked all the kernels closely but I'm sure they can shine under different GEMM shapes + tuning knobs. Having an assemble of kernels is useful for autotuning - it is effortless thanks to agents now.
Show more
I couldn't find a public performant MXFP8 GEMM on MI355X, so I just asked @KimiDevs K3 to write one, and it beats torch._scaled_mm by >2x times 🤯
This is an adaptation of the official Gluon MXFP4 example. FlyDSL adaptation is WIP.
@AnushElangovan you guys should try K3 for kernel engineering if you haven't!
Show more
I couldn't find a public performant MXFP8 GEMM on MI355X, so I just asked @KimiDevs K3 to write one, and it beats torch._scaled_mm by >2x times 🤯
This is an adaptation of the official Gluon MXFP4 example. FlyDSL adaptation is WIP.
@AnushElangovan you guys should try K3 for kernel engineering if you haven't!
Show more
Guess I have no choice but to use K3
Recently I tried using K3 for Triton kernels and I was surprised by how good it was. The generated kernels are definitely very different from how Codex would have done it.
Not a proper experiment but I asked Codex (Sol medium) and K3 (high) to work a the same problem. K3 came out on top by a large margin.
@Kimi_Moonshot @KimiDevs
Kimi-K3 was served (locally) with @vllm_project by @inferact 😎
Show more
looks like writing PTX is the most portable way to maintain perf (well ig ptxas can regress but that's out of non-NVIDIA folks' hands)
Thank you @pranjalssh for the shoutout 🥰
Some thoughts:
- Reading lane_id is interesting... Usually I do warp_uniform() on warp_id (broadcast from lane0) + elect.sync. Wonder if there's perf advantage (at the mercy of ptxas)
- ftw
Show more
Sequel for my H100 blog is out now.
We implement Blackwell matmul for NVFP4 from scratch, and outperform cuBLAS by 4.7% at N=8192.
This time we have something better than Hilbert curves - that exact-fits for the Nvidia hardware.
Full blog post with all details:
Show more
X the everything app: PyTorch issue report
Together with the new primitives API, CuteDSL 4.7.0 now has explicit host-side TMA creation, mirroring CUDA C++'s cuTensorMapEncodeTiled
no more make_tiled_tma_atom and tma_partition 🥳
Thanks @Leik0w0 for sharing with me this
Show more
Can confirm, @rogerw0108 doesnt sleep
The @vllm_project maintainers at @inferact 🚀 are some of the most cracked engineers in the world. They’re building one of the inference engines that powers much of the world’s intelligence—and doing so with remarkable dedication, kindness, and hard work.
Show more
TIL: PyTorch conv2d has unfused bias 🙃
@tenderizzation whyyy???
We are getting TORCH_COMPILE_LLM_URL/API_KEY
@tenderizzation just a matter of time until LLM is an inductor backend as well
more reasons to avoid convnet in 2026
K3 is a good model for local LLM
