feat: add wgmma f16acc minimal cpp path
Summary
This PR adds a minimal Hopper WGMMA F16-acc HGEMM TN path for C++ benchmark flow only.
Scope
- Add F16 accumulator WGMMA TN kernel implementation.
- Add Makefile build entry for standalone C++ binary.
- Keep PR intentionally minimal for easier review.
Included Files
- kernels/hgemm/wgmma/hgemm_wgmma_fp16acc_stages_tn.cu
- /kernels/hgemm/makefile
Not Included
- No PyTorch binding changes in this PR.
- No Python API exposure changes in this PR.
- No default kernel selection changes.
Link issue
#416
Kernel Notes
- Layout: TN.
- WGMMA: m64n128k16.
- Pipeline: TMA + warp specialization + multi-stage.
- Target: SM90 (Hopper), CUDA 12.x.
Validation
- Local standalone build passes.
- Local binary run passes.
- Correctness check on representative small square cases reports max error = 0.
- Local performance run is stable; peak observed around 572.8 TFLOPS in current environment.
Why Minimal First
- Keep review scope focused on core CUDA path and build integration.
- Pybind integration will be submitted in a follow-up PR after this kernel path is reviewed.
Checklist
- [x] Build success for standalone C++ path
- [x] Runtime execution verified
- [x] Correctness checked against reference flow
- [x] Minimal scope maintained (no pybind changes)
合并状态:未合并 关闭于 2026-04-07 5 条评论