NVIDIA L4 Tensor Core study reveals hand-written PTX speedups

Researchers Matt J. Borowski and Blazej Osinski published a study examining the performance tradeoffs of hand-written PTX Tensor Core kernels versus NVIDIA's WMMA C++ API on the L4 GPU (Ada architecture, SM89). Their controlled single-GPU tests compared double-buffered WMMA baselines against hand-written PTX GEMM kernels across FP16, INT8, and INT4 precisions for matrix sizes from N=512 to N=8192.

The results showed no speedup for FP16 due to instruction-packing overhead offsetting gains. However, INT8 kernels achieved 1.4x-1.8x speedups from lower instruction counts and improved global-memory coalescing, while INT4 kernels saw 2.9x-4.3x gains by leveraging native mma.sync.m16n8k64.s4 execution to avoid WMMA's software-emulated path. At N=8192, INT8 reached 34.4x and INT4 hit 98.7x relative to the FP16 WMMA baseline. Performance correlated more with memory-system behavior (global-load coalescing, DRAM-active cycles) than Tensor Core utilization or occupancy.