Fix deep-ep sm_80 fatbin build: real cuda_fp8.h + unconditional smem_size in patch

[ETOgaosion]

Fix deep-ep x86_64 sm_80 fatbin build (two patch bugs exposed by the #10 full rebuild)

After #10 merged, the fail-closed manifest logic forced a full deep-ep rebuild and the x86_64 wheel failed compiling csrc/kernels/intranode.cu for the sm_80 gencode. This path was never compiled before #10: the old skip logic skipped the x86_64 deep-ep build because an sm_90-only wheel already existed in the release, so the sm_80 cubin added by #8 was never actually exercised in CI. (aarch64 9.0;10.0 built green.)

Both bugs are in the build-time-only patch ci/patches/enable_deep_ep_sm80.py; the DeepEP submodule is untouched.

Fix 1 — fake FP8 typedefs collide with the real cuda_fp8.h (commit c97733c)

Errors (71) in the sm_80 device pass:

cuda_fp8.h(185): error: invalid redeclaration of type name "__nv_fp8_interpretation_t" (declared at configs.cuh)
cuda_fp8.h(478): error: invalid redeclaration of type name "__nv_fp8x4_e4m3" (declared at configs.cuh)
... cuda_fp8.hpp: argument of type "int" is incompatible with parameter of type "__nv_fp8_interpretation_t"

The sm_80-only DISABLE_SM90_FEATURES define makes upstream configs.cuh take its Ampere #else branch, which re-declares the FP8 API with fake typedefs (typedef int __nv_fp8_interpretation_t;, typedef uint8_t __nv_fp8_storage_t;, #define __NV_E4M3 0). In the same TU, torch/nvshmem headers transitively pull in the toolkit's real <cuda_fp8.h> → redeclaration / type-mismatch.

Fix: configs.cuh now unconditionally includes the real <cuda_fp8.h>. No sm_80 compilation unit references any FP8 symbol — the only FP8 call sites are in internode_ll.cu, whose sm_80 cubin is compiled out by this same patch (verified by grepping every sm_80-compiled TU; calculate_fp8_scales is pure float math). Host / sm_90 / sm_100 passes already used the real header.

Fix 2 — smem_size undefined at the launch macro sites (commit e65a9d1)

First branch CI run (33861703220) then failed with exactly:

intranode.cu(521): error: identifier "smem_size" is undefined
intranode.cu(941): error: identifier "smem_size" is undefined

(and zero FP8 errors — fix 1 confirmed). intranode.cu declares smem_size inside #ifndef DISABLE_SM90_FEATURES (dispatch + combine launch wrappers). Upstream's Ampere build defines SET_SHARED_MEMORY_FOR_TMA(kernel) to void() under that guard, so the missing constant never mattered; our runtime-dispatched macro (edit 3 of the patch) always expands to code referencing smem_size (the cudaFuncSetAttribute call is a runtime Hopper+ check). Fix: move both smem_size constants out of the guard; the patch fails loudly unless exactly the two expected guarded declarations are present.

Other DISABLE_SM90_FEATURES sites in sm_80-compiled TUs were audited: the 6 remaining guards in intranode.cu and the 2 in utils.cuh are device-kernel TMA/elect paths with upstream's own Ampere #else fallbacks; runtime.cu/layout.cu have no guards and no SET_SHARED_MEMORY_FOR_TMA calls; internode.cu/internode_ll.cu/pcie.cu are compiled out for sm_80 entirely.

Verification

• Patch applied twice against a clean DeepEP checkout copy: second run is a no-op (idempotent); resulting tree has one unconditional #include <cuda_fp8.h>, zero fake FP8 typedefs, and 6 (was 8) SM90 guards in intranode.cu.
• python -m unittest test_cuda_archs test_generate_matrix — 43/43 pass.
• Branch CI run 33861703220 (x86_64, workflow_dispatch = build only, no publish): FP8 errors eliminated; only the smem_size sites remained.
• Branch CI run 33862535136 (x86_64, same dispatch): in progress — must compile clean and pass the "Verify wheel CUDA arch coverage" gate (sm_80 + sm_90 + sm_100 in the fatbin).
• The patch's sha256 feeds the build_inputs fingerprint, so the deep-ep skip decision correctly flips to rebuild after merge.