Change rounding method in quantization process (triton backend fix for amd gpus)#54
Open
patientx wants to merge 1 commit into
Open
Change rounding method in quantization process (triton backend fix for amd gpus)#54patientx wants to merge 1 commit into
patientx wants to merge 1 commit into
Conversation
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.
|
✅ All contributors have signed the CLA. Thank you! This PR is ready to be merged. |
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Organization UI Review profile: ASSERTIVE Plan: Pro Plus Run ID: 📒 Files selected for processing (1)
📝 WalkthroughWalkthroughThe Triton INT8 row-wise quantization kernel now rounds ChangesINT8 quantization rounding
🚥 Pre-merge checks | ✅ 2✅ Passed checks (2 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
✨ Simplify 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. Comment |
Author
|
I have read and agree to the Contributor License Agreement |
comfy-legal
added a commit
to Comfy-Org/comfy-cla
that referenced
this pull request
Jun 26, 2026
Author
|
Just loading a model converted to int8-convrot without this patch results in this log from a failed run : |
Author
|
this PR #56 for eager works more reliably on amd rdna2 although triton still looks faster |
patientx
changed the title
upgrading from |
Author
there was a reason I kept using 3.6.0.post25 and I remembered that reason after trying the 3.7.0, yes this lets me use triton backend without a PATCH but it is still slower than node route -tried 8 gens on both- BUT more importantly ... That version + using native loader straight up stops generation in like 3rd or 4th time and after a few seconds whole PC freezes, I monitor vram etc. no temp raising, no gpu usage maxing etc. It just kills the comfy & windows. Only a hard reset solves it. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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.