[GH-ISSUE #144] CUDA backend cannot build with sm_70 (V100) in arch list: ptxas ldmatrix/cvt.pack errors + missing flash.h submodule #46

New issue

Open

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

bsakaguchi commented

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

Owner

Originally created by @Joevy-1112 on GitHub (Aug 30, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/144

CUDA backend cannot build with sm_70 (V100) in the arch list: ptxas errors on ldmatrix / cvt.pack + missing flash.h

Environment: Tesla V100 (sm_70), CUDA 12.8 toolkit (nvcc 12.8.61), Ubuntu 24.04, gcc 13.3.
Build command:

COMFY_CUDA_ARCHS="70;75-real;80-real" CUDA_HOME=... pip install . --no-build-isolation

Symptom

24 of 27 .cu files compile fine for sm_70 (including rms_rope.cu, adaln.cu, turing_int4/int8.cu, convrot_w4a4.cu, all sol_attn*.cu), but the build hard-fails on 3 sources:

1. ops/awq_w4a16.cu — ptxas cannot assemble compute_70:

ptxas .../awq_w4a16.compute_70.ptx, line 2026; error   : Feature 'ldmatrix' requires .target sm_75 or higher
ptxas .../awq_w4a16.compute_70.ptx, line 2026; error   : Modifier '.m8n8' requires .target sm_75 or higher
... (40+ lines)

ldmatrix is used directly in the file and via shared wrappers sage_attention/mma.cuh (ldmatrix_m8n8x2/x4/x4_trans) and ops/svdquant_utils.cuh (ldmatrix_x4) — none of these asm blocks are guarded by __CUDA_ARCH__, so this is a compile-time hard failure, not a runtime stub.

2. sage_attention/sage_attn_launcher.cu — ptxas:

ptxas .../sage_attn_launcher.compute_70.ptx, line 6082; error   : Feature 'cvt.pack' requires .target sm_72 or higher
... (20+ lines)

cvt.pack does not appear verbatim in the source — it is generated by the compiler from an intrinsic/packing macro on sm_72+, so it cannot be fixed by a simple grep-and-guard.

3. ops/flash_decode.cu — missing header (submodule-dependent):

fatal error: flash.h: 没有那个文件或目录

flash_decode.cu includes flash.h which lives in the third_party/flash-attention submodule (csrc/flash_attn/src/flash.h). A shallow clone without --recurse-submodules leaves it empty.

What works on sm_70

rms_rope.cu / adaln.cu compile and run correctly on V100 after building the remaining 24 sources with COMFY_CUDA_ARCHS="70;75-real;80-real".
Measured on V100 (driver 580.173.02): rms_rope_split_half 0.093 ms/iter vs eager 0.612 ms/iter (~6.6x); rms_adaln 0.139 ms/iter vs eager 0.240 ms/iter (~1.7x), 1000 iters, numerics finite.

Question for maintainers

How would you prefer to support sm_70 builds?

(a) Guard the ldmatrix/cvt.pack code paths with __CUDA_ARCH__ >= 750 / >= 720 (with a real fallback or trap on sm_70 — silently compiling away the asm without a fallback would produce wrong results), or
(b) Conditionally exclude the 3 sources from CUDA_SOURCES when the arch list contains 70 (documenting that awq_w4a16/sage_attn/flash_decode need sm_72+/sm_75+), or
(c) keep them in the build and let sm_70 users carry a local patch.

I have a working local build (the 3 sources excluded + bindings stubbed) with rms_rope/adaln verified on V100. I can share the patch or prepare a PR in whichever direction you prefer. Happy to do the __CUDA_ARCH__ guard work on mma.cuh / svdquant_utils.cuh if you confirm that's the desired shape — just want to avoid guessing about kernel fallback semantics in shared attention headers.

Originally created by @Joevy-1112 on GitHub (Aug 30, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/144 ## CUDA backend cannot build with sm_70 (V100) in the arch list: ptxas errors on ldmatrix / cvt.pack + missing flash.h **Environment**: Tesla V100 (sm_70), CUDA 12.8 toolkit (nvcc 12.8.61), Ubuntu 24.04, gcc 13.3. **Build command**: ``` COMFY_CUDA_ARCHS="70;75-real;80-real" CUDA_HOME=... pip install . --no-build-isolation ``` ### Symptom 24 of 27 `.cu` files compile fine for sm_70 (including `rms_rope.cu`, `adaln.cu`, `turing_int4/int8.cu`, `convrot_w4a4.cu`, all `sol_attn*.cu`), but the build hard-fails on 3 sources: **1. `ops/awq_w4a16.cu`** — ptxas cannot assemble compute_70: ``` ptxas .../awq_w4a16.compute_70.ptx, line 2026; error : Feature 'ldmatrix' requires .target sm_75 or higher ptxas .../awq_w4a16.compute_70.ptx, line 2026; error : Modifier '.m8n8' requires .target sm_75 or higher ... (40+ lines) ``` `ldmatrix` is used directly in the file and via shared wrappers `sage_attention/mma.cuh` (`ldmatrix_m8n8x2/x4/x4_trans`) and `ops/svdquant_utils.cuh` (`ldmatrix_x4`) — none of these asm blocks are guarded by `__CUDA_ARCH__`, so this is a **compile-time hard failure**, not a runtime stub. **2. `sage_attention/sage_attn_launcher.cu`** — ptxas: ``` ptxas .../sage_attn_launcher.compute_70.ptx, line 6082; error : Feature 'cvt.pack' requires .target sm_72 or higher ... (20+ lines) ``` `cvt.pack` does not appear verbatim in the source — it is generated by the compiler from an intrinsic/packing macro on sm_72+, so it cannot be fixed by a simple grep-and-guard. **3. `ops/flash_decode.cu`** — missing header (submodule-dependent): ``` fatal error: flash.h: 没有那个文件或目录 ``` `flash_decode.cu` includes `flash.h` which lives in the `third_party/flash-attention` **submodule** (`csrc/flash_attn/src/flash.h`). A shallow clone without `--recurse-submodules` leaves it empty. ### What works on sm_70 - `rms_rope.cu` / `adaln.cu` compile and run correctly on V100 after building the remaining 24 sources with `COMFY_CUDA_ARCHS="70;75-real;80-real"`. - Measured on V100 (driver 580.173.02): `rms_rope_split_half` 0.093 ms/iter vs eager 0.612 ms/iter (~6.6x); `rms_adaln` 0.139 ms/iter vs eager 0.240 ms/iter (~1.7x), 1000 iters, numerics finite. ### Question for maintainers How would you prefer to support sm_70 builds? - **(a)** Guard the ldmatrix/cvt.pack code paths with `__CUDA_ARCH__ >= 750` / `>= 720` (with a real fallback or trap on sm_70 — silently compiling away the asm without a fallback would produce wrong results), or - **(b)** Conditionally exclude the 3 sources from `CUDA_SOURCES` when the arch list contains `70` (documenting that `awq_w4a16`/`sage_attn`/`flash_decode` need sm_72+/sm_75+), or - **(c)** keep them in the build and let sm_70 users carry a local patch. I have a working local build (the 3 sources excluded + bindings stubbed) with rms_rope/adaln verified on V100. I can share the patch or prepare a PR in whichever direction you prefer. Happy to do the `__CUDA_ARCH__` guard work on `mma.cuh` / `svdquant_utils.cuh` if you confirm that's the desired shape — just want to avoid guessing about kernel fallback semantics in shared attention headers.

bsakaguchi commented

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

Author

Owner

@Joevy-1112 commented on GitHub (Aug 30, 2026):

Reference implementation (working local build, V100-verified):

sm70-build.patch: CMakeLists (exclude the 3 sm_72+/sm_75+ sources) + dlpack_bindings.cpp (5 binding call sites stubbed to throw instead of linking against the excluded kernels)
typing-compat.patch: list[int] -> typing.List[int] in eager/na.py + eager/sol_attn.py (torch 2.5.1 infer_schema rejects lowercase generics; unrelated to sm_70)

gist: https://gist.github.com/Joevy-1112/4f7df791e25e4aa11baad295fbc82637

rms_rope_split_half / rms_adaln verified running on V100 with this build (6.6x / 1.7x vs eager). Happy to turn this into a PR in the shape you prefer.

@Joevy-1112 commented on GitHub (Aug 30, 2026): Reference implementation (working local build, V100-verified): - `sm70-build.patch`: CMakeLists (exclude the 3 sm_72+/sm_75+ sources) + dlpack_bindings.cpp (5 binding call sites stubbed to throw instead of linking against the excluded kernels) - `typing-compat.patch`: `list[int]` -> `typing.List[int]` in eager/na.py + eager/sol_attn.py (torch 2.5.1 `infer_schema` rejects lowercase generics; unrelated to sm_70) gist: https://gist.github.com/Joevy-1112/4f7df791e25e4aa11baad295fbc82637 rms_rope_split_half / rms_adaln verified running on V100 with this build (6.6x / 1.7x vs eager). Happy to turn this into a PR in the shape you prefer.

bsakaguchi referenced this issue

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

[PR #46] [MERGED] Add apply_rope_split_half for models with this type of rope. #104