Search papers, labs, and topics across Lattice.
This study evaluates the performance benefits of hand-written PTX Tensor-Core GEMM kernels compared to the WMMA C++ API on an NVIDIA L4 GPU across multiple precisions (FP16, INT8, INT4) and matrix sizes. The findings reveal that while hand-written PTX kernels do not improve FP16 performance due to operand-packing overhead, they achieve significant speedups of 1.4x-1.8x for INT8 and 2.9x-4.3x for INT4, particularly at larger matrix sizes where memory-system behavior is more critical than Tensor Core utilization. This research highlights the specific scenarios where the complexity of hand-written PTX is beneficial, providing valuable insights for optimizing high-performance computing applications.
Hand-written PTX kernels can outperform WMMA by up to 98.7x for INT4 operations, revealing critical insights into when low-level optimizations are worth the added complexity.
High-performance Tensor Core kernels rely on a low-level PTX pipeline built from asynchronous data movement with cp.async, warp-level matrix loads with ldmatrix, and matrix multiply-accumulate operations with mma.sync. However, most application code accesses Tensor Cores indirectly through the WMMA C++ API. This paper asks a focused, practical question: when does replacing WMMA with hand-written PTX actually pay off? To answer this question, we conduct a controlled, single-GPU study on an NVIDIA L4 GPU (Ada, SM89), comparing double-buffered WMMA baselines with a family of hand-written PTX GEMM kernels across FP16, INT8, and INT4 arithmetic and square problem sizes from $N=512$ to $N=8192$. Every kernel is profiled with Nsight Compute across the full metric set, and PTX speedups are reported relative to the corresponding same-precision WMMA baseline. Hand-written PTX provides no end-to-end speedup for FP16, because its instruction-level gains are offset by operand-packing overhead. In contrast, the PTX kernels achieve consistent speedups of 1.4x-1.8x for INT8, driven primarily by lower instruction counts and better global-memory coalescing, and 2.9x-4.3x for INT4, where native mma.sync.m16n8k64.s4 execution avoids the software-emulated sequence used by the WMMA path. Relative to the FP16 WMMA baseline, the best quantized kernels reach 34.4x (INT8) and 98.7x (INT4) at $N=8192$. Across these experiments, occupancy is a poor predictor of throughput. For large matrices, performance instead tracks memory-system behavior -- particularly global-load coalescing and DRAM-active cycles -- more closely than Tensor Core utilization. These results identify the precisions and operating regimes in which the additional complexity of hand-written PTX is justified.
