[GH-ISSUE #258] SM75: MiniMax H3 VAE fp16_accumulation launches unsupported CUTLASS MMA and aborts #72

New issue

Open

opened 2026-10-26 22:19:57 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

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

Owner

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

Summary

On an RTX 2080 Ti (SM75), the MiniMax H3 FP16 video VAE crashes when ComfyUI is started with --fast fp16_accumulation and the comfy-kitchen CUDA backend is enabled.

The CUDA backend advertises fp16_conv3d and fp16_linear, but the H3 VAE reaches a CUTLASS FP16 MMA that is not implemented on this architecture. The process then terminates with an unspecified launch failure and the driver records Xid 43.

Removing --fast fp16_accumulation makes the same workflow complete with the CUDA backend enabled. Keeping the flag while disabling the comfy-kitchen CUDA backend also completes and is substantially faster on this machine.

This looks like a missing SM75 capability/launch guard: the CUDA implementation should decline this configuration and allow the documented cuDNN/eager fallback.

Environment

GPU: NVIDIA GeForce RTX 2080 Ti, 22528 MiB, compute capability 7.5
OS: Ubuntu 22.04.4 LTS, kernel 6.8.0-48-generic
NVIDIA driver: 580.178.04
ComfyUI: v0.42.0, commit 355cd6b14c92c19ac3188fe528e47a3f21ef5178
comfy-kitchen: 0.2.39
comfy-aimdo: 0.5.6
Python: 3.13.13
PyTorch: 2.11.0+cu130
Model: official MiniMax H3 I2V workflow with minimax_h3_video_vae_fp16.safetensors

Failing configuration

The relevant launch options are:

--fast fp16_accumulation
--use-ck-attention
--cuda-device 2
--disable-fast-disk
--disable-dynamic-vram

Dynamic VRAM is disabled in this reproduction, so this failure is separate from the earlier AIMDO HostBuffer.read_file_slice failure.

Relevant log output repeats the following message before the process aborts:

void cutlass::arch::Mma<cutlass::gemm::GemmShape<16, 8, 16>, 32,
cutlass::half_t, cutlass::layout::RowMajor,
cutlass::half_t, cutlass::layout::ColumnMajor,
cutlass::half_t, cutlass::layout::RowMajor,
cutlass::arch::OpMultiplyAdd>::operator()(...) const not implemented

CUDA error: unspecified launch failure

The Linux kernel then reports NVIDIA Xid 43 for the application channel.

Reproduction workflow

The official MiniMax H3 image-to-video workflow was reduced to:

256 x 160
22 frames at 24 fps
2 sampling steps
FP16 video VAE
FP32 audio VAE

The failure occurs in the H3 VAE path after DynamicVRAM has already been disabled.

A/B results

Configuration	Result	Execution time
CUDA backend enabled, `--fast fp16_accumulation` enabled	CUTLASS failure, process abort, Xid 43	failed
CUDA backend enabled, `--fast fp16_accumulation` removed	success, valid 256x160 H.264, 22 frames	417.772 s
CUDA backend disabled via `comfy_kitchen.disable_backend("cuda")`, fast flag retained	success, valid 256x160 H.264, 22 frames	104.547 s

No new Xid was recorded during either successful run.

Expected behavior

On SM75, unsupported CUTLASS FP16 H3 VAE shapes should be rejected before launch and fall back to cuDNN/eager, as described in Comfy-Org/comfy-kitchen#167 and Comfy-Org/ComfyUI#16187.

I can provide the full startup log and a minimal API workflow if needed.

Originally created by @ZanePoe on GitHub (Oct 19, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/258 ### Summary On an RTX 2080 Ti (SM75), the MiniMax H3 FP16 video VAE crashes when ComfyUI is started with `--fast fp16_accumulation` and the comfy-kitchen CUDA backend is enabled. The CUDA backend advertises `fp16_conv3d` and `fp16_linear`, but the H3 VAE reaches a CUTLASS FP16 MMA that is not implemented on this architecture. The process then terminates with an unspecified launch failure and the driver records Xid 43. Removing `--fast fp16_accumulation` makes the same workflow complete with the CUDA backend enabled. Keeping the flag while disabling the comfy-kitchen CUDA backend also completes and is substantially faster on this machine. This looks like a missing SM75 capability/launch guard: the CUDA implementation should decline this configuration and allow the documented cuDNN/eager fallback. ### Environment - GPU: NVIDIA GeForce RTX 2080 Ti, 22528 MiB, compute capability 7.5 - OS: Ubuntu 22.04.4 LTS, kernel 6.8.0-48-generic - NVIDIA driver: 580.178.04 - ComfyUI: v0.42.0, commit `355cd6b14c92c19ac3188fe528e47a3f21ef5178` - comfy-kitchen: 0.2.39 - comfy-aimdo: 0.5.6 - Python: 3.13.13 - PyTorch: 2.11.0+cu130 - Model: official MiniMax H3 I2V workflow with `minimax_h3_video_vae_fp16.safetensors` ### Failing configuration The relevant launch options are: ```text --fast fp16_accumulation --use-ck-attention --cuda-device 2 --disable-fast-disk --disable-dynamic-vram ``` Dynamic VRAM is disabled in this reproduction, so this failure is separate from the earlier AIMDO `HostBuffer.read_file_slice` failure. Relevant log output repeats the following message before the process aborts: ```text void cutlass::arch::Mma<cutlass::gemm::GemmShape<16, 8, 16>, 32, cutlass::half_t, cutlass::layout::RowMajor, cutlass::half_t, cutlass::layout::ColumnMajor, cutlass::half_t, cutlass::layout::RowMajor, cutlass::arch::OpMultiplyAdd>::operator()(...) const not implemented CUDA error: unspecified launch failure ``` The Linux kernel then reports NVIDIA Xid 43 for the application channel. ### Reproduction workflow The official MiniMax H3 image-to-video workflow was reduced to: - 256 x 160 - 22 frames at 24 fps - 2 sampling steps - FP16 video VAE - FP32 audio VAE The failure occurs in the H3 VAE path after DynamicVRAM has already been disabled. ### A/B results | Configuration | Result | Execution time | |---|---|---:| | CUDA backend enabled, `--fast fp16_accumulation` enabled | CUTLASS failure, process abort, Xid 43 | failed | | CUDA backend enabled, `--fast fp16_accumulation` removed | success, valid 256x160 H.264, 22 frames | 417.772 s | | CUDA backend disabled via `comfy_kitchen.disable_backend("cuda")`, fast flag retained | success, valid 256x160 H.264, 22 frames | 104.547 s | No new Xid was recorded during either successful run. ### Expected behavior On SM75, unsupported CUTLASS FP16 H3 VAE shapes should be rejected before launch and fall back to cuDNN/eager, as described in Comfy-Org/comfy-kitchen#167 and Comfy-Org/ComfyUI#16187. I can provide the full startup log and a minimal API workflow if needed.

bsakaguchi commented

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

Author

Owner

@potoior commented on GitHub (Oct 23, 2026):

Opened #272 with the fix: min_compute_capability=(8, 0) on the fp16_conv3d constraints (fp16_linear already had it) plus per-device sm_80 gates in the two Python entry points, with capability-mocked routing tests. On sm_75 the ops now decline and ComfyUI falls back to cuDNN/eager instead of aborting the context.

@potoior commented on GitHub (Oct 23, 2026): Opened #272 with the fix: `min_compute_capability=(8, 0)` on the fp16_conv3d constraints (fp16_linear already had it) plus per-device sm_80 gates in the two Python entry points, with capability-mocked routing tests. On sm_75 the ops now decline and ComfyUI falls back to cuDNN/eager instead of aborting the context.

bsakaguchi referenced this issue

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

[PR #74] [MERGED] Add HIP backend for AMD RDNA2-RDNA4 to the existing wheels #123

bsakaguchi referenced this issue

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

[PR #72] [CLOSED] Add HIP backend with WMMA kernels for AMD RDNA4 (gfx12) #125