kernels/hopper/tensor_core_wgmma/README.md

Hopper Tensor Core WGMMA Lane

This directory is reserved for the Hopper-native gap-closing kernel lane.

Current state:

• tensor_core_async is the owned WMMA guardrail lane.
• cuBLASLt is the control ceiling lane.
• tensor_core_wgmma is the next implementation lane intended to close the remaining device-side gap.

Minimum contract for the first implementation:

• target shape: 4096x4096x4096
• dtype: BF16 input, FP32 accumulation
• architecture: sm90
• correctness lane: 512x512x512 with reference enabled
• performance lane: 4096x4096x4096 with reference disabled

Design goals:

• use warpgroup MMA rather than per-warp WMMA
• move toward TMA-backed staging for the long-term feed path
• preserve a standalone harness so the lane can be benchmarked by kernel_lab
• keep the first fast path simple: no generalized epilogue work unless it is proven necessary

Promotion rule:

• only promote a tensor_core_wgmma variant if it beats the current owned SM90 baseline without correctness regressions