feat(kda): add persistent SM90 bwd intra CUDA kernel

[zheyang0825]

📌 Description

Integrates the optimized persistent CUDA C++ KDA intra-chunk backward kernel under csrc/kda/sm90/bwd and dispatches supported safe_gate=True, BF16, K=128, chunk-size-64 workloads to it. Unsupported inputs continue to use the existing Triton implementation.

The kernel uses warp-level mma.sync.m16n8k8. It is organized under SM90 to match the current repository layout and is compiled for SM90, SM100, and SM103.

Also adds:

• a low-level Python entry point for direct validation and profiling
• correctness, determinism, dense-batch, ragged-varlen, dispatch/fallback, device, and launch-boundary tests
• benchmark and NCU profiling scripts
• benchmark documentation for FLA v0.5.0 and v0.4.2

🔍 Related Issues

None.

🚀 Pull Request Checklist

✅ Pre-commit Checks

• I have installed pre-commit.
• I have installed the hooks.
• I have run pre-commit run --all-files and fixed all reported issues.

🧪 Tests

• Tests have been added or updated as needed.
• All targeted tests are passing.

Validation on NVIDIA L20X (SM90), CUDA 12.9, PyTorch 2.9.1+cu129:

python -m pytest tests/test_kda_sm90_bwd_intra.py -v
9 passed in 13.94s

compute-sanitizer --tool memcheck ... benchmarks/ncu_profile_kda_bwd_intra.py
ERROR SUMMARY: 0 errors

⚡ Performance

Same-machine median latency, 25 ms warmup and 100 ms measurement windows:

• FLA v0.5.0: 1.699x geometric-mean speedup
• FLA v0.4.2: 1.612x geometric-mean speedup

The full eight-shape tables and reproduction steps are in BENCHMARK_KDA_BWD_INTRA_SM90.md.

Reviewer Notes

Three review/fix passes were completed before opening this PR:

1. API/dispatch correctness: hardened CUDA/device selection and added CPU rejection coverage.
2. CUDA launch/runtime safety: checked runtime calls and kernel launch failures; validated the ragged path with compute-sanitizer.
3. Boundaries/maintainability/docs: rejected empty/overflowing launch grids, removed dead scheduler state, and corrected Triton benchmark timing terminology.

The optimized path intentionally supports only safe_gate=True, BF16 Q/K/beta, FP32 gate/gradient inputs, K=128, and chunk size 64. Other cases retain the Triton fallback.