Comfy-Org/comfy-kitchen

Fork 0

mirror of https://github.com/Comfy-Org/comfy-kitchen.git synced 2026-11-05 07:47:40 +00:00

[GH-ISSUE #92] eager backend advertises int8_linear on MPS but dispatches to CUDA-only torch._int_mm #29

New issue

Open

opened 2026-10-26 22:19:22 +00:00 by bsakaguchi · 7 comments

bsakaguchi commented

2026-10-26 22:19:22 +00:00

Owner

Originally created by @namikazi25 on GitHub (Aug 3, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/92

Describe the bug

int8 quantized checkpoints fail immediately on Apple Silicon. The eager backend lists int8_linear under native ops on MPS, but fast_int8_mm calls torch._int_mm, which has no MPS kernel. It dies at the first linear layer of step 0.

Interestingly, the nvfp4 text encoder in the same workflow loads and runs fine. It's classified as emulated. So the emulation path works on MPS. The problem looks like just the native/emulated classification being wrong for int8 on this device.

Reproduce

Files (from the MiniMax H3 template):

diffusion_models/minimax_h3_fl2va_pruned_int8_convrot.safetensors
text_encoders/qwen3vl_32b_minimax_h3_nvfp4_awq.safetensors

Load and run any sampler. Fails at 0/20.

Startup shows:

Found comfy_kitchen backend cuda: available=False (Extension file not found: .../backends/cuda/_C.abi3.so)
Found comfy_kitchen backend triton: available=False (No module named 'triton')
Found comfy_kitchen backend hip: available=False (PyTorch ROCm/HIP runtime not available)
Found comfy_kitchen backend eager: available=True

Model load shows:

Detected mixed precision quantization
Native ops: int8_tensorwise, convrot_w4a4 , emulated ops: float8_e4m3fn, mxfp8, nvfp4, float8_e5m2

Traceback

comfy/ops.py:1313 in _forward
    return torch.nn.functional.linear(input, weight, bias)
comfy_kitchen/tensor/base.py:362 in __torch_dispatch__
comfy_kitchen/tensor/int8.py:287 in _handle_int8_linear_tensorwise
comfy_kitchen/backends/eager/quantization.py:1222 in _op_int8_linear
comfy_kitchen/backends/eager/quantization.py:1022 in int8_linear
    result = _int8_matmul_accumulate(x_8, weight.T.contiguous())
comfy_kitchen/backends/eager/quantization.py:779 in _int8_matmul_accumulate
    result = fast_int8_mm(a, b)
comfy_kitchen/backends/eager/quantization.py:754 in fast_int8_mm
    return torch._int_mm(lhs, rhs)

NotImplementedError: The operator 'aten::_int_mm' is not currently
implemented for the MPS device.

Upstream op tracker: [pytorch/pytorch#141287]

Environment

macOS 26.1, Apple Silicon, 48GB unified
torch 2.10.0, Python 3.12.12
comfy-kitchen 0.2.26, ComfyUI 0.30.0
Device mps, vram state SHARED

Notes

Not a regression. [#9](https://github.com/Comfy-Org/comfy-kitchen/issues/9) scoped the int8 path to Nvidia 8.0-8.9 with torch._int_mm as the matmul, so MPS was never considered. Just a gap.

The README says out-of-domain requests fall back to torch/eager, and describes HIP degrading on RDNA2 by not advertising GEMMs it can't run. Same idea seems like it should apply here. Eager is the fallback tier, so it shouldn't advertise something it can't service on the current device.

Two ways I can see to fix it:

1. Don't advertise int8_linear on non-CUDA devices, so dispatch falls through to the normal torch path. Probably the cleaner one given the HIP precedent.

2. Fall back inside fast_int8_mm:

def fast_int8_mm(lhs, rhs):
    if lhs.device.type != "cuda":
        return torch.matmul(
            lhs.to(torch.float32), rhs.to(torch.float32)
        ).to(torch.int32)
    return torch._int_mm(lhs, rhs)

fp32 holds integers exactly to 2^24, so this should be bit-identical to _int_mm for realistic K with int8 inputs. Loses the int8 speedup but stays on GPU. (PYTORCH_ENABLE_MPS_FALLBACK=1 "works" but round-trips to CPU on every linear in every block on every step, so it's not really usable.)

Happy to open a PR for either. Let me know which you'd prefer.

Originally created by @namikazi25 on GitHub (Aug 3, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/92 ### Describe the bug int8 quantized checkpoints fail immediately on Apple Silicon. The eager backend lists `int8_linear` under native ops on MPS, but `fast_int8_mm` calls `torch._int_mm`, which has no MPS kernel. It dies at the first linear layer of step 0. Interestingly, the nvfp4 text encoder in the same workflow loads and runs fine. It's classified as emulated. So the emulation path works on MPS. The problem looks like just the native/emulated classification being wrong for int8 on this device. ### Reproduce Files (from the MiniMax H3 template): - `diffusion_models/minimax_h3_fl2va_pruned_int8_convrot.safetensors` - `text_encoders/qwen3vl_32b_minimax_h3_nvfp4_awq.safetensors` Load and run any sampler. Fails at 0/20. Startup shows: ``` Found comfy_kitchen backend cuda: available=False (Extension file not found: .../backends/cuda/_C.abi3.so) Found comfy_kitchen backend triton: available=False (No module named 'triton') Found comfy_kitchen backend hip: available=False (PyTorch ROCm/HIP runtime not available) Found comfy_kitchen backend eager: available=True ``` Model load shows: ``` Detected mixed precision quantization Native ops: int8_tensorwise, convrot_w4a4 , emulated ops: float8_e4m3fn, mxfp8, nvfp4, float8_e5m2 ``` ### Traceback ``` comfy/ops.py:1313 in _forward return torch.nn.functional.linear(input, weight, bias) comfy_kitchen/tensor/base.py:362 in __torch_dispatch__ comfy_kitchen/tensor/int8.py:287 in _handle_int8_linear_tensorwise comfy_kitchen/backends/eager/quantization.py:1222 in _op_int8_linear comfy_kitchen/backends/eager/quantization.py:1022 in int8_linear result = _int8_matmul_accumulate(x_8, weight.T.contiguous()) comfy_kitchen/backends/eager/quantization.py:779 in _int8_matmul_accumulate result = fast_int8_mm(a, b) comfy_kitchen/backends/eager/quantization.py:754 in fast_int8_mm return torch._int_mm(lhs, rhs) NotImplementedError: The operator 'aten::_int_mm' is not currently implemented for the MPS device. ``` Upstream op tracker: [[pytorch/pytorch#141287](https://github.com/pytorch/pytorch/issues/141287)] ### Environment - macOS 26.1, Apple Silicon, 48GB unified - torch 2.10.0, Python 3.12.12 - comfy-kitchen 0.2.26, ComfyUI 0.30.0 - Device mps, vram state SHARED ### Notes Not a regression. [[#9](https://github.com/Comfy-Org/comfy-kitchen/issues/9)](https://github.com/Comfy-Org/comfy-kitchen/issues/9) scoped the int8 path to Nvidia 8.0-8.9 with `torch._int_mm` as the matmul, so MPS was never considered. Just a gap. The README says out-of-domain requests fall back to torch/eager, and describes HIP degrading on RDNA2 by not advertising GEMMs it can't run. Same idea seems like it should apply here. Eager is the fallback tier, so it shouldn't advertise something it can't service on the current device. Two ways I can see to fix it: **1. Don't advertise `int8_linear` on non-CUDA devices**, so dispatch falls through to the normal torch path. Probably the cleaner one given the HIP precedent. **2. Fall back inside `fast_int8_mm`:** ```python def fast_int8_mm(lhs, rhs): if lhs.device.type != "cuda": return torch.matmul( lhs.to(torch.float32), rhs.to(torch.float32) ).to(torch.int32) return torch._int_mm(lhs, rhs) ``` fp32 holds integers exactly to 2^24, so this should be bit-identical to `_int_mm` for realistic K with int8 inputs. Loses the int8 speedup but stays on GPU. (`PYTORCH_ENABLE_MPS_FALLBACK=1` "works" but round-trips to CPU on every linear in every block on every step, so it's not really usable.) Happy to open a PR for either. Let me know which you'd prefer.

bsakaguchi commented

2026-10-26 22:19:23 +00:00

Author

Owner

@coderabbitai[bot] commented on GitHub (Aug 3, 2026):

🔗 Related PRs

Comfy-Org/comfy-kitchen#90 - Add optimized w4a8 with int8 codebook [open]

🧪 Issue enrichment is currently in open beta.

You can configure auto-planning by selecting labels in the issue_enrichment configuration.

To disable automatic issue enrichment, add the following to your .coderabbit.yaml:

issue_enrichment:
  auto_enrich:
    enabled: false

💬 Have feedback or questions? Drop into our discord!

@coderabbitai[bot] commented on GitHub (Aug 3, 2026):  <details> <summary>🔗 Related PRs</summary> Comfy-Org/comfy-kitchen#90 - Add optimized w4a8 with int8 codebook [open] </details> --- <details> <summary> 🧪 Issue enrichment is currently in open beta.</summary> You can configure auto-planning by selecting labels in the issue_enrichment configuration. To disable automatic issue enrichment, add the following to your `.coderabbit.yaml`: ```yaml issue_enrichment: auto_enrich: enabled: false ``` </details> 💬 Have feedback or questions? Drop into our [discord](https://discord.gg/coderabbit)!

bsakaguchi commented

2026-10-26 22:19:23 +00:00

Author

Owner

@namikazi25 commented on GitHub (Aug 3, 2026):

Confirmed the overlap with #90. Its eager backend
(comfy_kitchen/backends/eager/w4a8_int8.py) describes itself as a
"portable fallback for any torch device (CPU/CUDA/ROCm)" and ends in
return int8_linear(...), so it reaches fast_int8_mm and hits the same
_int_mm wall on MPS. Worth handling together rather than twice.

That also shifts my preference toward option 2. The int4 to int8 dequant
in that file is pure torch and runs fine on MPS. Only the terminal
int8_linear fails, so a single fallback in fast_int8_mm unblocks both
the existing int8 path and the new w4a8 one. Option 1 would need the
capability mask applied in two places.

#90 does validate devices and compute capability when registering
backend ops, so either direction has precedent in the codebase already.

@namikazi25 commented on GitHub (Aug 3, 2026): Confirmed the overlap with #90. Its eager backend (`comfy_kitchen/backends/eager/w4a8_int8.py`) describes itself as a "portable fallback for any torch device (CPU/CUDA/ROCm)" and ends in `return int8_linear(...)`, so it reaches `fast_int8_mm` and hits the same `_int_mm` wall on MPS. Worth handling together rather than twice. That also shifts my preference toward option 2. The int4 to int8 dequant in that file is pure torch and runs fine on MPS. Only the terminal `int8_linear` fails, so a single fallback in `fast_int8_mm` unblocks both the existing int8 path and the new w4a8 one. Option 1 would need the capability mask applied in two places. #90 does validate devices and compute capability when registering backend ops, so either direction has precedent in the codebase already.

bsakaguchi commented

2026-10-26 22:19:23 +00:00

Author

Owner

@AstroPig7 commented on GitHub (Aug 4, 2026):

@namikazi25 Did you test the fallback in fast_int8_mm for speed? I modified comfy_kitchen’s quantization routine with your patch and found it to be very, very slow.

@AstroPig7 commented on GitHub (Aug 4, 2026): @namikazi25 Did you test the fallback in **`fast_int8_mm`** for speed? I modified comfy_kitchen’s quantization routine with your patch and found it to be very, very slow.

bsakaguchi commented

2026-10-26 22:19:23 +00:00

Author

Owner

@Terralien-jp commented on GitHub (Aug 5, 2026):

@AstroPig7 I hit the same thing, and I think I can explain the slowness and offer a faster shape for the fix.

Why a `fast_int8_mm` fallback is slow

A fallback there has to keep producing the int32 accumulator, so the only exact option on MPS is an fp32 matmul. MPS fp32 matmul is ~4x slower than bf16, and that dominates everything else. Measured on M5 Pro / 64GB, torch 2.12.0, at a shape the H3 DiT actually uses (M=10032 tokens, K=3072):

op	N=3072	N=12288
bf16 matmul	6.98 ms	28.72 ms
fp32 matmul (operands already cast)	28.38 ms	—
int8 fallback via fp32, K-chunked for exactness	34.67 ms	147.65 ms

The int8→float cast itself is negligible (0.82 ms). Chunking K to keep the accumulation inside fp32's 24-bit mantissa adds ~4.5 ms. The rest is just fp32 being slow on this backend. So no amount of tuning inside fast_int8_mm gets close to bf16 — the integer domain is the problem, not the implementation.

What worked instead

Handle it one level up, in int8_linear: dequantize the weight to the activation dtype and run a plain F.linear. The weight stays int8 in memory (only the per-call operand is widened), and the activation no longer round-trips through int8 quantization, so this is more accurate than the int8 path, not less. ComfyUI already drives these models at bfloat16 (manual cast: torch.bfloat16), so bf16 is the target precision anyway — the int32 exactness buys nothing.

def _mps_int8_linear(x, weight, weight_scale, bias, out_dtype,
                     convrot=False, convrot_groupsize=256, input_act=None):
    x = _apply_input_act(x, input_act)
    weight = weight.to(device=x.device)
    scale = weight_scale.to(device=x.device, dtype=torch.float32).reshape(-1)
    if convrot:
        h = _build_hadamard(convrot_groupsize, device=x.device, dtype=x.dtype)
        x = _rotate_activation(x, h, convrot_groupsize)
    w = weight.to(x.dtype)
    w = w * (scale.to(x.dtype) if scale.numel() == 1 else scale.to(x.dtype).reshape(-1, 1))
    return torch.nn.functional.linear(
        x, w, None if bias is None else bias.to(x.dtype)).to(out_dtype)

The convrot Hadamard rotation still has to be applied to the activation exactly as the int8 path does it, since the stored weight is pre-rotated.

Accuracy: relative difference vs the CPU int8_linear reference is 0.0085 with convrot on and 0.0086 with it off. That is the reference's own per-row activation quantization error (1/127 ≈ 0.0079), i.e. the difference is the int8 path's error, not this path's.

End-to-end, MiniMax H3

minimax_h3_fl2va_pruned_int8_convrot + qwen3vl_32b_minimax_h3_int8_convrot, same seed and settings:

	fp32 fallback in `fast_int8_mm`	dequant-to-bf16 in `int8_linear`
608x352, 39f, 4 steps	175.0 s (24.8 s/it)	84.5 s (7.4 s/it)
640x640, 39f, 20 steps	22m43s (60.3 s/it)	8m27s (22.0 s/it)

3.4x on sampling. Output is visually indistinguishable at the same seed.

Since this sits in int8_linear rather than fast_int8_mm, it also covers the w4a8 path from #90 that @namikazi25 mentioned — that one ends in int8_linear too.

Two notes for whoever picks this up:

--gpu-only is not a workaround for the related text-encoder placement issue. text_encoder_device() never returns the GPU on Apple Silicon because vram_state is SHARED, and forcing it with --gpu-only put the 26GB TE on the GPU next to the 20GB DiT: swap hit 28GB and the run got slower (217 s vs 175 s).
MPS cannot hold the fp8 dtype at all (Trying to convert Float8_e4m3fn to the MPS backend but it does not have support for that dtype), so the fp8_scaled checkpoints are not an escape hatch either.

Environment: M5 Pro, 64GB unified, macOS 26.5, torch 2.12.0, comfy-kitchen 0.2.26, ComfyUI 0.30.0, eager backend only.

@Terralien-jp commented on GitHub (Aug 5, 2026): @AstroPig7 I hit the same thing, and I think I can explain the slowness and offer a faster shape for the fix. ## Why a `fast_int8_mm` fallback is slow A fallback there has to keep producing the int32 accumulator, so the only exact option on MPS is an fp32 matmul. **MPS fp32 matmul is ~4x slower than bf16**, and that dominates everything else. Measured on M5 Pro / 64GB, torch 2.12.0, at a shape the H3 DiT actually uses (M=10032 tokens, K=3072): | op | N=3072 | N=12288 | |---|---|---| | bf16 matmul | 6.98 ms | 28.72 ms | | fp32 matmul (operands already cast) | 28.38 ms | — | | int8 fallback via fp32, K-chunked for exactness | 34.67 ms | 147.65 ms | The int8→float cast itself is negligible (0.82 ms). Chunking K to keep the accumulation inside fp32's 24-bit mantissa adds ~4.5 ms. **The rest is just fp32 being slow on this backend.** So no amount of tuning inside `fast_int8_mm` gets close to bf16 — the integer domain is the problem, not the implementation. ## What worked instead Handle it one level up, in `int8_linear`: dequantize the weight to the activation dtype and run a plain `F.linear`. The weight stays int8 in memory (only the per-call operand is widened), and the activation no longer round-trips through int8 quantization, so this is **more** accurate than the int8 path, not less. ComfyUI already drives these models at bfloat16 (`manual cast: torch.bfloat16`), so bf16 is the target precision anyway — the int32 exactness buys nothing. ```python def _mps_int8_linear(x, weight, weight_scale, bias, out_dtype, convrot=False, convrot_groupsize=256, input_act=None): x = _apply_input_act(x, input_act) weight = weight.to(device=x.device) scale = weight_scale.to(device=x.device, dtype=torch.float32).reshape(-1) if convrot: h = _build_hadamard(convrot_groupsize, device=x.device, dtype=x.dtype) x = _rotate_activation(x, h, convrot_groupsize) w = weight.to(x.dtype) w = w * (scale.to(x.dtype) if scale.numel() == 1 else scale.to(x.dtype).reshape(-1, 1)) return torch.nn.functional.linear( x, w, None if bias is None else bias.to(x.dtype)).to(out_dtype) ``` The convrot Hadamard rotation still has to be applied to the activation exactly as the int8 path does it, since the stored weight is pre-rotated. **Accuracy**: relative difference vs the CPU `int8_linear` reference is 0.0085 with convrot on and 0.0086 with it off. That is the reference's own per-row activation quantization error (1/127 ≈ 0.0079), i.e. the difference is the int8 path's error, not this path's. ## End-to-end, MiniMax H3 `minimax_h3_fl2va_pruned_int8_convrot` + `qwen3vl_32b_minimax_h3_int8_convrot`, same seed and settings: | | fp32 fallback in `fast_int8_mm` | dequant-to-bf16 in `int8_linear` | |---|---|---| | 608x352, 39f, 4 steps | 175.0 s (24.8 s/it) | **84.5 s (7.4 s/it)** | | 640x640, 39f, 20 steps | 22m43s (60.3 s/it) | **8m27s (22.0 s/it)** | **3.4x on sampling.** Output is visually indistinguishable at the same seed. Since this sits in `int8_linear` rather than `fast_int8_mm`, it also covers the w4a8 path from #90 that @namikazi25 mentioned — that one ends in `int8_linear` too. Two notes for whoever picks this up: - `--gpu-only` is not a workaround for the related text-encoder placement issue. `text_encoder_device()` never returns the GPU on Apple Silicon because `vram_state` is `SHARED`, and forcing it with `--gpu-only` put the 26GB TE on the GPU next to the 20GB DiT: swap hit 28GB and the run got *slower* (217 s vs 175 s). - MPS cannot hold the fp8 dtype at all (`Trying to convert Float8_e4m3fn to the MPS backend but it does not have support for that dtype`), so the fp8_scaled checkpoints are not an escape hatch either. Environment: M5 Pro, 64GB unified, macOS 26.5, torch 2.12.0, comfy-kitchen 0.2.26, ComfyUI 0.30.0, eager backend only.

bsakaguchi commented

2026-10-26 22:19:23 +00:00

Author

Owner

@namikazi25 commented on GitHub (Aug 5, 2026):

Confirmed working on a second config. macOS 26.1, Apple Silicon 48GB, torch 2.10.0, comfy-kitchen 0.2.26, eager backend only.

I applied the dequant path as a runtime monkeypatch on a custom node rather than editing site-packages, so a pip upgrade doesn't quietly wipe it. fast_int8_mm left alone.

608x352, 39 frames, 20 steps, res_multistep. Sampling ran 6m56s at 20.84 s/it, 525s total with VAE decode. Output looks right. Same DiT and nvfp4 TE from the H3 template.

Worth flagging that's roughly 2.8x slower than @Terralien-jp's 7.4 s/it at the same resolution and frame count. Probably just hardware and torch version, since they're on M5 Pro with torch 2.12 and I'm on torch 2.10. Their 640x640 number tracks about the same ratio, so it looks like a uniform gap rather than anything specific to this path. Mentioning it in case someone else on older silicon sees the same and it turns out to be fixable.

One thing that cost me time and isn't obvious from the source: registry.get_implementation resolves via getattr on the backend module object, so if you only patch backends.eager.quantization.int8_linear it gets silently bypassed and you sit there wondering why nothing changed. You need backends.eager.int8_linear too.

Two small deviations I'd suggest from the snippet as posted, both just matching what the CUDA path already does. Keep the K-mismatch and scale-shape checks. And run _apply_input_act before the MPS branch, since swiglu halves K and skipping it breaks the dimension check.

Also w = weight.to(x.dtype) then w = w * scale allocates two [N,K] bf16 tensors. mul_ in place drops the transient from 36 MiB to 18 at K=N=3072, and 144 to 72 at N=12288. Free, no behaviour change.

Accuracy against the CPU int8_linear reference at M=10032, K=3072, N=3072: 0.00966 rel Frobenius with convrot off, 0.00964 with it on, max abs 0.035. Close to @Terralien-jp's 0.0085. The two convrot cases differing by 2e-5 is what convinced me the rotation matches the existing path properly.

I also went down a dead end that's probably worth recording so nobody else repeats it. torch._weight_int8pack_mm exists and runs on MPS in torch 2.10, which looks like it should give a proper fused native path. It doesn't help. K=3072, bf16, mean of 20 after 5 warmup:

M	N	cast+F.linear	int8pack_mm	speedup
10032	3072	13.52 ms	13.29 ms	1.02x
10032	12288	54.29 ms	53.91 ms	1.01x
1	3072	0.415 ms	0.020 ms	21.2x
1	12288	1.667 ms	0.073 ms	22.7x

It's a decode kernel. At M=1 the cast is the whole cost because there's no GEMM to hide it behind. At diffusion M you're bandwidth-bound on the [M,N] output and the [N,K] weight read anyway, so the cast just amortises away. Only argument left for it is memory, since it kills the transient entirely. But it's a private ATen op with no version guarantee and a 0.0026 rel diff from the cast path, so I don't think it's worth it unless someone actually hits OOM.

I should admit I spent a while assuming the fix belonged in fast_int8_mm before it clicked that by the time you're down there you're already committed to returning an int32 accumulator, which forces fp32, which is the slow thing. Going up a level is the whole trick.

Two other notes. The bug survives a ComfyUI update, aimdo 0.4.11 to 0.4.13 and frontend 1.47 to 1.48, comfy-kitchen still 0.2.26.

And the eager backend already advertises dequantize_int8_convrot_weight and dequantize_int8_simple_dtype on MPS. So the primitive this needs is sitting right there in the same backend, working. It just never gets reached from int8_linear.

</html>

@namikazi25 commented on GitHub (Aug 5, 2026): <html><head></head><body>Confirmed working on a second config. macOS 26.1, Apple Silicon 48GB, torch 2.10.0, comfy-kitchen 0.2.26, eager backend only. I applied the dequant path as a runtime monkeypatch on a custom node rather than editing site-packages, so a pip upgrade doesn't quietly wipe it. <code>fast_int8_mm</code> left alone. 608x352, 39 frames, 20 steps, res_multistep. Sampling ran 6m56s at 20.84 s/it, 525s total with VAE decode. Output looks right. Same DiT and nvfp4 TE from the H3 template. Worth flagging that's roughly 2.8x slower than @Terralien-jp's 7.4 s/it at the same resolution and frame count. Probably just hardware and torch version, since they're on M5 Pro with torch 2.12 and I'm on torch 2.10. Their 640x640 number tracks about the same ratio, so it looks like a uniform gap rather than anything specific to this path. Mentioning it in case someone else on older silicon sees the same and it turns out to be fixable. One thing that cost me time and isn't obvious from the source: <code>registry.get_implementation</code> resolves via <code>getattr</code> on the backend module object, so if you only patch <code>backends.eager.quantization.int8_linear</code> it gets silently bypassed and you sit there wondering why nothing changed. You need <code>backends.eager.int8_linear</code> too. Two small deviations I'd suggest from the snippet as posted, both just matching what the CUDA path already does. Keep the K-mismatch and scale-shape checks. And run <code>_apply_input_act</code> before the MPS branch, since swiglu halves K and skipping it breaks the dimension check. Also <code>w = weight.to(x.dtype)</code> then <code>w = w * scale</code> allocates two [N,K] bf16 tensors. <code>mul_</code> in place drops the transient from 36 MiB to 18 at K=N=3072, and 144 to 72 at N=12288. Free, no behaviour change. Accuracy against the CPU <code>int8_linear</code> reference at M=10032, K=3072, N=3072: 0.00966 rel Frobenius with convrot off, 0.00964 with it on, max abs 0.035. Close to @Terralien-jp's 0.0085. The two convrot cases differing by 2e-5 is what convinced me the rotation matches the existing path properly. I also went down a dead end that's probably worth recording so nobody else repeats it. <code>torch._weight_int8pack_mm</code> exists and runs on MPS in torch 2.10, which looks like it should give a proper fused native path. It doesn't help. K=3072, bf16, mean of 20 after 5 warmup: M | N | cast+F.linear | int8pack_mm | speedup -- | -- | -- | -- | -- 10032 | 3072 | 13.52 ms | 13.29 ms | 1.02x 10032 | 12288 | 54.29 ms | 53.91 ms | 1.01x 1 | 3072 | 0.415 ms | 0.020 ms | 21.2x 1 | 12288 | 1.667 ms | 0.073 ms | 22.7x It's a decode kernel. At M=1 the cast is the whole cost because there's no GEMM to hide it behind. At diffusion M you're bandwidth-bound on the [M,N] output and the [N,K] weight read anyway, so the cast just amortises away. Only argument left for it is memory, since it kills the transient entirely. But it's a private ATen op with no version guarantee and a 0.0026 rel diff from the cast path, so I don't think it's worth it unless someone actually hits OOM. I should admit I spent a while assuming the fix belonged in <code>fast_int8_mm</code> before it clicked that by the time you're down there you're already committed to returning an int32 accumulator, which forces fp32, which is the slow thing. Going up a level is the whole trick. Two other notes. The bug survives a ComfyUI update, aimdo 0.4.11 to 0.4.13 and frontend 1.47 to 1.48, comfy-kitchen still 0.2.26. And the eager backend already advertises <code>dequantize_int8_convrot_weight</code> and <code>dequantize_int8_simple_dtype</code> on MPS. So the primitive this needs is sitting right there in the same backend, working. It just never gets reached from <code>int8_linear</code>.</body></html>

bsakaguchi commented

2026-10-26 22:19:24 +00:00

Author

Owner

@MrLavender commented on GitHub (Aug 7, 2026):

Similar problem with AMD RDNA2 (gfx1030).

The nvfp4 text encoder step works fine but it crashes on the next step because torch._int_mm tries to use hipBLASLt which isn't available on RDNA2.

Note that I have set TORCH_BLAS_PREFER_HIPBLASLT=0.

Stack trace

rocblaslt error: Cannot read "/venv/lib/python3.12/site-packages/torch/lib/hipblaslt/library/TensileLibrary_lazy_gfx1030.dat": No such file or directory

rocblaslt error: Could not load "/venv/lib/python3.12/site-packages/torch/lib/hipblaslt/library/TensileLibrary_lazy_gfx1030.dat"
  0%|                                                                       | 0/20 [00:01<?, ?it/s]
[ERROR] !!! Exception during processing !!! CUDA error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblasLtMatmulAlgoGetHeuristic( ltHandle, computeDesc.descriptor(), Adesc.descriptor(), Bdesc.descriptor(), Cdesc.descriptor(), Cdesc.descriptor(), preference.descriptor(), 1, &heuristicResult, &returnedResult)`
[ERROR] Traceback (most recent call last):
  File "/comfyui/execution.py", line 545, in execute
    output_data, output_ui, has_subgraph, has_pending_tasks = await get_output_data(prompt_id, unique_id, obj, input_data_all, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data)
                                                              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/execution.py", line 344, in get_output_data
    return_values = await _async_map_node_over_list(prompt_id, unique_id, obj, input_data_all, obj.FUNCTION, allow_interrupt=True, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/execution.py", line 318, in _async_map_node_over_list
    await process_inputs(input_dict, i)
  File "/comfyui/execution.py", line 306, in process_inputs
    result = f(**inputs)
             ^^^^^^^^^^^
  File "/comfyui/comfy_api/internal/__init__.py", line 149, in wrapped_func
    return method(locked_class, **inputs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy_api/latest/_io.py", line 1935, in EXECUTE_NORMALIZED
    to_return = cls.execute(*args, **kwargs)
                ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy_extras/nodes_custom_sampler.py", line 1049, in execute
    samples = guider.sample(noise.generate_noise(latent), latent_image, sampler, sigmas, denoise_mask=noise_mask, callback=callback, disable_pbar=disable_pbar, seed=noise.seed)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1333, in sample
    output = executor.execute(noise, latent_image, sampler, sigmas, denoise_mask, callback, disable_pbar, seed, latent_shapes=latent_shapes)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1260, in outer_sample
    output = self.inner_sample(noise, latent_image, device, sampler, sigmas, denoise_mask, callback, disable_pbar, seed, latent_shapes=latent_shapes)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1235, in inner_sample
    samples = executor.execute(self, sigmas, extra_args, callback, noise, latent_image, denoise_mask, disable_pbar)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1005, in sample
    samples = self.sampler_function(model_k, noise, sigmas, extra_args=extra_args, callback=k_callback, disable=disable_pbar, **self.extra_options)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 124, in decorate_context
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/k_diffusion/sampling.py", line 1460, in sample_res_multistep
    return res_multistep(model, x, sigmas, extra_args=extra_args, callback=callback, disable=disable, s_noise=s_noise, noise_sampler=noise_sampler, eta=0., cfg_pp=False)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 124, in decorate_context
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/k_diffusion/sampling.py", line 1418, in res_multistep
    denoised = model(x, sigmas[i] * s_in, **extra_args)
               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 640, in __call__
    out = self.inner_model(x, sigma, model_options=model_options, seed=seed)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1208, in __call__
    return self.outer_predict_noise(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1215, in outer_predict_noise
    ).execute(x, timestep, model_options, seed)
      ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 1218, in predict_noise
    return sampling_function(self.inner_model, x, timestep, self.conds.get("negative", None), self.conds.get("positive", None), self.cfg, model_options=model_options, seed=seed)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 620, in sampling_function
    out = calc_cond_batch(model, conds, x, timestep, model_options)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 211, in calc_cond_batch
    return _calc_cond_batch_outer(model, conds, x_in, timestep, model_options)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 219, in _calc_cond_batch_outer
    return executor.execute(model, conds, x_in, timestep, model_options)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/samplers.py", line 335, in _calc_cond_batch
    output = model.apply_model(input_x, timestep_, **c).chunk(batch_chunks)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/model_base.py", line 196, in apply_model
    return comfy.patcher_extension.WrapperExecutor.new_class_executor(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/model_base.py", line 240, in _apply_model
    model_output = self.diffusion_model(xc, t, context=context, control=control, transformer_options=transformer_options, **extra_conds)
                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ldm/minimax/model.py", line 499, in forward
    return comfy.patcher_extension.WrapperExecutor.new_class_executor(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/patcher_extension.py", line 113, in execute
    return self.original(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ldm/minimax/model.py", line 629, in _forward
    h = block(h, t_emb, mod_segments, rope_freqs, transformer_options=transformer_options)
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ldm/minimax/model.py", line 270, in forward
    x = _mod_gate(x, gate_msa, self.attn(h, rope_freqs=rope_freqs, transformer_options=transformer_options), mod_segments)
                               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ldm/minimax/model.py", line 158, in forward
    q, k, v = self.qkv_proj(x).split(self.heads * self.head_dim, dim=-1)
              ^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ops.py", line 1405, in forward
    output = self.forward_comfy_cast_weights(
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ops.py", line 1342, in forward_comfy_cast_weights
    x = self._forward(input, weight, bias)
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/comfy/ops.py", line 1313, in _forward
    return torch.nn.functional.linear(input, weight, bias)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/tensor/base.py", line 362, in __torch_dispatch__
    return op_handlers[parent_cls](qt, args, kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/tensor/int8.py", line 287, in _handle_int8_linear_tensorwise
    return torch.ops.comfy_kitchen.int8_linear(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/_ops.py", line 1269, in __call__
    return self._op(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/_library/custom_ops.py", line 347, in backend_impl
    result = self._backend_fns[device_type](*args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/_compile.py", line 54, in inner
    return disable_fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/_dynamo/eval_frame.py", line 1263, in _fn
    return fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^
  File "/venv/lib/python3.12/site-packages/torch/_library/custom_ops.py", line 382, in wrapped_fn
    return fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 1222, in _op_int8_linear
    return impl(**kwargs)
           ^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 1022, in int8_linear
    result = _int8_matmul_accumulate(x_8, weight.T.contiguous())
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 779, in _int8_matmul_accumulate
    result = fast_int8_mm(a, b)
             ^^^^^^^^^^^^^^^^^^
  File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 754, in fast_int8_mm
    return torch._int_mm(lhs, rhs)
           ^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: CUDA error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblasLtMatmulAlgoGetHeuristic( ltHandle, computeDesc.descriptor(), Adesc.descriptor(), Bdesc.descriptor(), Cdesc.descriptor(), Cdesc.descriptor(), preference.descriptor(), 1, &heuristicResult, &returnedResult)`

EDIT: Tested both of the solutions posted above and they work (obviously I had to remove the device.type != "cuda" check because AMD gpus are cuda in PyTorch). The int8_linear fix is only about 7-8% faster on my RX 6800 16GB.

Using the default T2V example that comes with ComfyUI;

fast_int8_mm    37:37<00:00, 112.87s/it
int8_linear     35:09<00:00, 105.48s/it

EDIT2: Found https://github.com/patientx/ComfyUI-INT8-Fast-ROCM which has a faster implementation for RDNA2 using custom triton kernels. Still, it's only 17-18% faster than int8_linear on the default T2V template @ 0.2 megapixels;

int8_linear     20/20 [20:50<00:00, 62.53s/it]
INT8-Fast-ROCM  20/20 [17:12<00:00, 51.64s/it]

And I guess that's as about good as I'm going get on this hardware without LoRA/cache/other speedup hacks. Tbh I'm still amazed that I can run this SOTA model at all!

@MrLavender commented on GitHub (Aug 7, 2026): Similar problem with AMD RDNA2 (gfx1030). The nvfp4 text encoder step works fine but it crashes on the next step because `torch._int_mm` tries to use hipBLASLt which isn't available on RDNA2. Note that I have set `TORCH_BLAS_PREFER_HIPBLASLT=0`. <details><summary>Stack trace</summary> ``` rocblaslt error: Cannot read "/venv/lib/python3.12/site-packages/torch/lib/hipblaslt/library/TensileLibrary_lazy_gfx1030.dat": No such file or directory rocblaslt error: Could not load "/venv/lib/python3.12/site-packages/torch/lib/hipblaslt/library/TensileLibrary_lazy_gfx1030.dat" 0%| | 0/20 [00:01<?, ?it/s] [ERROR] !!! Exception during processing !!! CUDA error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblasLtMatmulAlgoGetHeuristic( ltHandle, computeDesc.descriptor(), Adesc.descriptor(), Bdesc.descriptor(), Cdesc.descriptor(), Cdesc.descriptor(), preference.descriptor(), 1, &heuristicResult, &returnedResult)` [ERROR] Traceback (most recent call last): File "/comfyui/execution.py", line 545, in execute output_data, output_ui, has_subgraph, has_pending_tasks = await get_output_data(prompt_id, unique_id, obj, input_data_all, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/execution.py", line 344, in get_output_data return_values = await _async_map_node_over_list(prompt_id, unique_id, obj, input_data_all, obj.FUNCTION, allow_interrupt=True, execution_block_cb=execution_block_cb, pre_execute_cb=pre_execute_cb, v3_data=v3_data) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/execution.py", line 318, in _async_map_node_over_list await process_inputs(input_dict, i) File "/comfyui/execution.py", line 306, in process_inputs result = f(**inputs) ^^^^^^^^^^^ File "/comfyui/comfy_api/internal/__init__.py", line 149, in wrapped_func return method(locked_class, **inputs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy_api/latest/_io.py", line 1935, in EXECUTE_NORMALIZED to_return = cls.execute(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy_extras/nodes_custom_sampler.py", line 1049, in execute samples = guider.sample(noise.generate_noise(latent), latent_image, sampler, sigmas, denoise_mask=noise_mask, callback=callback, disable_pbar=disable_pbar, seed=noise.seed) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1333, in sample output = executor.execute(noise, latent_image, sampler, sigmas, denoise_mask, callback, disable_pbar, seed, latent_shapes=latent_shapes) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1260, in outer_sample output = self.inner_sample(noise, latent_image, device, sampler, sigmas, denoise_mask, callback, disable_pbar, seed, latent_shapes=latent_shapes) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1235, in inner_sample samples = executor.execute(self, sigmas, extra_args, callback, noise, latent_image, denoise_mask, disable_pbar) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1005, in sample samples = self.sampler_function(model_k, noise, sigmas, extra_args=extra_args, callback=k_callback, disable=disable_pbar, **self.extra_options) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 124, in decorate_context return func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/k_diffusion/sampling.py", line 1460, in sample_res_multistep return res_multistep(model, x, sigmas, extra_args=extra_args, callback=callback, disable=disable, s_noise=s_noise, noise_sampler=noise_sampler, eta=0., cfg_pp=False) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 124, in decorate_context return func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/k_diffusion/sampling.py", line 1418, in res_multistep denoised = model(x, sigmas[i] * s_in, **extra_args) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 640, in __call__ out = self.inner_model(x, sigma, model_options=model_options, seed=seed) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1208, in __call__ return self.outer_predict_noise(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1215, in outer_predict_noise ).execute(x, timestep, model_options, seed) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 1218, in predict_noise return sampling_function(self.inner_model, x, timestep, self.conds.get("negative", None), self.conds.get("positive", None), self.cfg, model_options=model_options, seed=seed) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 620, in sampling_function out = calc_cond_batch(model, conds, x, timestep, model_options) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 211, in calc_cond_batch return _calc_cond_batch_outer(model, conds, x_in, timestep, model_options) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 219, in _calc_cond_batch_outer return executor.execute(model, conds, x_in, timestep, model_options) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/samplers.py", line 335, in _calc_cond_batch output = model.apply_model(input_x, timestep_, **c).chunk(batch_chunks) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/model_base.py", line 196, in apply_model return comfy.patcher_extension.WrapperExecutor.new_class_executor( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/model_base.py", line 240, in _apply_model model_output = self.diffusion_model(xc, t, context=context, control=control, transformer_options=transformer_options, **extra_conds) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ldm/minimax/model.py", line 499, in forward return comfy.patcher_extension.WrapperExecutor.new_class_executor( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/patcher_extension.py", line 113, in execute return self.original(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ldm/minimax/model.py", line 629, in _forward h = block(h, t_emb, mod_segments, rope_freqs, transformer_options=transformer_options) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ldm/minimax/model.py", line 270, in forward x = _mod_gate(x, gate_msa, self.attn(h, rope_freqs=rope_freqs, transformer_options=transformer_options), mod_segments) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ldm/minimax/model.py", line 158, in forward q, k, v = self.qkv_proj(x).split(self.heads * self.head_dim, dim=-1) ^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ops.py", line 1405, in forward output = self.forward_comfy_cast_weights( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ops.py", line 1342, in forward_comfy_cast_weights x = self._forward(input, weight, bias) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/comfy/ops.py", line 1313, in _forward return torch.nn.functional.linear(input, weight, bias) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/tensor/base.py", line 362, in __torch_dispatch__ return op_handlers[parent_cls](qt, args, kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/tensor/int8.py", line 287, in _handle_int8_linear_tensorwise return torch.ops.comfy_kitchen.int8_linear( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/_ops.py", line 1269, in __call__ return self._op(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/_library/custom_ops.py", line 347, in backend_impl result = self._backend_fns[device_type](*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/_compile.py", line 54, in inner return disable_fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/_dynamo/eval_frame.py", line 1263, in _fn return fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^ File "/venv/lib/python3.12/site-packages/torch/_library/custom_ops.py", line 382, in wrapped_fn return fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 1222, in _op_int8_linear return impl(**kwargs) ^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 1022, in int8_linear result = _int8_matmul_accumulate(x_8, weight.T.contiguous()) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 779, in _int8_matmul_accumulate result = fast_int8_mm(a, b) ^^^^^^^^^^^^^^^^^^ File "/comfyui/venv/lib/python3.12/site-packages/comfy_kitchen/backends/eager/quantization.py", line 754, in fast_int8_mm return torch._int_mm(lhs, rhs) ^^^^^^^^^^^^^^^^^^^^^^^ RuntimeError: CUDA error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblasLtMatmulAlgoGetHeuristic( ltHandle, computeDesc.descriptor(), Adesc.descriptor(), Bdesc.descriptor(), Cdesc.descriptor(), Cdesc.descriptor(), preference.descriptor(), 1, &heuristicResult, &returnedResult)` ``` </details> **EDIT:** Tested both of the solutions posted above and they work (obviously I had to remove the `device.type != "cuda"` check because AMD gpus are cuda in PyTorch). The `int8_linear` fix is only about 7-8% faster on my RX 6800 16GB. Using the default T2V example that comes with ComfyUI; ``` fast_int8_mm 37:37<00:00, 112.87s/it int8_linear 35:09<00:00, 105.48s/it ``` **EDIT2:** Found https://github.com/patientx/ComfyUI-INT8-Fast-ROCM which has a faster implementation for RDNA2 using custom triton kernels. Still, it's only 17-18% faster than `int8_linear` on the default T2V template @ 0.2 megapixels; ``` int8_linear 20/20 [20:50<00:00, 62.53s/it] INT8-Fast-ROCM 20/20 [17:12<00:00, 51.64s/it] ``` And I guess that's as about good as I'm going get on this hardware without LoRA/cache/other speedup hacks. Tbh I'm still amazed that I can run this SOTA model at all!

bsakaguchi commented

2026-10-26 22:19:24 +00:00

Author

Owner

@Maxim-Mazurok commented on GitHub (Aug 30, 2026):

I can reproduce this on Apple Silicon with TRELLIS.2 using trellis_2_int8_convrot.safetensors.

Environment:

ComfyUI 0.34.2
comfy-kitchen 0.2.31
PyTorch 2.12.1
Python 3.13.12
macOS 26.6.2
MPS, 32 GB unified memory

Sampling fails immediately at step 0 in comfy_kitchen.backends.eager.quantization.fast_int8_mm() because torch._int_mm has no MPS implementation. This confirms the issue is still present with the newer comfy-kitchen version and the TRELLIS.2 INT8 ConvRot checkpoint.

PYTORCH_ENABLE_MPS_FALLBACK=1 is the suggested workaround, although I have not yet measured its performance for this workflow.

@Maxim-Mazurok commented on GitHub (Aug 30, 2026): I can reproduce this on Apple Silicon with TRELLIS.2 using `trellis_2_int8_convrot.safetensors`. Environment: - ComfyUI 0.34.2 - comfy-kitchen 0.2.31 - PyTorch 2.12.1 - Python 3.13.12 - macOS 26.6.2 - MPS, 32 GB unified memory Sampling fails immediately at step 0 in `comfy_kitchen.backends.eager.quantization.fast_int8_mm()` because `torch._int_mm` has no MPS implementation. This confirms the issue is still present with the newer comfy-kitchen version and the TRELLIS.2 INT8 ConvRot checkpoint. `PYTORCH_ENABLE_MPS_FALLBACK=1` is the suggested workaround, although I have not yet measured its performance for this workflow.

bsakaguchi referenced this issue

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

[PR #29] [CLOSED] Adding c10d ops and additional aten ops for DTensor and FSDP2 #89

No labels

pull-request

No milestone

No project

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#29

No description provided.

Rows
Columns

[GH-ISSUE #92] eager backend advertises int8_linear on MPS but dispatches to CUDA-only torch._int_mm #29

Describe the bug

Reproduce

Traceback

Environment

Notes

Why a fast_int8_mm fallback is slow

What worked instead

End-to-end, MiniMax H3

Why a `fast_int8_mm` fallback is slow