# Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4

arXiv: https://arxiv.org/abs/2608.10103 · HTML: https://mlsanity.com/paper/2608.10103/

## Key facts

- No citing papers in the map yet (data through 14 Oct 2026).
- Builds on 5 of its 15 references that are in the map.
- First posted on arXiv 10 Aug 2026.

## Builds on (most cited first, 5 of 5)

- [GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers](https://mlsanity.com/paper/2210.17323/) (2022)
- [LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale](https://mlsanity.com/paper/2208.07339/) (2022)
- [NVIDIA Tensor Core Programmability, Performance & Precision](https://mlsanity.com/paper/1803.04014/) (2018)
- [Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking](https://mlsanity.com/paper/1804.06826/) (2018)
- [Dissecting Tensor Cores via Microbenchmarks: Latency, Throughput and Numeric Behaviors](https://mlsanity.com/paper/2206.02874/) (2022)

## Abstract (arXiv)

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.