[PR #54] Change rounding method in quantization process (triton backend fix for amd gpus) #109

New issue

Open

opened 2026-10-26 22:20:28 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 22:20:28 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/Comfy-Org/comfy-kitchen/pull/54
Author: @patientx
Created: 6/26/2026
Status: 🔄 Open

Base: main ← Head: patch-1

📝 Commits (1)

3cb5be6 Change rounding method in quantization process

📊 Changes

1 file changed (+1 additions, -1 deletions)

View changed files

📝 comfy_kitchen/backends/triton/quantization.py (+1 -1)

📄 Description

fails on ROCm:
AttributeError("module 'triton.language.extra.hip.libdevice' has no attribute 'rint'") The generic from triton.language.extra import libdevice import resolves to triton.language.extra.hip.libdevice on AMD GPUs, which doesn't expose rint. This is the same gap as triton-lang#7135, where libdevice.round is declared but throws the same AttributeError when actually compiled on HIP — the documented workaround there is the same swap used below. Repro: Load an INT8/ConvRot-quantized checkpoint with the native UNETLoader on an AMD GPU with the triton backend enabled (COMFY_KITCHEN_BACKEND=triton or equivalent). Fails on first call into int8_linear → triton_quantize_rowwise.

tl.floor(x + 0.5) is portable across both CUDA and HIP backends (no branching needed), and the rounding difference vs. rint's round-half-to-even is negligible for INT8 quantization scales. Confirmed against current main (v0.2.12) — this is the only occurrence of libdevice.rint/libdevice.round in the codebase, so this single-line change should resolve ROCm compatibility for this path.

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/Comfy-Org/comfy-kitchen/pull/54 **Author:** [@patientx](https://github.com/patientx) **Created:** 6/26/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `patch-1` --- ### 📝 Commits (1) - [`3cb5be6`](https://github.com/Comfy-Org/comfy-kitchen/commit/3cb5be6faeee2b0410e599fbcf9cb2b1cb2ad8e7) Change rounding method in quantization process ### 📊 Changes **1 file changed** (+1 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `comfy_kitchen/backends/triton/quantization.py` (+1 -1) </details> ### 📄 Description fails on ROCm: AttributeError("module 'triton.language.extra.hip.libdevice' has no attribute 'rint'") The generic from triton.language.extra import libdevice import resolves to triton.language.extra.hip.libdevice on AMD GPUs, which doesn't expose rint. This is the same gap as triton-lang#7135, where libdevice.round is declared but throws the same AttributeError when actually compiled on HIP — the documented workaround there is the same swap used below. Repro: Load an INT8/ConvRot-quantized checkpoint with the native UNETLoader on an AMD GPU with the triton backend enabled (COMFY_KITCHEN_BACKEND=triton or equivalent). Fails on first call into int8_linear → triton_quantize_rowwise. tl.floor(x + 0.5) is portable across both CUDA and HIP backends (no branching needed), and the rounding difference vs. rint's round-half-to-even is negligible for INT8 quantization scales. Confirmed against current main (v0.2.12) — this is the only occurrence of libdevice.rint/libdevice.round in the codebase, so this single-line change should resolve ROCm compatibility for this path. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi added the

pull-request

label

2026-10-26 22:20:28 +00:00

bsakaguchi referenced this issue

2026-10-26 22:20:54 +00:00

[PR #109] [CLOSED] Fix HIP SageAttention argument binding #152