Comfy-Org / Comfy-Org/comfy-kitchen
INT8 matmul autotune uses NVIDIA-tuned configs on AMD RDNA
- Dominant language
- Python
- Stars
- 220
- Forks
- 91
- Avg merge
- 1d 7h
- Merged PRs (30d)
- 12
Description
The two INT8 matmul Triton kernels ship 6 autotune configs tuned for NVIDIA
(`num_stages=3–4`, `block_n=256`). On AMD RDNA these are suboptimal — RDNA prefers shallow
software pipelining (`num_stages=2`) and deeper `block_k`.
A `triton.testing.do_bench` sweep on gfx1100 (RDNA3) and gfx1201 (RDNA4) shows an RDNA-tuned
`num_stages=2` pool is ~1.03–1.11× faster per shape across typical DiT/SD INT8 GEMM shapes.
**Proposal:** add `_int8_autotune_configs()` that returns an RDNA pool when `torch.version.hip`
is set, keeping the NVIDIA list unchanged for CUDA. INT8 GEMM is int32-exact, so this changes
speed only (outputs byte-identical). PR ready.
Contributor guide
Research direction
Locate the two INT8 matmul Triton kernels and inspect their existing six NVIDIA autotune configs, then trace the torch.version.hip entry point. Done means an RDNA pool is selected on HIP while the NVIDIA list remains unchanged on CUDA, with byte-identical outputs and the reported gfx1100/gfx1201 benchmark improvement.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python, pytorch
- Domain
- machine-learning, performance
- Issue type
- Feature
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 55/100