[PR #113] [CLOSED] Fix Triton HIP backend crash: libdevice.rint missing on ROCm #155

New issue

Closed

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

bsakaguchi commented

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

Owner

📋 Pull Request Information

Original PR: https://github.com/Comfy-Org/comfy-kitchen/pull/113
Author: @wlkns-dev
Created: 8/12/2026
Status: ❌ Closed

Base: main ← Head: fix/hip-libdevice-rint

📝 Commits (1)

d9882a5 Fix Triton HIP backend crash: libdevice.rint missing on ROCm

📊 Changes

2 files changed (+26 additions, -3 deletions)

View changed files

📝 comfy_kitchen/backends/triton/quantization.py (+12 -1)
📝 comfy_kitchen/backends/triton/w4a8_int8.py (+14 -2)

📄 Description

Summary

triton.language.extra.hip.libdevice has no rint() (round-half-to-even), only nearbyint() with equivalent rounding semantics (differs only in floating-point exception flags, irrelevant for these int8 quantization kernels). The CUDA libdevice does have rint().
The generic triton.language.extra.libdevice dispatcher module always exposes rint as an attribute regardless of the active backend — it resolves the concrete implementation at kernel trace/compile time, not at Python import time. So a naive hasattr(libdevice, "rint") check on that dispatcher is always True and can't be used to detect the missing symbol; the concrete backend submodule (triton.language.extra.hip.libdevice) must be checked instead.
Without this fix, any AsymW4A8Int8 model (e.g. quantized diffusion models loaded via UNETLoader) crashes on the first sampling step on ROCm with:
```
AttributeError: module 'triton.language.extra.hip.libdevice' has no attribute 'rint'
```

Test plan

Tested on an AMD Radeon RX 9060 XT (gfx1200, ROCm 6.4, PyTorch 2.8.0+rocm6.4, Triton 3.4.0).
Before this fix: sampling crashed immediately on the first step with the AttributeError above.
After this fix: a full 20-step MiniMax H3 image-to-video generation (w4a8-quantized diffusion model) completes successfully end-to-end, no crash, no numerical/OOM issues observed.
CUDA path unaffected — torch.version.hip is None on CUDA systems, so the cuda.libdevice branch (unchanged behavior, still uses rint) is taken.

_{🔄 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/113 **Author:** [@wlkns-dev](https://github.com/wlkns-dev) **Created:** 8/12/2026 **Status:** ❌ Closed **Base:** `main` ← **Head:** `fix/hip-libdevice-rint` --- ### 📝 Commits (1) - [`d9882a5`](https://github.com/Comfy-Org/comfy-kitchen/commit/d9882a555877e722da5f3b45f5251620cfa9a082) Fix Triton HIP backend crash: libdevice.rint missing on ROCm ### 📊 Changes **2 files changed** (+26 additions, -3 deletions) <details> <summary>View changed files</summary> 📝 `comfy_kitchen/backends/triton/quantization.py` (+12 -1) 📝 `comfy_kitchen/backends/triton/w4a8_int8.py` (+14 -2) </details> ### 📄 Description ## Summary - `triton.language.extra.hip.libdevice` has no `rint()` (round-half-to-even), only `nearbyint()` with equivalent rounding semantics (differs only in floating-point exception flags, irrelevant for these int8 quantization kernels). The CUDA `libdevice` does have `rint()`. - The generic `triton.language.extra.libdevice` dispatcher module always exposes `rint` as an attribute regardless of the active backend — it resolves the concrete implementation at kernel trace/compile time, not at Python import time. So a naive `hasattr(libdevice, "rint")` check on that dispatcher is always `True` and can't be used to detect the missing symbol; the concrete backend submodule (`triton.language.extra.hip.libdevice`) must be checked instead. - Without this fix, any `AsymW4A8Int8` model (e.g. quantized diffusion models loaded via `UNETLoader`) crashes on the first sampling step on ROCm with: ``` AttributeError: module 'triton.language.extra.hip.libdevice' has no attribute 'rint' ``` ## Test plan - Tested on an AMD Radeon RX 9060 XT (`gfx1200`, ROCm 6.4, PyTorch `2.8.0+rocm6.4`, Triton `3.4.0`). - Before this fix: sampling crashed immediately on the first step with the `AttributeError` above. - After this fix: a full 20-step MiniMax H3 image-to-video generation (w4a8-quantized diffusion model) completes successfully end-to-end, no crash, no numerical/OOM issues observed. - CUDA path unaffected — `torch.version.hip is None` on CUDA systems, so the `cuda.libdevice` branch (unchanged behavior, still uses `rint`) is taken. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi

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