ITADN

feat: add wgmma f16acc minimal cpp path

#417Pull Requestlnxtree 创建于 2026-04-07
L
lnxtreecommented
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 条评论