Comfy-Org/comfy-kitchen
1
0
Fork 0
mirror of https://github.com/Comfy-Org/comfy-kitchen.git synced 2026-11-04 23:37:42 +00:00
Code Issues 118 Projects Packages Wiki Activity

[PR #238] [Feat] [RDNA1] Add RDNA1 INT8 Attention backend separate from the existing HIP backend #250

New issue
Open
opened 2026-10-26 22:22:05 +00:00 by bsakaguchi · 0 comments
bsakaguchi commented 2026-10-26 22:22:05 +00:00
Owner
Copy link

📋 Pull Request Information

Original PR: https://github.com/Comfy-Org/comfy-kitchen/pull/238
Author: @Zakhrov
Created: 10/10/2026
Status: 🔄 Open

Base: main ← Head: rdna1-backend

📝 Commits (10+)

  • 5bd4f17 Add RDNA1 backend separate from the HIP backend
  • 9c485ba Enhance README.md for rdna1 backend: clarify separation from HIP and outline unique requirements
  • 62d61dc Refactor RDNA1 backend: Remove unneeded WMMA and FP8 structures, update GEMM and quantization kernels
  • 7509aab Add test for small weights to ensure scaled rows remain in fp16 range
  • 80b856c Add new operations for depthwise convolution, direct convolution, GEMV, and packed scaling
  • 3c8754b Merge branch 'main' of https://github.com/Comfy-Org/comfy-kitchen into rdna1-backend
  • 19b8e11 Add Winograd F(2x2, 3x3) convolution implementation for fp16 in rdna1 backend
  • ba50e81 Add q4_k_linear implementation and related quantization functions for GGUF Q4_K weights
  • e946118 Merge branch 'main' of https://github.com/Comfy-Org/comfy-kitchen into rdna1-backend
  • a1e2f74 Add support for residual connections and input activation in w4a8_int8_linear

📊 Changes

96 files changed (+60819 additions, -103 deletions)

View changed files

📝 MANIFEST.in (+2 -0)
📝 README.md (+150 -45)
📝 comfy_kitchen/__init__.py (+230 -4)
➕ comfy_kitchen/backends/amd.py (+28 -0)
📝 comfy_kitchen/backends/eager/__init__.py (+82 -1)
📝 comfy_kitchen/backends/eager/conv3d.py (+89 -0)
📝 comfy_kitchen/backends/eager/quantization.py (+70 -0)
➕ comfy_kitchen/backends/rdna1/CMakeLists.txt (+263 -0)
➕ comfy_kitchen/backends/rdna1/README.md (+199 -0)
➕ comfy_kitchen/backends/rdna1/__init__.py (+4110 -0)
➕ comfy_kitchen/backends/rdna1/dlpack_bindings.cpp (+3079 -0)
➕ comfy_kitchen/backends/rdna1/embed_hsaco.cmake (+18 -0)
➕ comfy_kitchen/backends/rdna1/epilogue.h (+129 -0)
➕ comfy_kitchen/backends/rdna1/fp8_utils.h (+183 -0)
➕ comfy_kitchen/backends/rdna1/gemm_common.h (+100 -0)
➕ comfy_kitchen/backends/rdna1/gemm_simt.h (+853 -0)
➕ comfy_kitchen/backends/rdna1/hadamard.h (+954 -0)
➕ comfy_kitchen/backends/rdna1/launchers.h (+99 -0)
➕ comfy_kitchen/backends/rdna1/miopen_winograd/Conv_Winograd_v30_3_1_fp32_f2x3_dilation2.s (+42 -0)
➕ comfy_kitchen/backends/rdna1/miopen_winograd/Conv_Winograd_v30_3_1_fp32_f2x3_stride1.s (+42 -0)

...and 76 more files

📄 Description

What this PR does

Adds an optional rdna1 backend to comfy-kitchen for RDNA1 (gfx1010, e.g. RX 5600M / 5700 XT), built for that
one target from its own copy of the HIP sources. It registers only for gfx1010,
declines tensors on any other device, and leaves the hip backend and every other
GPU untouched.

Why

On gfx1010 the eager backend cannot run INT8 models at all: its INT8 matmul is
torch._int_mm, which PyTorch sends to hipBLASLt, and ROCm ships no hipBLASLt kernels
for gfx1010. It also has no INT8 attention, so ComfyUI's --use-ck-attention refuses
to start, and Triton is no alternative since it miscompiles on this target. This
backend supplies those kernels itself and never calls torch._int_mm.

Hardware constraints

gfx1010 is a "transition" architecture between the older Vega and Navi 20 chips. It has:

  • no matrix cores (WMMA was introduced with RDNA3): every GEMM and attention tile is computed in vector
    arithmetic;
  • no dot-product instructions (Vega, gfx1011, gfx1012, and gfx103x aka RDNA2 have them ) (v_dot4/sdot4): int8 products go through
    v_mad_i32_i24, or through fp16 FMAs, in which int8 values are exact;
  • no bf16 arithmetic: bf16 inputs are widened;
  • wave32 natively (built with -mno-wavefrontsize64).

PyTorch's ROCm libraries support it (marked as stable in ROCm 10.1), but gfx1010 does not have composable-kernel support, so it cannot use hipBLASLt, and is restricted to the older rocBLAS backend.

Why not extend HIP?

  • gfx1010 needs different algorithms. It has no WMMA, no dot-product instructions, no bf16 arithmetic, and is wave32. The hip kernels are built around WMMA tiles, so RDNA1 needs its own GEMMs (gemm_simt.h, the packed fp16 FMA GEMM, the implicit-GEMM conv3d) and a software tile policy for attention.
  • hip has no tier it fits. architectures.json has elementwise_only, wmma_gfx11 and wmma_gfx12. A fourth, non-WMMA tier would thread #if branches through every GEMM and attention kernel the supported cards share.
  • Those shared branches have already broken gfx1010 silently. #if GFX12 ... #else code that assumed the else branch meant gfx11 gave wrong masked int8 attention on gfx1010. CI has no RDNA1 hardware, so every HIP port could reintroduce that class of bug.
  • No regression risk for supported cards. This PR changes no file under backends/hip. rdna1 registers only when a gfx1010 is visible and declines tensors on any other device.
  • The op set differs. rdna1 drops flash decode and GatedDeltaNet decode (they need bf16) and adds gemm_f16_packed and transpose_cast. As a separate backend, the missing ops fall through to eager via the registry instead of needing per-op arch guards inside hip.
  • It is easy to drop. COMFY_KITCHEN_DISABLE_RDNA1=1 turns it off, and removing it is mostly deleting one directory and its tests, plus the small dispatch shim in backends/amd.py.

What this backend provides

  • Quantized GEMMs (fp8, int8, int4/int6, ConvRot, SVDQuant, AWQ) and the fp16 linear
    on a thread-level, register-blocked GEMM (gemm_simt.h); short M packs into 16-
    or 32-row tiles and splits K.
  • fp16_packed_linear / fp16_packed_conv3d: fp16 weights on v_pk_fma_f16, partial
    sums folded into fp32 every 32 products, activations of any range scaled by a power
    of two. The conv is an implicit GEMM with no im2col workspace. ConvRot INT8 linears
    with group size 256 run on the same packed GEMM.
  • NA3D, Sage INT8 and Sol attention on a software 16x16 tile policy (mma.h), plus an
    fp16-FMA Sage kernel for unmasked head_dim 64/128 and int8_block_sparse_attention.
  • Elementwise kernels: fp8/int8 quantize, RoPE, RMS-RoPE, AdaLN, GroupNorm+SiLU+pad.

Not provided: NA2D, flash decode and the GatedDeltaNet decode (the last two need bf16
arithmetic), and the NVFP4 and MXFP8 layouts. Those fall through to eager. Where the
fused ConvRot quantizer cannot take K, the activation is rotated eagerly and quantized
onto this backend's GEMM.

Dispatch

The backend registers as rdna1 when a gfx1010 is visible, and every op's call rule
declines tensors on any other device, so it sits beside hip in one process. Where
it registers, priority is rdna1 -> hip -> eager; Triton is left out. Calls that
bypass the registry (fp8 scaled_mm_v2, Sage attention, sol_attn_chunked) resolve
per device through comfy_kitchen/backends/amd.py. COMFY_KITCHEN_DISABLE_RDNA1=1
removes the backend.

Testing

Built with:

cd comfy-kitchen
COMFY_HIP_ARCHS="gfx1010" python setup.py build_ext --inplace
HIP_VISIBLE_DEVICES=0 python -m pytest -o addopts="" tests/test_rdna1.py \
    tests/test_fp16_packed_linear.py tests/test_fp16_conv3d.py tests/test_int8_attention.py

The INT8 references in tests/test_rdna1.py run eager with an exact fp32 matmul in
place of torch._int_mm.

Benchmark: Z-Image Turbo on upstream ComfyUI

Stock upstream ComfyUI on an RX 5600M (6 GB), 2026-10-10. A is the best configuration
the eager backend can run; B is this backend.

A: fp8, eager, sub-quadratic attention B: int8_convrot, rdna1, --use-ck-attention
Cold run (model loading included) 173.0 s 105.6 s
Warm runs (new seed) 139.5 s, 143.8 s 73.9 s, 75.5 s
Sampling 26.7-28.1 s/it 10.7-10.8 s/it
Peak VRAM 5.2 GB 3.5 GB

B samples 2.5x faster and peaks 1.7 GB lower. Both images of a seed show the same
composition.

The int8_convrot models do not run in configuration A: the first sampling step fails
in eager's int8_linear with HIPBLAS_STATUS_INVALID_VALUE from torch._int_mm
(rocblaslt error: Cannot read "TensileLibrary_lazy_gfx1010.dat"). fp8 is therefore
the eager baseline.

Where B's time goes

A B run with the GPU synchronized around each call, per 4-step image:

Op Calls Time Runs on
ConvRot INT8 linears (int8_linear) 680 33.8 s rdna1, packed fp16 GEMM
Attention (int8_attention_from_prequantized) 136 5.8 s rdna1
RMS-RoPE 136 0.3 s rdna1
Unquantized linears on the GPU 24 under 0.1 s torch
VAE decode 22.5 s CPU (--cpu-vae)

Upstream ComfyUI's own int8_linear and attention calls reach these kernels; nothing
in ComfyUI was changed. The VAE decode and, on a cold run, the text encoder run on the
CPU in both configurations, which is why the wall-clock ratio is smaller than the
sampling ratio.

Setup

  • ComfyUI 3a52aae (upstream master), no custom nodes, Python 3.14.
  • torch 2.12.0+rocm10.1.0, torchvision 0.27.0+rocm10.1.0, amd-torch-device-gfx1010
    and rocm-sdk-devel 10.1.0 from https://stable.repo.amd.com/rocm/whl-next/.
  • Both: --disable-dynamic-vram --lowvram --cpu-vae, HIP_VISIBLE_DEVICES=0. A adds
    --use-quad-cross-attention and uses comfy-kitchen 0.2.37 from PyPI, where only the
    eager backend is active on gfx1010. B adds --use-ck-attention and uses this branch
    built against the same ROCm.
  • ComfyUI's Z-Image Turbo template (cfg 1, res_multistep/simple, AuraFlow shift 3)
    at 4 steps, 1024x1024, batch 1, one fresh server per configuration.
  • A: z-image-turbo_fp8_scaled_e4m3fn_KJ.safetensors (Kijai/Z-Image_comfy_fp8_scaled)
    and qwen_3_4b_fp8_mixed.safetensors (Comfy-Org/z_image_turbo).
  • B: z_image_turbo_int8_convrot.safetensors (Comfy-Org/z_image_turbo) and
    qwen3-4b_int8_convrot_fp16emixed.safetensors (martin-rizzo/Qwen3-4B-INT8-ConvRot-ComfyUI).
  • Both: ae.safetensors (Comfy-Org/z_image_turbo).

Caveats

  • Upstream's default dynamic VRAM does not load a model on this GPU, so both
    configurations need --disable-dynamic-vram.
  • --lowvram runs the text encoder on the CPU, where this backend declines the
    tensors; the int8_convrot encoder is dequantized by eager there.
  • Both configurations keep about 2.2 GB of the diffusion model on the GPU and stream
    the other 3.7 GB from host RAM each step.
  • VRAM is device-wide mem_info_vram_used sampled at 20 Hz, which includes PyTorch's
    caching allocator; the card holds 20 MB at idle.
  • One cold and two warm runs per configuration.

🔄 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/238 **Author:** [@Zakhrov](https://github.com/Zakhrov) **Created:** 10/10/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `rdna1-backend` --- ### 📝 Commits (10+) - [`5bd4f17`](https://github.com/Comfy-Org/comfy-kitchen/commit/5bd4f1770a047f477d2440fb9882217cc4a6ca4a) Add RDNA1 backend separate from the HIP backend - [`9c485ba`](https://github.com/Comfy-Org/comfy-kitchen/commit/9c485bac09c4900ab3e83e721f1c95dd60c673a8) Enhance README.md for rdna1 backend: clarify separation from HIP and outline unique requirements - [`62d61dc`](https://github.com/Comfy-Org/comfy-kitchen/commit/62d61dcba77e55c8b992dd4009ee0c6ba6fa10bd) Refactor RDNA1 backend: Remove unneeded WMMA and FP8 structures, update GEMM and quantization kernels - [`7509aab`](https://github.com/Comfy-Org/comfy-kitchen/commit/7509aab8300ffc082bb746508b128c8d527bded0) Add test for small weights to ensure scaled rows remain in fp16 range - [`80b856c`](https://github.com/Comfy-Org/comfy-kitchen/commit/80b856cdf3a46f825de7ee9b996c637bb6ef3ddb) Add new operations for depthwise convolution, direct convolution, GEMV, and packed scaling - [`3c8754b`](https://github.com/Comfy-Org/comfy-kitchen/commit/3c8754b92b98f826b8f9150746a0aaf80d556e84) Merge branch 'main' of https://github.com/Comfy-Org/comfy-kitchen into rdna1-backend - [`19b8e11`](https://github.com/Comfy-Org/comfy-kitchen/commit/19b8e11d79ec32060da3c447563bc1659be0d9c7) Add Winograd F(2x2, 3x3) convolution implementation for fp16 in rdna1 backend - [`ba50e81`](https://github.com/Comfy-Org/comfy-kitchen/commit/ba50e8194407f6f25539e27fb62026998ef882d0) Add q4_k_linear implementation and related quantization functions for GGUF Q4_K weights - [`e946118`](https://github.com/Comfy-Org/comfy-kitchen/commit/e946118a250fdc7e4d5dbdf99862f42d3990f599) Merge branch 'main' of https://github.com/Comfy-Org/comfy-kitchen into rdna1-backend - [`a1e2f74`](https://github.com/Comfy-Org/comfy-kitchen/commit/a1e2f7450ef83b2dfdad7a15a77bf617e5f119f2) Add support for residual connections and input activation in w4a8_int8_linear ### 📊 Changes **96 files changed** (+60819 additions, -103 deletions) <details> <summary>View changed files</summary> 📝 `MANIFEST.in` (+2 -0) 📝 `README.md` (+150 -45) 📝 `comfy_kitchen/__init__.py` (+230 -4) ➕ `comfy_kitchen/backends/amd.py` (+28 -0) 📝 `comfy_kitchen/backends/eager/__init__.py` (+82 -1) 📝 `comfy_kitchen/backends/eager/conv3d.py` (+89 -0) 📝 `comfy_kitchen/backends/eager/quantization.py` (+70 -0) ➕ `comfy_kitchen/backends/rdna1/CMakeLists.txt` (+263 -0) ➕ `comfy_kitchen/backends/rdna1/README.md` (+199 -0) ➕ `comfy_kitchen/backends/rdna1/__init__.py` (+4110 -0) ➕ `comfy_kitchen/backends/rdna1/dlpack_bindings.cpp` (+3079 -0) ➕ `comfy_kitchen/backends/rdna1/embed_hsaco.cmake` (+18 -0) ➕ `comfy_kitchen/backends/rdna1/epilogue.h` (+129 -0) ➕ `comfy_kitchen/backends/rdna1/fp8_utils.h` (+183 -0) ➕ `comfy_kitchen/backends/rdna1/gemm_common.h` (+100 -0) ➕ `comfy_kitchen/backends/rdna1/gemm_simt.h` (+853 -0) ➕ `comfy_kitchen/backends/rdna1/hadamard.h` (+954 -0) ➕ `comfy_kitchen/backends/rdna1/launchers.h` (+99 -0) ➕ `comfy_kitchen/backends/rdna1/miopen_winograd/Conv_Winograd_v30_3_1_fp32_f2x3_dilation2.s` (+42 -0) ➕ `comfy_kitchen/backends/rdna1/miopen_winograd/Conv_Winograd_v30_3_1_fp32_f2x3_stride1.s` (+42 -0) _...and 76 more files_ </details> ### 📄 Description # What this PR does Adds an optional `rdna1` backend to comfy-kitchen for RDNA1 (gfx1010, e.g. RX 5600M / 5700 XT), built for that one target from its own copy of the HIP sources. It registers only for gfx1010, declines tensors on any other device, and leaves the `hip` backend and every other GPU untouched. ## Why On gfx1010 the eager backend cannot run INT8 models at all: its INT8 matmul is `torch._int_mm`, which PyTorch sends to hipBLASLt, and ROCm ships no hipBLASLt kernels for gfx1010. It also has no INT8 attention, so ComfyUI's `--use-ck-attention` refuses to start, and Triton is no alternative since it miscompiles on this target. This backend supplies those kernels itself and never calls `torch._int_mm`. ### Hardware constraints gfx1010 is a "transition" architecture between the older Vega and Navi 20 chips. It has: - **no matrix cores (WMMA was introduced with RDNA3)**: every GEMM and attention tile is computed in vector arithmetic; - **no dot-product instructions (Vega, gfx1011, gfx1012, and gfx103x aka RDNA2 have them )** (`v_dot4`/`sdot4`): int8 products go through `v_mad_i32_i24`, or through fp16 FMAs, in which int8 values are exact; - **no bf16 arithmetic**: bf16 inputs are widened; - **wave32** natively (built with `-mno-wavefrontsize64`). PyTorch's ROCm libraries support it (marked as stable in ROCm 10.1), but gfx1010 does not have composable-kernel support, so it cannot use hipBLASLt, and is restricted to the older rocBLAS backend. ### Why not extend HIP? - **gfx1010 needs different algorithms.** It has no WMMA, no dot-product instructions, no bf16 arithmetic, and is wave32. The `hip` kernels are built around WMMA tiles, so RDNA1 needs its own GEMMs (`gemm_simt.h`, the packed fp16 FMA GEMM, the implicit-GEMM conv3d) and a software tile policy for attention. - **`hip` has no tier it fits.** `architectures.json` has `elementwise_only`, `wmma_gfx11` and `wmma_gfx12`. A fourth, non-WMMA tier would thread `#if` branches through every GEMM and attention kernel the supported cards share. - **Those shared branches have already broken gfx1010 silently.** `#if GFX12 ... #else` code that assumed the else branch meant gfx11 gave wrong masked int8 attention on gfx1010. CI has no RDNA1 hardware, so every HIP port could reintroduce that class of bug. - **No regression risk for supported cards.** This PR changes no file under `backends/hip`. `rdna1` registers only when a gfx1010 is visible and declines tensors on any other device. - **The op set differs.** `rdna1` drops flash decode and GatedDeltaNet decode (they need bf16) and adds `gemm_f16_packed` and `transpose_cast`. As a separate backend, the missing ops fall through to eager via the registry instead of needing per-op arch guards inside `hip`. - **It is easy to drop.** `COMFY_KITCHEN_DISABLE_RDNA1=1` turns it off, and removing it is mostly deleting one directory and its tests, plus the small dispatch shim in `backends/amd.py`. ## What this backend provides - Quantized GEMMs (fp8, int8, int4/int6, ConvRot, SVDQuant, AWQ) and the fp16 linear on a thread-level, register-blocked GEMM (`gemm_simt.h`); short M packs into 16- or 32-row tiles and splits K. - `fp16_packed_linear` / `fp16_packed_conv3d`: fp16 weights on `v_pk_fma_f16`, partial sums folded into fp32 every 32 products, activations of any range scaled by a power of two. The conv is an implicit GEMM with no im2col workspace. ConvRot INT8 linears with group size 256 run on the same packed GEMM. - NA3D, Sage INT8 and Sol attention on a software 16x16 tile policy (`mma.h`), plus an fp16-FMA Sage kernel for unmasked head_dim 64/128 and `int8_block_sparse_attention`. - Elementwise kernels: fp8/int8 quantize, RoPE, RMS-RoPE, AdaLN, GroupNorm+SiLU+pad. Not provided: NA2D, flash decode and the GatedDeltaNet decode (the last two need bf16 arithmetic), and the NVFP4 and MXFP8 layouts. Those fall through to eager. Where the fused ConvRot quantizer cannot take K, the activation is rotated eagerly and quantized onto this backend's GEMM. ## Dispatch The backend registers as `rdna1` when a gfx1010 is visible, and every op's call rule declines tensors on any other device, so it sits beside `hip` in one process. Where it registers, priority is `rdna1 -> hip -> eager`; Triton is left out. Calls that bypass the registry (fp8 `scaled_mm_v2`, Sage attention, `sol_attn_chunked`) resolve per device through `comfy_kitchen/backends/amd.py`. `COMFY_KITCHEN_DISABLE_RDNA1=1` removes the backend. ## Testing Built with: ```bash cd comfy-kitchen COMFY_HIP_ARCHS="gfx1010" python setup.py build_ext --inplace HIP_VISIBLE_DEVICES=0 python -m pytest -o addopts="" tests/test_rdna1.py \ tests/test_fp16_packed_linear.py tests/test_fp16_conv3d.py tests/test_int8_attention.py ``` The INT8 references in `tests/test_rdna1.py` run eager with an exact fp32 matmul in place of `torch._int_mm`. ### Benchmark: Z-Image Turbo on upstream ComfyUI Stock upstream ComfyUI on an RX 5600M (6 GB), 2026-10-10. A is the best configuration the eager backend can run; B is this backend. | | A: fp8, eager, sub-quadratic attention | B: int8_convrot, rdna1, `--use-ck-attention` | | --- | --- | --- | | Cold run (model loading included) | 173.0 s | 105.6 s | | Warm runs (new seed) | 139.5 s, 143.8 s | 73.9 s, 75.5 s | | Sampling | 26.7-28.1 s/it | 10.7-10.8 s/it | | Peak VRAM | 5.2 GB | 3.5 GB | B samples 2.5x faster and peaks 1.7 GB lower. Both images of a seed show the same composition. The int8_convrot models do not run in configuration A: the first sampling step fails in eager's `int8_linear` with `HIPBLAS_STATUS_INVALID_VALUE` from `torch._int_mm` (`rocblaslt error: Cannot read "TensileLibrary_lazy_gfx1010.dat"`). fp8 is therefore the eager baseline. #### Where B's time goes A B run with the GPU synchronized around each call, per 4-step image: | Op | Calls | Time | Runs on | | --- | --- | --- | --- | | ConvRot INT8 linears (`int8_linear`) | 680 | 33.8 s | rdna1, packed fp16 GEMM | | Attention (`int8_attention_from_prequantized`) | 136 | 5.8 s | rdna1 | | RMS-RoPE | 136 | 0.3 s | rdna1 | | Unquantized linears on the GPU | 24 | under 0.1 s | torch | | VAE decode | | 22.5 s | CPU (`--cpu-vae`) | Upstream ComfyUI's own `int8_linear` and attention calls reach these kernels; nothing in ComfyUI was changed. The VAE decode and, on a cold run, the text encoder run on the CPU in both configurations, which is why the wall-clock ratio is smaller than the sampling ratio. #### Setup - ComfyUI `3a52aae` (upstream master), no custom nodes, Python 3.14. - torch 2.12.0+rocm10.1.0, torchvision 0.27.0+rocm10.1.0, `amd-torch-device-gfx1010` and `rocm-sdk-devel` 10.1.0 from `https://stable.repo.amd.com/rocm/whl-next/`. - Both: `--disable-dynamic-vram --lowvram --cpu-vae`, `HIP_VISIBLE_DEVICES=0`. A adds `--use-quad-cross-attention` and uses comfy-kitchen 0.2.37 from PyPI, where only the eager backend is active on gfx1010. B adds `--use-ck-attention` and uses this branch built against the same ROCm. - ComfyUI's Z-Image Turbo template (cfg 1, `res_multistep`/`simple`, AuraFlow shift 3) at 4 steps, 1024x1024, batch 1, one fresh server per configuration. - A: `z-image-turbo_fp8_scaled_e4m3fn_KJ.safetensors` (Kijai/Z-Image_comfy_fp8_scaled) and `qwen_3_4b_fp8_mixed.safetensors` (Comfy-Org/z_image_turbo). - B: `z_image_turbo_int8_convrot.safetensors` (Comfy-Org/z_image_turbo) and `qwen3-4b_int8_convrot_fp16emixed.safetensors` (martin-rizzo/Qwen3-4B-INT8-ConvRot-ComfyUI). - Both: `ae.safetensors` (Comfy-Org/z_image_turbo). #### Caveats - Upstream's default dynamic VRAM does not load a model on this GPU, so both configurations need `--disable-dynamic-vram`. - `--lowvram` runs the text encoder on the CPU, where this backend declines the tensors; the int8_convrot encoder is dequantized by eager there. - Both configurations keep about 2.2 GB of the diffusion model on the GPU and stream the other 3.7 GB from host RAM each step. - VRAM is device-wide `mem_info_vram_used` sampled at 20 Hz, which includes PyTorch's caching allocator; the card holds 20 MB at idle. - One cold and two warm runs per configuration. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
dev/adaptive-w3w5
dev/w3a8-gateup
dev/relayout-combine
dev/hip-residual-inplace
dev/qwen-adaptive-tree-pr
dev/h3-attn-layout-ci-ccache
dev/h3-attn-layout
dev/h3-quantizer-perf
dev/amd-resident-linear-fusions
dev/ada-unified-tree
dev/flash-decode-review
dev/qwen-decode-series
dev/bf16-delta-state
dev/int8-kv-decode
dev/qwen38-w6a8
matt/be-3580-agents-md
chore/node-runtime-refresh-2026-09-28
dev/cache-registry
feature/147-sageattention-wheels
fix/cla-api-failure-handling
docs/torch-min-version-2.7
alexisrolland-patch-2
alexisrolland-patch-1
dev/optim_int4
add-cla-workflow
dev/hip1
refactor-1
v0.2.40
v0.2.39
v0.2.38
v0.2.37
v0.2.36
v0.2.35
v0.2.34
v0.2.33
v0.2.32
v0.2.31
v0.2.30
v0.2.29
v0.2.28
v0.2.27
v0.2.26
v0.2.25
v0.2.24
v0.2.23
v0.2.22
v0.2.21
v0.2.20
v0.2.19
v0.2.18
v0.2.17
v0.2.16
v0.2.15
v0.2.14
v0.2.13
v0.2.12
v0.2.11
v0.2.10
v0.2.9
v0.2.8
v0.2.7
v0.2.6
v0.2.5
v0.2.4
v0.2.3
v0.2.2
v0.2.1
v0.2.0
v0.1.6
v0.1.5
v0.1.4
v0.1.3
v0.1.2
v0.1.1
v0.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
Comfy-Org/comfy-kitchen#250
No description provided.