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

[PR #272] Decline fp16-accumulate CUTLASS kernels on sub-sm_80 GPUs #278

New issue

Open

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

bsakaguchi commented

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

Owner

📋 Pull Request Information

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

Base: main ← Head: sm75-fp16-conv3d-capability-gate

📝 Commits (2)

d57a90c Add long key sequence tests for int8 attention
8d09387 Decline fp16-accumulate CUTLASS kernels on sub-sm_80 GPUs

📊 Changes

4 files changed (+127 additions, -0 deletions)

View changed files

📝 comfy_kitchen/backends/cuda/__init__.py (+5 -0)
📝 tests/test_fp16_conv3d.py (+39 -0)
📝 tests/test_fp16_linear.py (+33 -0)
📝 tests/test_int8_attention.py (+50 -0)

📄 Description

Fixes #258

Problem

On sm_75 (e.g. RTX 2080 Ti), starting ComfyUI with --fast fp16_accumulation routes the MiniMax H3 VAE to the CUTLASS fp16-accumulate conv3d kernel, which is a cutlass::arch::Sm80 kernel (multistage cp.async pipelines). Launching it on sm_75 aborts the CUDA context (Xid 43 in #258).

fp16_linear already declines on sub-sm_80 via its registry min_compute_capability=(8, 0) (added together with the kernels in #167), but the fp16_conv3d constraints entry was missing the field, and neither Python entry point checked the device itself:

get_capable_backend("fp16_conv3d", kwargs) with the registry capability forced to (7, 5) still selected the CUDA backend (reproduced on main)
direct callers of the CUDA backend bypass the registry entirely

Change

add min_compute_capability=(8, 0) to the fp16_conv3d constraints entry (fp16_conv3d_out inherits it through with_out_param)
add a per-device sm_80 gate to the supported checks of fp16_linear and _cutlass_fp16_conv3d, the same double-guard pattern sol_attn uses — callers that bypass the registry now fall back to torch instead of aborting

Sub-sm_80 devices now route to the eager backend (plain torch.conv3d / cuBLAS), the documented fallback requested in #258.

Tests

Added TestSm75CapabilityGate to tests/test_fp16_conv3d.py and tests/test_fp16_linear.py:

with the registry capability mocked to (7, 5): fp16_conv3d routes to eager (previously cuda), the top-level ck.fp16_conv3d returns torch's conv bit-exact, and the direct entries decline
on real sm_80+ hardware both ops still select the CUDA backend

Verified on an NVIDIA Thor (sm_110, torch 2.11.0+cu130): the 7 new tests pass and the full test_fp16_conv3d.py / test_fp16_linear.py / test_constraints.py suites still pass (101 passed). One failure in test_gemm_bias_contract.py reproduces identically on clean main on this architecture and is unrelated to this change.

No SM75 hardware was available; the fix is verified at the routing layer (capability-mocked), plus the kernels being sm_80-only by construction (cp.async multistage pipelines).

_{🔄 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/272 **Author:** [@potoior](https://github.com/potoior) **Created:** 10/23/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `sm75-fp16-conv3d-capability-gate` --- ### 📝 Commits (2) - [`d57a90c`](https://github.com/Comfy-Org/comfy-kitchen/commit/d57a90c2992b08279514aa76bde4bdbb082988cb) Add long key sequence tests for int8 attention - [`8d09387`](https://github.com/Comfy-Org/comfy-kitchen/commit/8d09387e6252806444841b9ff547ed453a63d508) Decline fp16-accumulate CUTLASS kernels on sub-sm_80 GPUs ### 📊 Changes **4 files changed** (+127 additions, -0 deletions) <details> <summary>View changed files</summary> 📝 `comfy_kitchen/backends/cuda/__init__.py` (+5 -0) 📝 `tests/test_fp16_conv3d.py` (+39 -0) 📝 `tests/test_fp16_linear.py` (+33 -0) 📝 `tests/test_int8_attention.py` (+50 -0) </details> ### 📄 Description Fixes #258 **Problem** On sm_75 (e.g. RTX 2080 Ti), starting ComfyUI with `--fast fp16_accumulation` routes the MiniMax H3 VAE to the CUTLASS fp16-accumulate conv3d kernel, which is a `cutlass::arch::Sm80` kernel (multistage cp.async pipelines). Launching it on sm_75 aborts the CUDA context (Xid 43 in #258). `fp16_linear` already declines on sub-sm_80 via its registry `min_compute_capability=(8, 0)` (added together with the kernels in #167), but the fp16_conv3d constraints entry was missing the field, and neither Python entry point checked the device itself: - `get_capable_backend("fp16_conv3d", kwargs)` with the registry capability forced to (7, 5) still selected the CUDA backend (reproduced on main) - direct callers of the CUDA backend bypass the registry entirely **Change** - add `min_compute_capability=(8, 0)` to the fp16_conv3d constraints entry (`fp16_conv3d_out` inherits it through `with_out_param`) - add a per-device sm_80 gate to the `supported` checks of `fp16_linear` and `_cutlass_fp16_conv3d`, the same double-guard pattern sol_attn uses — callers that bypass the registry now fall back to torch instead of aborting Sub-sm_80 devices now route to the eager backend (plain `torch.conv3d` / cuBLAS), the documented fallback requested in #258. **Tests** Added `TestSm75CapabilityGate` to `tests/test_fp16_conv3d.py` and `tests/test_fp16_linear.py`: - with the registry capability mocked to (7, 5): fp16_conv3d routes to eager (previously cuda), the top-level `ck.fp16_conv3d` returns torch's conv bit-exact, and the direct entries decline - on real sm_80+ hardware both ops still select the CUDA backend Verified on an NVIDIA Thor (sm_110, torch 2.11.0+cu130): the 7 new tests pass and the full `test_fp16_conv3d.py` / `test_fp16_linear.py` / `test_constraints.py` suites still pass (101 passed). One failure in `test_gemm_bias_contract.py` reproduces identically on clean main on this architecture and is unrelated to this change. No SM75 hardware was available; the fix is verified at the routing layer (capability-mocked), plus the kernels being sm_80-only by construction (cp.async multistage pipelines). --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi added the

pull-request

label

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

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

No description provided.

Rows
Columns