[GH-ISSUE #280] INT8 linears on sm_121 (GB10 / DGX Spark) run at ~37% of tensor-core peak via cutlass_int8_dequant; a fused-dequant Triton GEMM reaches ~88%, bit-identical #75

New issue

Open

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

bsakaguchi commented

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

Owner

Originally created by @kanthai on GitHub (Oct 26, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/280

Summary

On NVIDIA GB10 (DGX Spark, sm_121), comfy-kitchen 0.2.37 sends INT8 tensorwise linears to _C.cutlass_int8_dequant. It sustains about 72 TOPS, slower than the BF16 GEMM on the same chip. The mma.sync INT8 peak at the clock we run (1989 MHz) is 196 TOPS. A simple Triton GEMM with an int32 accumulator and the dequant (x_scale·w_scale + bias) fused into the epilogue reaches about 170 TOPS (88% of peak). Its output is bit-identical to the CUTLASS kernel. End to end, this makes INT8 checkpoints faster than BF16/FP8 on GB10 instead of slower.

Related: #98 (same class of problem — CUTLASS INT8 dequant picking a poor config for tall MiniMax H3 shapes — on an RTX A6000) and #237 (sm86 tile chooser; its description notes the generic select_fused_int8_config thresholds were fitted on Ada and Blackwell). GB10 is Blackwell but with a very different memory system (unified LPDDR5x), which may be why the Blackwell fit misses here.

Question for maintainers: would you take a Triton fallback for sm_121, or would you rather add an sm_121 entry to the CUTLASS tile chooser (like #237 does for sm86)? We can share the sweep data either way.

Environment

comfy-kitchen 0.2.37 (0.2.40 routes the same way), torch 2.14.0+cu130, triton 3.8.0, ComfyUI 0.41.3, NVIDIA DGX Spark: GB10 with 48 SMs, sm_121, unified LPDDR5x (~273 GB/s).

Measured mma.sync peaks at 1989 MHz: NVFP4 391, INT8 = FP8 196, BF16 98 TOPS.

Where the time goes (installed sources)

comfy_kitchen/backends/cuda/__init__.py:321: _cuda_device_supports_cutlass_int8_dequant returns True for every major >= 8, so sm_121 takes the CUTLASS path.
The _C.cutlass_int8_dequant(...) call sites are at about lines 1009, 1192 and 2221 (int8 GEMM, scaled_mm, int8_linear).
Other routes we tried:
- COMFY_KITCHEN_DISABLE_CUTLASS=1 (cuBLAS int32 GEMM + separate dequant) reaches about 105 TOPS but needs about 3 GB of int32 scratch for H3 shapes. That round-trip is expensive at ~273 GB/s.
- Fusing the dequant into the epilogue is the main win on this memory system.

Kernel used

Triton, an int32 accumulator over K, fp32 epilogue: acc * x_scale[m] * w_scale[n] (+ bias[n]) → the output dtype.

Config: BM=128, BN=128, BK=64, num_warps=8, num_stages=4, grouped launch G=8, with int64 row/column offsets. Without int64, M·N or M·K past 2^31 elements wraps in int32. That happens with CFG batches or long video clips (M ≥ 74.9k at N = 28672).

The sweep found this config best, or within 1%, on every H3 linear shape at 29k and 36k tokens. Residual-fused calls, M < 128, K not a multiple of 16 and non-contiguous tensors stay on CUTLASS (so anything CUTLASS declines still takes comfy-kitchen's stock cuBLAS route).

Results (same graphs, same seeds)

Workload	CUTLASS (stock)	Triton fused	Notes
MiniMax H3 ref2v 8 s, 1344x768, 8 steps, sampling (2-node split)	369.7 s	302.9 s (−18%)	latent sha256, all 192 decoded frames and audio identical
MiniMax H3 10 s clip, sampling	507.4 s	422.2 s
Krea2 fine-tune 1344x768 8 steps, int8_convrot (stock ComfyUI)	18.2 s/img (fp8 default)	13.2 s; 10.0 s with TorchCompileModel
Qwen Image 2.1 int8_convrot, 25 steps, 1024², 2-node Monarch	16.3 s (bf16)	11.1 s	31–40 dB vs bf16
Qwen Image 2.1 turbo int8_convrot, 8 steps	5.1 s (bf16)	3.6 s

Without the faster kernel, INT8 on GB10 is slower than BF16. With it, Comfy-Org's int8_convrot releases are the fastest and the most faithful option on this hardware: in our tests, NVFP4 W4A4 visibly changes content. For example, a Krea2 fine-tune made INT8 from the official BF16 gets 21.7 dB PSNR vs BF16, against 17.0 dB for FP8, at 9.6 vs 17.8 s/img.

For contrast, on an RTX PRO 2000 Blackwell (sm_120, 34 SMs, PCIe), stock CUTLASS INT8 is already good: 1308 ms/forward vs Triton 1509. The gap seems specific to GB10 (low memory bandwidth, unified memory), not to all of Blackwell.

Exactness

INT8×INT8 accumulation in int32 is exact. The epilogue computes in fp32 in the same order as the CUTLASS epilogue (acc * xs * ws + bias). The evidence that the outputs match:

unit tests over random and real H3 shapes: torch.equal against _C.cutlass_int8_dequant;
two independent audits (the second re-checked M up to 75000×28672, i.e. M·N > 2³¹, K = 16…65536, fp16/fp32 output, tensor-wide scales, no-bias: identical for comfy-kitchen's call pattern);
a full H3 render with identical latent hash, frames and audio.

How we deploy it now

We deploy it as a .pth-loaded proxy around comfy_kitchen.backends.cuda._C that swaps only cutlass_int8_dequant, with the same signature. It is version-guarded to comfy-kitchen 0.2.37 and falls back to stock on any Triton import or build failure. We would much rather this lived upstream behind a capability check (for example sm_121 / integrated GPUs, or a measured-throughput gate).

Source of the proxy and kernel, the benchmark tools and the full write-up: https://github.com/kanthai/gb10-comfyui-tuning (patches/ck_int8_patch.py, docs/02-int8-triton-gemm.md).

Originally created by @kanthai on GitHub (Oct 26, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/280 ## Summary On NVIDIA GB10 (DGX Spark, sm_121), comfy-kitchen 0.2.37 sends INT8 tensorwise linears to `_C.cutlass_int8_dequant`. It sustains about **72 TOPS**, slower than the BF16 GEMM on the same chip. The `mma.sync` INT8 peak at the clock we run (1989 MHz) is 196 TOPS. A simple Triton GEMM with an int32 accumulator and the dequant (x_scale·w_scale + bias) fused into the epilogue reaches about **170 TOPS** (88% of peak). Its output is **bit-identical** to the CUTLASS kernel. End to end, this makes INT8 checkpoints faster than BF16/FP8 on GB10 instead of slower. Related: #98 (same class of problem — CUTLASS INT8 dequant picking a poor config for tall MiniMax H3 shapes — on an RTX A6000) and #237 (sm86 tile chooser; its description notes the generic `select_fused_int8_config` thresholds were fitted on Ada and Blackwell). GB10 is Blackwell but with a very different memory system (unified LPDDR5x), which may be why the Blackwell fit misses here. Question for maintainers: would you take a Triton fallback for sm_121, or would you rather add an sm_121 entry to the CUTLASS tile chooser (like #237 does for sm86)? We can share the sweep data either way. ## Environment comfy-kitchen 0.2.37 (0.2.40 routes the same way), torch 2.14.0+cu130, triton 3.8.0, ComfyUI 0.41.3, NVIDIA DGX Spark: GB10 with 48 SMs, sm_121, unified LPDDR5x (~273 GB/s). Measured `mma.sync` peaks at 1989 MHz: NVFP4 391, INT8 = FP8 196, BF16 98 TOPS. ## Where the time goes (installed sources) - `comfy_kitchen/backends/cuda/__init__.py:321`: `_cuda_device_supports_cutlass_int8_dequant` returns True for every `major >= 8`, so sm_121 takes the CUTLASS path. - The `_C.cutlass_int8_dequant(...)` call sites are at about lines 1009, 1192 and 2221 (int8 GEMM, `scaled_mm`, `int8_linear`). - Other routes we tried: - `COMFY_KITCHEN_DISABLE_CUTLASS=1` (cuBLAS int32 GEMM + separate dequant) reaches about 105 TOPS but needs about 3 GB of int32 scratch for H3 shapes. That round-trip is expensive at ~273 GB/s. - Fusing the dequant into the epilogue is the main win on this memory system. ## Kernel used Triton, an int32 accumulator over K, fp32 epilogue: `acc * x_scale[m] * w_scale[n] (+ bias[n])` → the output dtype. Config: BM=128, BN=128, BK=64, num_warps=8, num_stages=4, grouped launch G=8, with **int64 row/column offsets**. Without int64, M·N or M·K past 2^31 elements wraps in int32. That happens with CFG batches or long video clips (M ≥ 74.9k at N = 28672). The sweep found this config best, or within 1%, on every H3 linear shape at 29k and 36k tokens. Residual-fused calls, M < 128, K not a multiple of 16 and non-contiguous tensors stay on CUTLASS (so anything CUTLASS declines still takes comfy-kitchen's stock cuBLAS route). ## Results (same graphs, same seeds) | Workload | CUTLASS (stock) | Triton fused | Notes | |---|---|---|---| | MiniMax H3 ref2v 8 s, 1344x768, 8 steps, sampling (2-node split) | 369.7 s | 302.9 s (**−18%**) | latent sha256, all 192 decoded frames and audio identical | | MiniMax H3 10 s clip, sampling | 507.4 s | 422.2 s | | | Krea2 fine-tune 1344x768 8 steps, int8_convrot (stock ComfyUI) | 18.2 s/img (fp8 default) | 13.2 s; 10.0 s with TorchCompileModel | | | Qwen Image 2.1 int8_convrot, 25 steps, 1024², 2-node Monarch | 16.3 s (bf16) | 11.1 s | 31–40 dB vs bf16 | | Qwen Image 2.1 turbo int8_convrot, 8 steps | 5.1 s (bf16) | 3.6 s | | Without the faster kernel, INT8 on GB10 is slower than BF16. With it, Comfy-Org's int8_convrot releases are the fastest *and* the most faithful option on this hardware: in our tests, NVFP4 W4A4 visibly changes content. For example, a Krea2 fine-tune made INT8 from the official BF16 gets 21.7 dB PSNR vs BF16, against 17.0 dB for FP8, at 9.6 vs 17.8 s/img. For contrast, on an RTX PRO 2000 Blackwell (sm_120, 34 SMs, PCIe), stock CUTLASS INT8 is already good: 1308 ms/forward vs Triton 1509. The gap seems specific to GB10 (low memory bandwidth, unified memory), not to all of Blackwell. ## Exactness INT8×INT8 accumulation in int32 is exact. The epilogue computes in fp32 in the same order as the CUTLASS epilogue (`acc * xs * ws + bias`). The evidence that the outputs match: - unit tests over random and real H3 shapes: `torch.equal` against `_C.cutlass_int8_dequant`; - two independent audits (the second re-checked M up to 75000×28672, i.e. M·N > 2³¹, K = 16…65536, fp16/fp32 output, tensor-wide scales, no-bias: identical for comfy-kitchen's call pattern); - a full H3 render with identical latent hash, frames and audio. ## How we deploy it now We deploy it as a `.pth`-loaded proxy around `comfy_kitchen.backends.cuda._C` that swaps only `cutlass_int8_dequant`, with the same signature. It is version-guarded to comfy-kitchen 0.2.37 and falls back to stock on any Triton import or build failure. We would much rather this lived upstream behind a capability check (for example sm_121 / integrated GPUs, or a measured-throughput gate). Source of the proxy and kernel, the benchmark tools and the full write-up: https://github.com/kanthai/gb10-comfyui-tuning (`patches/ck_int8_patch.py`, `docs/02-int8-triton-gemm.md`).

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

No description provided.

Rows
Columns