Comfy-Org/comfy-kitchen

Fork 0

mirror of https://github.com/Comfy-Org/comfy-kitchen.git synced 2026-11-04 23:37:42 +00:00

[GH-ISSUE #136] INT8 GEMM kernels index the output in int32: illegal memory access once m*n >= 2**31 (MiniMax H3 at high resolution x length) #44

New issue

Open

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

bsakaguchi commented

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

Owner

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

Summary

Both Triton INT8 GEMM kernels in comfy_kitchen/backends/triton/quantization.py
(_int8_matmul_dequant_kernel and _int8_matmul_dequant_per_row_kernel) compute the
output address with int32 arithmetic. As soon as m * n >= 2**31 the offset wraps and
the kernel faults with CUDA error: an illegal memory access was encountered, which
aborts the whole ComfyUI process.

offs_am = (pid_m * block_m + tl.arange(0, block_m)) % m      # int32
offs_bn = (pid_n * block_n + tl.arange(0, block_n)) % n      # int32
...
c_ptrs = c_ptr + stride_cm * offs_am[:, None] + stride_cn * offs_bn[None, :]

stride_cm is n, so the largest store offset is (m-1)*n + (n-1) and it crosses
2**31 exactly when m*n does. m is the packed sequence length and n the projection
width, which is why the failure depends on resolution x length rather than on either
one alone.

This is what takes MiniMax H3 down on a 24 GB card above a certain size. It is independent
of dynamic VRAM, pinned memory and the attention backend.

Minimal reproduction (no model, no weights)

One m per process, because the CUDA context is unusable after the fault:

import sys, torch
from comfy_kitchen.backends.triton.quantization import int8_linear

m = int(sys.argv[1])
n = int(sys.argv[2]) if len(sys.argv) > 2 else 21504
k = 5376
x  = torch.randn(m, k, device='cuda', dtype=torch.bfloat16)
w  = torch.randint(-127, 127, (n, k), device='cuda', dtype=torch.int8)
ws = torch.rand(n, device='cuda', dtype=torch.float32) * 0.01
out = int8_linear(x, w, ws, out_dtype=torch.bfloat16)
torch.cuda.synchronize()
print("OK", tuple(out.shape))

n = 21504, k = 5376 are the real MiniMax H3 QKV projection dimensions (the same shapes
as in #98).

Measured boundary, three different n:

n	last good m	first failing m	m*n at the boundary
21504	99,865	99,900	2.147e9
10752	199,000	200,500	2.147e9
5376	399,000	400,500	2.147e9

The crossing sits at m*n = 2**31 in all three cases. Below the boundary the results are
numerically fine (~1 % against a bf16 reference, i.e. plain quantization error), so this is
a hard fault at a fixed index limit and not a gradual accuracy problem.

Stack trace

With CUDA_LAUNCH_BLOCKING=1 the fault is pinned to the kernel launch itself rather than
to a later synchronize:

  File "comfy_kitchen\backends\triton\quantization.py", line 1078, in int8_linear
    _int8_matmul_dequant_per_row_kernel[grid](
  File "triton\runtime\autotuner.py", line 164, in _bench
    return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))
  File "triton\runtime\jit.py", line 744, in run
    kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, ...)
  File "triton\backends\nvidia\driver.py", line 752, in __call__
    self.launch(gridX, gridY, gridZ, stream, function, ...)
RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered

Without it the error surfaces one frame later, in triton/testing.py do_bench -> di.synchronize(), which makes it look like an autotuner problem. It is not — the
autotuner is simply the first thing that launches the kernel.

Real-world impact (MiniMax H3, RTX 3090 24 GB)

Text-to-video with minimax_h3_fl2va_pruned_int8_convrot and --enable-triton-backend:

size	pixels x frames	before	with the fix below
1344x768, 243 frames (10 s)	251 M	ok	ok
1344x768, 294 frames (12 s)	303 M	crash after 43 s	ok, 410 s
1344x768, 362 frames (15 s)	374 M	crash after 43 s	ok, 620 s

1024x576 at 15 s — more frames, fewer pixels — always worked, which matches the m*n law
rather than a frame-count or VRAM limit.

Ruled out here by direct test: --disable-dynamic-vram, --disable-pinned-memory,
comfy-kitchen 0.2.30 vs 0.2.31 (the kernel file is byte-identical between the two), and the
w4a8 and GGUF weight variants — w4a8 goes through the same int8_linear and fails at the
same shapes.

Suggested fix

Casting the offsets to int64 in both kernels is enough:

offs_am = ((pid_m * block_m + tl.arange(0, block_m)) % m).to(tl.int64)
offs_bn = ((pid_n * block_n + tl.arange(0, block_n)) % n).to(tl.int64)

With that applied locally, m = 200,000, n = 21504 runs clean and the output is
numerically correct (~1 % against the reference, the same as below the old boundary), and
both H3 sizes above complete end to end. I have not measured the throughput cost of the
wider index arithmetic; if it turns out to matter, gating the int64 path on
m * n >= 2**31 would leave the fast path untouched.

Related: #98 (same tall-M H3 shapes, but about config selection and performance) and #64
(same error text on AMD gfx1151, different backend).

Environment

RTX 3090 24 GB, driver 595.97, Windows 11
comfy-kitchen 0.2.31 (and 0.2.30, identical kernel file)
ComfyUI 0.32.0, torch 2.10.0+cu128, triton-windows 3.6.0.post26, Python 3.12.11

