Port the full FlashAttention-3 Hopper forward kernel set (arXiv:2407.08608):
head dims 64/128/256/512, fp16 + bf16, dense and GQA-packed, causal/local/
varlen, with the warp-specialized producer/consumer pipeline, TMA async copy
and wgmma of the reference hopper/ kernels.
- Kernels compile ONLY under the `cuda` feature, cross-compiled for sm_90a
(nvcc -gencode arch=compute_90a,code=sm_90a) on any CUDA host. Without the
feature the crate is a pure-Rust stub, so ROCm/Metal/Vulkan/CPU/non-Hopper
CUDA builds are unaffected. Un-excluded as a stub-by-default workspace member.
- cuda_compat.h: force-included shim restoring the unversioned
PFN_cuTensorMapEncodeTiled/Im2col typedefs CUTLASS 3.5 needs under CUDA 13;
touches no kernel. All 49 sources compile clean for sm_90a on nvcc 13.0.
- Numerical-equivalence harness: FA3 vs f32 CPU oracle within a scale-relative
bound across head dims/dtypes/causal/GQA/varlen, #[ignore]d for Hopper.
- publish.yml publishes the crate after hanzo-ml.
fp8 (e4m3) left as a documented hook. Perf UNVALIDATED (no Hopper silicon).
Hanzo Dev
561164956e