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

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

Start with the CUDA_SOURCES selection and the failing files ops/awq_w4a16.cu, sage_attention/sage_attn_launcher.cu, and ops/flash_decode.cu; inspect the shared asm in sage_attention/mma.cuh and ops/svdquant_utils.cuh. Check the third_party/flash-attention submodule layout and build with sm_70 enabled. Done means the maintainers' chosen fallback or source-exclusion approach produces a successful sm_70 build without silently incorrect kernels.

Written by the indexing model from the issue text.