Originally created by @mjungnickel18 on GitHub (Aug 24, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/136 ## Summary Both Triton INT8 GEMM kernels in `comfy_kitchen/backends/triton/quantization.py` (`_int8_matmul_dequant_kernel` and `_int8_matmul_dequant_per_row_kernel`) compute the output address with **int32** arithmetic. As soon as `m * n >= 2**31` the offset wraps and the kernel faults with `CUDA error: an illegal memory access was encountered`, which aborts the whole ComfyUI process. ```python offs_am = (pid_m * block_m + tl.arange(0, block_m)) % m # int32 offs_bn = (pid_n * block_n + tl.arange(0, block_n)) % n # int32 ... c_ptrs = c_ptr + stride_cm * offs_am[:, None] + stride_cn * offs_bn[None, :] ``` `stride_cm` is `n`, so the largest store offset is `(m-1)*n + (n-1)` and it crosses `2**31` exactly when `m*n` does. `m` is the packed sequence length and `n` the projection width, which is why the failure depends on **resolution x length** rather than on either one alone. This is what takes MiniMax H3 down on a 24 GB card above a certain size. It is independent of dynamic VRAM, pinned memory and the attention backend. ## Minimal reproduction (no model, no weights) One `m` per process, because the CUDA context is unusable after the fault: ```python import sys, torch from comfy_kitchen.backends.triton.quantization import int8_linear m = int(sys.argv[1]) n = int(sys.argv[2]) if len(sys.argv) > 2 else 21504 k = 5376 x = torch.randn(m, k, device='cuda', dtype=torch.bfloat16) w = torch.randint(-127, 127, (n, k), device='cuda', dtype=torch.int8) ws = torch.rand(n, device='cuda', dtype=torch.float32) * 0.01 out = int8_linear(x, w, ws, out_dtype=torch.bfloat16) torch.cuda.synchronize() print("OK", tuple(out.shape)) ``` `n = 21504, k = 5376` are the real MiniMax H3 QKV projection dimensions (the same shapes as in #98). Measured boundary, three different `n`: | n | last good m | first failing m | m*n at the boundary | |---|---|---|---| | 21504 | 99,865 | 99,900 | 2.147e9 | | 10752 | 199,000 | 200,500 | 2.147e9 | | 5376 | 399,000 | 400,500 | 2.147e9 | The crossing sits at `m*n = 2**31` in all three cases. Below the boundary the results are numerically fine (~1 % against a bf16 reference, i.e. plain quantization error), so this is a hard fault at a fixed index limit and not a gradual accuracy problem. ## Stack trace With `CUDA_LAUNCH_BLOCKING=1` the fault is pinned to the kernel launch itself rather than to a later synchronize: ``` File "comfy_kitchen\backends\triton\quantization.py", line 1078, in int8_linear _int8_matmul_dequant_per_row_kernel[grid]( File "triton\runtime\autotuner.py", line 164, in _bench return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8)) File "triton\runtime\jit.py", line 744, in run kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, ...) File "triton\backends\nvidia\driver.py", line 752, in __call__ self.launch(gridX, gridY, gridZ, stream, function, ...) RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered ``` Without it the error surfaces one frame later, in `triton/testing.py do_bench -> di.synchronize()`, which makes it look like an autotuner problem. It is not — the autotuner is simply the first thing that launches the kernel. ## Real-world impact (MiniMax H3, RTX 3090 24 GB) Text-to-video with `minimax_h3_fl2va_pruned_int8_convrot` and `--enable-triton-backend`: | size | pixels x frames | before | with the fix below | |---|---|---|---| | 1344x768, 243 frames (10 s) | 251 M | ok | ok | | 1344x768, 294 frames (12 s) | 303 M | **crash after 43 s** | ok, 410 s | | 1344x768, 362 frames (15 s) | 374 M | **crash after 43 s** | ok, 620 s | 1024x576 at 15 s — more frames, fewer pixels — always worked, which matches the `m*n` law rather than a frame-count or VRAM limit. Ruled out here by direct test: `--disable-dynamic-vram`, `--disable-pinned-memory`, comfy-kitchen 0.2.30 vs 0.2.31 (the kernel file is byte-identical between the two), and the w4a8 and GGUF weight variants — w4a8 goes through the same `int8_linear` and fails at the same shapes. ## Suggested fix Casting the offsets to int64 in both kernels is enough: ```python offs_am = ((pid_m * block_m + tl.arange(0, block_m)) % m).to(tl.int64) offs_bn = ((pid_n * block_n + tl.arange(0, block_n)) % n).to(tl.int64) ``` With that applied locally, `m = 200,000, n = 21504` runs clean and the output is numerically correct (~1 % against the reference, the same as below the old boundary), and both H3 sizes above complete end to end. I have not measured the throughput cost of the wider index arithmetic; if it turns out to matter, gating the int64 path on `m * n >= 2**31` would leave the fast path untouched. Related: #98 (same tall-M H3 shapes, but about config selection and performance) and #64 (same error text on AMD gfx1151, different backend). ## Environment - RTX 3090 24 GB, driver 595.97, Windows 11 - comfy-kitchen 0.2.31 (and 0.2.30, identical kernel file) - ComfyUI 0.32.0, torch 2.10.0+cu128, triton-windows 3.6.0.post26, Python 3.12.11

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

No description provided.

Rows
Columns