RDNA-aware autotune configs for the INT8 matmul kernels

[liminfei-amd]

What

The two INT8 matmul kernels ship 6 NVIDIA-tuned autotune configs (num_stages=3–4, block_n=256).
On AMD RDNA, num_stages=2 with deeper block_k is faster. This adds _int8_autotune_configs(),
which returns an RDNA-tuned pool when torch.version.hip is set and keeps the NVIDIA list unchanged
for CUDA. Both kernels share the resulting _INT8_MATMUL_CONFIGS.

Why it's safe

• Used only by the two INT8 matmul kernels (the only @triton.autotune blocks in the package).
• NVIDIA path unchanged (torch.version.hip gate).
• INT8 GEMM is int32-exact → changes speed only; outputs byte-identical.

Perf & testing

~1.03–1.11× per shape on gfx1100 (RDNA3) and gfx1201 (RDNA4). pytest tests/test_int8.py →
37 passed / 5 cuda-only failures on AMD (unchanged).

Fix: #68

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Organization UI

Review profile: ASSERTIVE

Plan: Pro Plus

Run ID: eedcbe10-47a8-45e9-abf9-76711d4f4096

📥 Commits

Reviewing files that changed from the base of the PR and between 29b5574 and a61d505.

📒 Files selected for processing (1)

• comfy_kitchen/backends/triton/quantization.py

---

📝 Walkthrough

Walkthrough

Changes

INT8 autotuning configuration

Layer / File(s)	Summary
Platform-specific autotune configuration and kernel wiring comfy_kitchen/backends/triton/quantization.py	Adds _int8_autotune_configs() with separate NVIDIA and HIP Triton configurations, caches the selected list in _INT8_MATMUL_CONFIGS, and uses it for _int8_matmul_dequant_kernel autotuning.

🚥 Pre-merge checks | ✅ 2

✅ Passed checks (2 passed)

Check name	Status	Explanation
Linked Issues check	✅ Passed	The PR implements RDNA-aware INT8 autotune configs for the Triton kernels while preserving the existing CUDA config pool, as requested.
Out of Scope Changes check	✅ Passed	The changes stay focused on INT8 matmul autotuning and AMD RDNA tuning, with no unrelated code paths introduced.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

✨ Simplify code

• Create PR with simplified code

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[liminfei-amd]

Thanks for merging Comfy-Org/ComfyUI#14862 (default-enabling the Triton backend on ROCm) — much
appreciated, and the quick review turnaround was great!

This PR is the companion on the kernel side: it tunes the INT8 matmul autotune configs for RDNA
(~1.03–1.11× faster on gfx1100/1201), so AMD users now on the default Triton path get faster
kernels too. It's a small, arch-gated change (NVIDIA configs untouched; INT8 GEMM is int32-exact so
outputs are byte-identical). @comfyanonymous — would you have a moment to take a look when
convenient? No rush at all, and thanks again!

[LuXuxue]

What about add waves_per_eu into the configs?
I scaned configs on my gfx1103 for anima and noobai models, and this is the configs i'm using:

        configs = [
        triton.Config({'block_m': 64, 'block_n': 64, 'block_k': 128, 'group_size_m': 8, 'waves_per_eu': 1}, num_warps=2, num_stages=2), #Anima01 Anima03
        triton.Config({'block_m': 128, 'block_n': 256, 'block_k': 128, 'group_size_m': 8, 'waves_per_eu': 1}, num_warps=8, num_stages=2), #Anima04
        triton.Config({'block_m': 64, 'block_n': 128, 'block_k': 64, 'group_size_m': 8, 'waves_per_eu': 4}, num_warps=4, num_stages=2), #SDXL05
        triton.Config({'block_m': 64, 'block_n': 64, 'block_k': 64, 'group_size_m': 8, 'waves_per_eu': 4}, num_warps=2, num_stages=2), #SDXL07
        triton.Config({'block_m': 64, 'block_n': 128, 'block_k': 128, 'group_size_m': 8, 'waves_per_eu': 1}, num_warps=8, num_stages=2), #SDXL08
        triton.Config({'block_m': 64, 'block_n': 256, 'block_k': 64, 'group_size_m': 8, 'waves_per_eu': 3}, num_warps=4, num_stages=2), #SDXL11
        triton.Config({'block_m': 64, 'block_n': 256, 'block_k': 64, 'group_size_m': 8, 'waves_per_eu': 2}, num_warps=8, num_stages=2), #SDXL12
        ]

gfx1103 is a apu, limited by memory bandwidth, dgpu may need to use different configs.

[liminfei-amd]

Thanks @LuXuxue! I checked waves_per_eu on our hardware — it helps on APUs (gfx1151: up to ~1.15x) and is within noise on dGPUs (gfx1201). So I've added a few waves_per_eu options to the RDNA autotune pool, letting the autotuner pick per shape/arch rather than hardcoding per-model configs. Thanks for the data!

[patientx]

I know this might not be the place to ask this but your recent change in comfyui's enabling triton backend for ROCM Comfy-Org/ComfyUI#14862 totally killed the int8 model usage for me. RDNA2, triton-windows 3.7.1.post27. I was using comfy without triton-backend and with custom node everything worked perfectly now neither native load diffusion model nor the custom node works, comfy just freezes and only a reset helps.