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

• URL: http://arxiv.org/abs/2608.10103v1
• Date: Mon, 10 Aug 2026 18:13:00 GMT
• Title: Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4
• Abstract summary: Core kernels rely on a low-level PTX pipeline built from asynchronous data movement.<n>Most application code accesses Cores indirectly through the WMMA C++ API.<n>This paper asks a focused, practical question: when does replacing WMMA with hand-written PTX actually pay off?
• Score: 0.7265954868708899
• License: http://arxiv.org/licenses/nonexclusive-distrib/1.0/
• Abstract: 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.

Related papers

• Realizing Native INT8 Compute for Diffusion Transformers on Consumer GPUs: A Fused INT8 GEMM Kernel for Ideogram 4.0 [0.08599681538174887]
  Post-training INT8 (W8A8) quantization of diffusion transformers is widely deployed as a speed optimization.<n>On consumer Ampere GPU it is frequently slower than the FP8 and NF4 alternatives it is meant to beat.<n>We close this gap with a single fused Triton INT8 GEMM.<n>INT8 thus goes from the slowest variant to the fastest.
  arXiv  Detail & Related papers  (2026-06-12T16:19:03Z)
• APEX4: Efficient Pure W4A4 LLM Inference via Intra-SM Compute Rebalancing [7.218138615728212]
  W4A4 quantization promises full utilization of INT4 Cores, yet group dequantization overhead on Cores has driven existing systems to mixed-precision fallbacks.<n>We present the first systematic study of how intra-SM compute balance governs this bottleneck.
  arXiv  Detail & Related papers  (2026-06-07T18:01:55Z)
• FairyFuse: Multiplication-Free LLM Inference on CPUs via Fused Ternary Kernels [2.339485034137561]
  We present FairyFuse, an inference system that enables multiplication-free execution on commodity CPUs.<n>Roofline analysis shows that 16x weight compression shifts memory-bound GEMV toward the compute regime on bandwidth-limited CPUs.<n>End-to-end, FairyFuse achieves 32.4 tokens per second on a single Intel Xeon 8558P, outperforming llama.
  arXiv  Detail & Related papers  (2026-04-22T02:42:10Z)
• Scaling DoRA: High-Rank Adaptation via Factored Norms and Fused Kernels [83.99688944263843]
  Weight-De Low-Rank Adaptation (DoRA) extends LoRA by decoupling weight magnitude from direction.<n>At d_in = 8192 and rank r = 384, a single module's norm requires about 512 MB of transient working memory in bf16.<n>A factored norm decomposes the squared norm into base, cross, and Gram terms computable through O(d_out r + r2) intermediates, eliminating the dense product.
  arXiv  Detail & Related papers  (2026-03-23T17:57:24Z)
• INT v.s. FP: A Comprehensive Study of Fine-Grained Low-bit Quantization Formats [51.72056104795248]
  Modern AI hardware, such as Nvidia's Blackwell architecture, is increasingly embracing low-precision floating-point (FP) formats.<n>This paper systematically investigates the trade-offs between FP and integer (INT) formats.<n>We reveal a critical performance crossover: while FP excels in coarse-grained quantization, the comparison at fine-grained (block-wise) levels is more nuanced.
  arXiv  Detail & Related papers  (2025-10-29T15:11:53Z)
• Speeding Up MACE: Low-Precision Tricks for Equivarient Force Fields [51.95157731126864]
  Machine-learning force fields can deliver accurate molecular dynamics (MD) at high computational cost.<n>This thesis aims to make MACE cheaper and faster by identifying computational bottlenecks and evaluating low-precision execution policies.
  arXiv  Detail & Related papers  (2025-10-23T14:02:34Z)
• Bridging the Gap Between Promise and Performance for Microscaling FP4 Quantization [77.67818998672516]
  We present the first comprehensive study of MXFP4 and NVFP4 for post-training quantization.<n>We introduce Micro-Rotated-GPTQ (MR-GPTQ), a variant of the classic GPTQ quantization algorithm.<n>We show that MR-GPTQ matches or outperforms state-of-the-art accuracy.
  arXiv  Detail & Related papers  (2025-09-27T09:22:21Z)
• APT-LLM: Exploiting Arbitrary-Precision Tensor Core Computing for LLM Acceleration [5.075697428779204]
  Large language models (LLMs) have revolutionized AI applications, yet their enormous computational demands severely limit deployment and real-time performance.<n>This is primarily due to the limited support for the GPU Cores, inefficient memory management, and inflexible kernel optimizations.<n>We propose a comprehensive acceleration scheme for arbitrary precision LLMs, namely APT-LLM.
  arXiv  Detail & Related papers  (2025-08-26T14:48:29Z)
• MicroMix: Efficient Mixed-Precision Quantization with Microscaling Formats for Large Language Models [3.305409455598179]
  Quantization significantly accelerates inference in large language models (LLMs)<n>Recent advances in weight-activation quantization have primarily focused on mapping both weights and activations to the INT4 format.<n>We propose MicroMix, a co-designed mixed-precision quantization algorithm and matrix multiplication kernel based on Microscaling (MX) data formats.
  arXiv  Detail & Related papers  (2025-08-04T12:22:39Z)
• MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models [58.3342517278868]
  This paper describes the design of Mixed-precision AutoRegressive LINear kernels. It shows that batchsizes up to 16-32 can be supported with close to maximum ($4times$) quantization speedup. MarLIN accomplishes this via a combination of techniques, such as asynchronous memory access, complex task scheduling and pipelining.
  arXiv  Detail & Related papers  (2024-08-21T16:10:41Z)
• vTensor: Flexible Virtual Tensor Management for Efficient LLM Serving [53.972175896814505]
  Large Language Models (LLMs) are widely used across various domains, processing millions of daily requests. Large Language Models (LLMs) are widely used across various domains, processing millions of daily requests.
  arXiv  Detail & Related papers  (2024-07-22T14:37:58Z)

This list is automatically generated from the titles and abstracts of the papers in this site.

This site does not guarantee the quality of this site (including all information) and is not responsible for any consequences.