Paper deep dive
Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4
Matt J. Borowski, Blazej Osinski
Intelligence
Status: succeeded | Model: Gemma-4-26B-A4B | Prompt: intel-v1 | Confidence: 90%
Last extracted: 8/13/2026, 4:29:22 AM
Summary
This study evaluates the performance of hand-written PTX Tensor Core GEMM kernels versus the WMMA C++ API on an NVIDIA L4 GPU. It finds that hand-written PTX offers no speedup for FP16 due to packing overhead but provides significant speedups for INT8 (1.4x-1.8x) and INT4 (2.9x-4.3x) by reducing instruction counts and avoiding software emulation. Performance at large matrix sizes is driven by memory coalescing and DRAM latency rather than Tensor Core utilization or occupancy.
Entities (12)
Relation Signals (10)
NVIDIA L4 → runs → PTX
confidence 95% · controlled, single-GPU study on an NVIDIA L4 GPU ... hand-written PTX GEMM kernels
Performance → correlateswith → global-load coalescing
confidence 90% · performance instead tracks memory-system behavior -- particularly global-load coalescing
PTX → nospeedupfor → FP16
confidence 90% · Hand-written PTX provides no end-to-end speedup for FP16
PTX → outperforms → WMMA
confidence 90% · PTX kernels achieve consistent speedups ... for INT8 ... and INT4
Performance → poorpredictorof → Occupancy
confidence 90% · occupancy is a poor predictor of throughput
Nsight Compute → profiles → PTX
confidence 90% · Every kernel is profiled with Nsight Compute
int4_ptx_mma_k64 → achievesspeedupfor → INT4
confidence 85% · int4_ptx_mma_k64 ... 4.27x speedup vs INT4-WMMA at N=8192
int8_ptx_mma_k32 → achievesspeedupfor → INT8
confidence 85% · int8_ptx_mma_k32 ... 1.79x speedup vs INT8-WMMA at N=8192
PTX → avoidsemulationfor →
Cypher Suggestions (0)
No Cypher suggestions yet.
Abstract
Abstract:High-performance Tensor Core kernels rely on a low-level PTX pipeline built from asynchronous data movement with this http URL, warp-level matrix loads with ldmatrix, and matrix multiply-accumulate operations with this http URL. 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.
Tags
Links
- Source: https://arxiv.org/abs/2608.10103v1
- Canonical: https://arxiv.org/abs/2608.10103v1
Trouble viewing inline? Open PDF directly →
Full Text
30,554 characters extracted from source content.
Expand or collapse full text
Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4 Matt J. Borowski, MSc Blazej Osinski, PhD 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=512N=512 to N=8192N=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.4×1.4×–1.8×1.8× for INT8, driven primarily by lower instruction counts and better global-memory coalescing, and 2.9×2.9×–4.3×4.3× 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.4×34.4× (INT8) and 98.7×98.7× (INT4) at N=8192N=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. Index Terms—GPU computing, Tensor Cores, PTX, GEMM, quantization, INT4, INT8, mma.sync, cp.async, Nsight Compute, NVIDIA Ada, LLM inference. I Introduction General matrix multiplication (GEMM) dominates the compute cost of modern deep-learning inference, and on NVIDIA GPUs the fastest path for GEMM is the Tensor Core. Since the Volta generation, Tensor Cores have been exposed to developers at two very different levels of abstraction. The high-level WMMA (warp-matrix-multiply-accumulate) C++ API is convenient and portable, but it fixes fragment shapes, load patterns, and accumulator layouts, and it hides the underlying instruction stream. The low-level PTX path—cp.async for asynchronous global-to-shared copies, ldmatrix for cooperative shared-to-register fragment loads, and mma.sync for the matrix multiply itself—gives the programmer direct control over tile shape, staging depth, and register mapping, at the cost of considerable complexity. A widely held assumption is that hand-written PTX is uniformly faster than the WMMA API. In practice the picture is far more nuanced: whether PTX helps depends strongly on the numeric precision, the problem size, and, above all, on whether the kernel is compute-bound or memory-bound. This paper provides a controlled, reproducible answer for one concrete platform—the NVIDIA L4 inference GPU (Ada, compute capability 8.9)—by profiling a coherent family of kernels under identical conditions. We make the following contributions. First, we implement a matched set of PTX GEMM kernels across FP16, INT8, and INT4 that share a common tiling strategy and differ only in the instruction-level choices under study (Section I-A). Second, we benchmark each kernel against a double-buffered WMMA baseline of the same precision, over N=512 to N=8192, using Nsight Compute for full hardware-counter attribution (Section IV). Third, we isolate the mechanisms behind each result: instruction count and coalescing for INT8, avoidance of WMMA INT4 software emulation for INT4, and the compute-to-memory-bound transition that flattens all FP16 differences at large N. Finally, we perform a focused ablation over the best INT4 kernel family—loader split, cache-eviction policy, and operand layout—to show that the chosen configuration sits in a genuine local optimum (see Section IV-D). As part of these contributions, the complete source code of every kernel studied here—together with the build system, benchmarking harness, and all Nsight Compute profiles—is openly available as the article’s reproducibility artifact (see the artifact availability statement at the end of Section VII). Taken together, the results yield a practical rule of thumb: PTX is worth the engineering effort precisely when it removes instruction overhead that the WMMA path cannot avoid—most dramatically for INT4, where WMMA emulates the operation in software—and is not worth it when the kernel is already bandwidth-bound or when the compiler’s WMMA lowering is already near-optimal, as with FP16. These findings are directly relevant to serving quantized open-weight LLMs on a single commodity inference GPU. I Background I-A Tensor Cores, WMMA, and PTX mma.sync A Tensor Core performs a small matrix multiply-accumulate, D=AB+CD=AB+C, through a single warp-collective instruction. The WMMA API exposes this operation through opaque fragment types and the load_matrix_sync, mma_sync, and store_matrix_sync calls, using the m16n16k16 shape for FP16 and INT8. The PTX layer instead exposes the underlying mma.sync.aligned instructions with explicit shapes such as m16n8k16 and m16n8k64, together with ldmatrix for loading the register fragments held by each lane. Working at the PTX level lets a kernel choose the K-tile width, decompose one WMMA step into several narrower MMA instructions, and control exactly how many instructions are issued per unit of arithmetic work. I-B Asynchronous staging: cp.async and pipelining All kernels studied here overlap DRAM→ -memory prefetch with tile-level compute. The cp.async instruction copies a 16-byte vector from global to shared memory without occupying registers or blocking the issuing warp; commit_group and wait_group bracket the in-flight copies. A two-stage (double-buffered) pipeline keeps one shared-memory buffer filling while the other is consumed by ldmatrix + mma.sync; a three-stage variant (wait_group 1) keeps an additional buffer in flight to hide longer L2 latency at the cost of extra shared-memory pressure. I-C Quantization and the memory wall Lowering operand precision from FP16 to INT8 or INT4 shrinks the bytes moved per multiply-accumulate by 2× and 4× respectively, and shrinks the working set proportionally. For inference-scale GEMM this is decisive: once operands are narrow enough for the active tiles to reside in L2, the kernel escapes the DRAM bottleneck that dominates FP16 at large N. The experiments below quantify this directly and show that quantization reduces memory-bandwidth consumption by roughly 5×–100× while maintaining or improving cache locality. I-D Platform: NVIDIA L4 (Ada, SM89) All measurements were taken on an NVIDIA L4, an Ada-generation datacenter inference GPU (compute capability 8.9) with fourth-generation Tensor Cores supporting FP16, INT8, and INT4 MMA. The optimization patterns transfer to other Tensor-Core GPUs such as A100 (SM80) and H100 (SM90), though the exact gains are hardware- and compiler-dependent; only the SM target passed to the compiler changes. I Methodology I-A Kernel design space Every kernel implements a tiled GEMM with the warp- and block-level tiling parameters fixed by a shared configuration header (WMMA_M = WMMA_N = 16, four warp-tiles in X and two in Y, eight warps per block). Within each precision family, exactly one WMMA-API kernel serves as the baseline and the PTX variants change one dimension of the design at a time: the SRAM→ load path (ldmatrix width vs. manual scalar packing), the mma.sync tile shape (K = 8/16/32/64), the accumulator type, and the pipeline depth (two- vs. three-stage). This one-variable-at-a-time discipline is what allows each observed speedup or slowdown to be attributed to a specific mechanism. Listing 1 shows the core of the fastest large-N INT4 kernel (int4_ptx_mma_k64): a double-buffered main loop that issues cp.async prefetches for the next K-tile, performs two m16n8k64 INT4 MMA calls on the current fragments, then rotates buffers. Listing 2 isolates the three PTX primitives—ldmatrix.x4, ldmatrix.x2, and mma.sync.m16n8k64.s4—that distinguish the native path from WMMA. Listing 3 shows how the three-stage variant keeps one additional copy group in flight with wait_group 1. ⬇ // int4_ptx_mma_k64 -- double-buffered main loop (K-tile = 64 int4) for (int k = WMMA_K; k < K; k += WMMA_K) const int next = 1 - buf; // Stage next A/B tile: 16-byte cp.async vectors, .ca cache policy for (int i = lane_id; i < (WMMA_M * K_BYTES) / 16; i += THREADS_PER_WARP) const int row = (i * 16) / K_BYTES; const int byte_col = (i * 16) % K_BYTES; const unsigned dst = __cvta_generic_to_shared( &As[next][warp_id][row][byte_col]); asm volatile( "cp.async.ca.shared.global [%0], [%1], 16;" :: "r"(dst), "l"(&A_b[(tile_row + row) * (K / 2) + byte_col + (k / 2)])); // ... identical staging for B ... asm volatile("cp.async.commit_group;"); mma_int4_k64(rc0, ra, rb0); // compute on current fragments mma_int4_k64(rc1, ra, rb1); asm volatile("cp.async.wait_group 0;"); __syncthreads(); buf = next; // rotate double buffer ldmatrix_a_k64(ra, As[buf][warp_id], lane_id); ldmatrix_b_k64(rb0, Bs[buf][warp_id], lane_id, 0); ldmatrix_b_k64(rb1, Bs[buf][warp_id], lane_id, 8); Listing 1: Overlap of cp.async prefetch with mma.sync compute in the double-buffered INT4 k64 kernel. ⬇ // (a) A fragment: ldmatrix.x4 loads a 16x64 int4 tile (as 16x16 b16) asm volatile( "ldmatrix.sync.aligned.m8n8.x4.shared.b16 " "%0,%1,%2,%3, [%4];" : "=r"(t0), "=r"(t1), "=r"(t2), "=r"(t3) : "r"(addr)); // (b) B fragment: ldmatrix.x2 per n8-half (8x64 int4 tile) asm volatile( "ldmatrix.sync.aligned.m8n8.x2.shared.b16 %0,%1, [%2];" : "=r"(rb[0]), "=r"(rb[1]) : "r"(addr)); // (c) Native INT4 Tensor-Core multiply-accumulate (no WMMA emulation) asm volatile( "mma.sync.aligned.m16n8k64.row.col.s32.s4.s4.s32 " "%0,%1,%2,%3, %4,%5,%6,%7, %8,%9, %0,%1,%2,%3;" : "+r"(c0), "+r"(c1), "+r"(c2), "+r"(c3) : "r"(ra[0]), "r"(ra[1]), "r"(ra[2]), "r"(ra[3]), "r"(rb[0]), "r"(rb[1])); Listing 2: The three PTX primitives that define the native INT4 path: cooperative ldmatrix fragment loads and the m16n8k64.s4 MMA. I-B Benchmark and profiling setup Each kernel is compiled for CUDA_ARCH=89 and run over square sizes N ∈ 512, 1024, 2048, 4096, 8192. Wall-clock timings are collected from repeated launches; detailed hardware counters are collected with Nsight Compute (ncu –set full), which reports throughput, occupancy, cache hit rates, instruction counts, coalescing efficiency, and warp-state statistics. Speedup is always defined relative to the same-precision WMMA baseline (fp16_wmma, int8_wmma, int4_wmma); a separate summary also expresses every kernel relative to the FP16 WMMA baseline to capture the end-to-end value of quantization plus PTX. The study is organized as four runs: Run 1 covers the FP16 family, Run 2 the INT8 family, Run 3 the INT4 base kernels, and Run 4 an ablation over the best INT4 (k64) family. Table I summarizes the outcome of each run before the detailed analysis. TABLE I: AT-A-GLANCE RESULTS: BEST KERNEL AND DOMINANT MECHANISM PER RUN Run Best kernel(s) Speedup vs. same-precision WMMA Key finding Run 1 ⋅· FP16 fp16_wmma 1.0× PTX does not improve wall time; local gains are offset by extra instruction / packing overhead. Run 2 ⋅· INT8 int8_ptx_mma_k32 1.4×–1.8× k32 wins via fewer executed instructions and near-perfect global-load coalescing. Run 3 ⋅· INT4 int4_ptx_3stage (small N); int4_ptx_mma_k64 (large N) 2.9×–4.3× Both avoid WMMA INT4 software emulation; 3stage loses L1 locality as N grows, k64 keeps it. Run 4 ⋅· INT4 k64 int4_ptx_mma_k64_ x4_x2nontrans_ca 2.9×–4.3× Loader split, cache policy and layout do not beat the baseline; transposed B collapses coalescing. TABLE I: HEADLINE PER-KERNEL METRICS AT N = 8192 Kernel Prec. Duration (ms) AI TOPS DRAM % L2 hit % L1/TEX % fp16_wmma fp16 16500 2.73 0.067 54.0 65.9 26.7 int8_wmma int8 858.3 5.46 1.281 52.2 81.9 33.6 int8_ptx_mma_k32 int8 480.1 5.46 2.290 52.5 65.6 61.5 int4_wmma int4 713.0 10.92 1.542 0.88 99.7 80.3 int4_ptx_3stage int4 197.6 10.92 5.564 3.16 99.8 14.5 int4_ptx_mma_k64 int4 167.1 10.92 6.581 3.73 99.5 61.6 TABLE I: BEST-PER-PRECISION WALL-CLOCK DURATION (MS) VS. PROBLEM SIZE N Kernel Role 512 1024 2048 4096 8192 fp16_wmma Base & Optimal 0.184 1.430 11.850 102.350 16500.000 int8_wmma Base 0.152 1.110 8.550 68.470 858.300 int8_ptx_mma_k32 Optimal 0.109 0.726 5.410 42.000 480.060 int4_wmma Base 0.195 1.450 11.400 90.110 712.960 int4_ptx_3stage Optimal (small N) 0.056 0.382 2.910 24.260 197.630 int4_ptx_mma_k64 Optimal (large N) 0.068 0.395 2.810 21.520 167.100 TABLE IV: SPEEDUP VS. THE FP16 WMMA BASELINE (HIGHER IS FASTER) Kernel Role 512 1024 2048 4096 8192 int8_wmma Base 1.2× 1.3× 1.4× 1.5× 19.2× int8_ptx_mma_k32 Optimal 1.7× 2.0× 2.2× 2.4× 34.4× int4_wmma Base 0.9× 1.0× 1.0× 1.1× 23.1× int4_ptx_3stage Optimal (small N) 3.3× 3.7× 4.1× 4.2× 83.5× int4_ptx_mma_k64 Optimal (large N) 2.7× 3.6× 4.2× 4.8× 98.7× IV Results and Analysis We report each run in turn. Throughout, positive speedup means faster than the relevant WMMA baseline. Two regimes recur: a compute-bound regime at small-to-medium N where instruction efficiency dominates, and a memory-bound regime at large N where bandwidth and coalescing dominate. The transition between them is sharpest for FP16, whose large working set overflows L2 between N=4096 and N=8192. IV-A Run 1 — FP16: PTX does not beat the compiler Six FP16 kernels were profiled. Across all sizes, no hand-written PTX variant is faster than the compiler-optimized WMMA baseline in wall time; measured ratios stay within roughly ± 5% (Table V). The reason is visible in the counters: in the compute-bound regime (N ≤ 4096) all kernels sustain 1,280–1,500 GFLOPS, and although fp16_ptx_manual_pack reaches the highest SM throughput (~57%) it also carries the most packing overhead, costing up to +22% at N=512. Once the ~537 MB combined working set at N=8192 overflows L2, every kernel stalls on DRAM: L1/TEX throughput collapses from ~90% to ~38%, GFLOPS roughly halve, and all differences shrink below 1%. The instruction mix becomes irrelevant because bandwidth is the sole bottleneck (Fig. 1). A secondary observation concerns the FP16-accumulator variant: using FP16 rather than FP32 accumulators halves the accumulator register count and lifts achieved occupancy to 70–82% (vs. 56–66% for the others), yet GFLOPS do not increase. This is direct evidence that occupancy alone does not drive throughput when the kernel is not limited by a warp-count shortage—a theme that recurs in Run 2. Figure 1: FP16 (Run 1): GFLOPS, SM throughput, achieved occupancy, and L1/TEX throughput vs. N. The L2 capacity cliff between N=4096 and N=8192 separates the compute-bound and memory-bound regimes. TABLE V: RUN 1 — FP16 RAW DURATIONS (MS) VS. N Kernel 512 1024 2048 4096 8192 fp16_wmma (baseline) 680.155 684.000 813.093 1539.726 14098.618 fp16_ptx_mma 710.899 698.062 782.720 1553.984 14150.838 fp16_ptx_k8 685.578 724.600 801.493 1540.406 14135.938 fp16_ptx_fp16acc 664.341 700.616 789.057 1555.479 14553.427 fp16_ptx_3stage 689.617 683.957 794.762 1577.620 14828.855 fp16_ptx_manual_pack 673.823 699.585 789.764 1574.092 14312.770 TABLE VI: RUN 1 — FP16 KERNEL DESIGN MATRIX Kernel SRAM→ mma shape Acc Pipeline Notes fp16_wmma wmma::load_matrix _sync m16n16k16 f32 2-stage WMMA baseline; no explicit PTX fp16_ptx_mma ldmatrix.x4 / .x2.trans m16n8k16 ×2 f32 2-stage first pure-PTX kernel fp16_ptx_k8 ldmatrix.x2 / .x1.trans m16n8k8 ×4 f32 2-stage narrower K; 4 MMA per K-step fp16_ptx_fp16acc ldmatrix.x4 / .x2.trans m16n8k16 ×2 f16 2-stage half the accumulator registers fp16_ptx_3stage ldmatrix.x4 / .x2.trans m16n8k16 ×2 f32 3-stage wait_group 1; extra SRAM buffer fp16_ptx_manual _pack scalar ld.shared + mov.b32 m16n8k16 ×2 f32 2-stage no ldmatrix; exposes packing cost IV-B Run 2 — INT8: k32 wins on instructions and coalescing Six INT8 kernels were profiled (Tables VII–VIII). int8_ptx_mma_k32 is the fastest at every size, from 23% faster than int8_wmma at N=512 to 43% faster at N=8192. Its advantage is rooted in instruction count: by decomposing each K=32 tile step into two tightly unrolled m16n8k16 MMA calls it executes 25–43% fewer instructions than WMMA, and its global-load coalescing is near-perfect—only 0.4% wasted sectors at N=8192, versus roughly 50% for every other kernel. The other variants illuminate the mechanism by contrast. int8_ptx_mma_k16 is slower than WMMA at small–medium sizes (uncoalesced loads waste 31 of every 32 bytes per sector) but recovers at N=8192 once a 76% L1 hit rate absorbs the excess traffic. int8_ptx_3stage degrades sharply at N=4096 (+37%) because its triple-buffer schedule saturates the MIO queue with shared-memory pressure. int8_ptx_manual_pack stays within 2–5% of WMMA and posts the highest IPC of the group, showing that a dense ALU packing sequence keeps the scheduler fed even without ldmatrix. The scalar int8_dp4a path is never competitive—up to 4.8× slower—because DP4A emits several times more instructions per unit of arithmetic and cannot use the Tensor Cores at all. Two counter-level findings generalize. First, occupancy is not the performance predictor: int8_wmma has the highest occupancy at every size yet is consistently beaten by the register-capped k32. Second, at N=8192 the Average DRAM Active Cycles metric is an almost exact fingerprint of wall time—k32 logs 1.57 B DRAM-active cycles versus WMMA’s 2.80 B, a −43.9%-43.9\% reduction that matches its −43%-43\% wall-clock speedup, and the same correspondence holds for every other kernel (Fig. 2). This is the clearest evidence that N=8192 is DRAM-latency-bound and that the ranking is set by coalescing quality. Figure 2: INT8 (Run 2): Average DRAM Active Cycles. Flat (compute-bound) for N≤ 4096, then diverges at N=8192, mirroring the speedup ranking almost 1:1—the clearest evidence that large-N performance is DRAM-latency-bound. Figure 3: INT8 (Run 2): achieved occupancy. The highest-occupancy kernel (int8_wmma) is not the fastest, confirming that occupancy is not the bottleneck in this regime. TABLE VII: RUN 2 — INT8 RAW DURATIONS VS. N Kernel 512 1024 2048 4096 8192 int8_wmma (baseline) 152.00 µs 1.11 ms 8.55 ms 68.47 ms 858.30 ms int8_ptx_mma_k32 108.67 µs 726.46 µs 5.41 ms 42.00 ms 480.06 ms int8_ptx_mma_k16 204.54 µs 1.33 ms 9.61 ms 72.93 ms 582.00 ms int8_ptx_manual_pack 159.55 µs 1.17 ms 8.93 ms 69.54 ms 617.68 ms int8_ptx_3stage 151.87 µs 1.15 ms 10.22 ms 93.88 ms 820.32 ms int8_dp4a 588.80 µs 4.65 ms 36.75 ms 296.86 ms 2360 ms TABLE VIII: RUN 2 — INT8 SPEEDUP VS. INT8_WMMA Kernel 512 1024 2048 4096 8192 int8_ptx_mma_k32 1.40× 1.53× 1.58× 1.63× 1.79× int8_ptx_mma_k16 0.74× 0.83× 0.89× 0.94× 1.47× int8_ptx_manual_pack 0.95× 0.95× 0.96× 0.98× 1.39× int8_ptx_3stage 1.00× 0.97× 0.84× 0.73× 1.05× int8_dp4a 0.26× 0.24× 0.23× 0.23× 0.36× IV-C Run 3 — INT4: native MMA beats WMMA emulation The largest gains appear in INT4 (Tables IX–XI). All PTX kernels are 2.2×–4.3× faster than int4_wmma because the WMMA INT4 path (wmma::experimental::precision::s4) is software-expanded, inflating instruction count and introducing heavy lane-dependent divergence. The native PTX kernels call mma.sync.m16n8k64.s4 directly and avoid that expansion entirely. At N=8192, int4_wmma runs with only 16.89 active threads per warp and 4.6 million divergent branches, while the winning kernels sustain the full 32 active threads per warp with zero divergent branches and execute roughly 6–7× fewer instructions per scheduler (Fig. 4). Among the PTX kernels the lead varies depending on the GEMM size. int4_ptx_3stage is fastest at N=512–1024 thanks to deeper prefetch overlap, but as N grows it becomes strongly L2-driven: by N=8192 its L1/TEX hit rate collapses to ~14.5% while it drives L2 throughput to ~95.6%. int4_ptx_mma_k64 instead retains a ~61.6% L1 hit rate at N=8192 and therefore lower large-size latency, making it the best kernel from N=2048 onward (Fig. 5). The crossover is thus a memory-hierarchy effect, not a difference in tensor-core utilization—both kernels use the identical m16n8k64 MMA. Figure 4: INT4 (Run 3): average active threads per warp. The native-MMA kernels hold the full 32/32; int4_wmma degrades to 16.89 because the WMMA s4 path is software-emulated and lane-divergent. Figure 5: INT4 (Run 3): L1/TEX hit rate vs. N. int4_ptx_mma_k64 preserves L1 locality (~61.6% at N=8192) where int4_ptx_3stage collapses (~14.5%), explaining the large-N crossover. TABLE IX: RUN 3 — INT4 RAW DURATIONS VS. N Kernel 512 1024 2048 4096 8192 int4_wmma (baseline) 195.42 µs 1.450 ms 11.40 ms 90.11 ms 712.96 ms int4_ptx_mma_k32 87.01 µs 0.579 ms 4.32 ms 33.99 ms 269.35 ms int4_ptx_mma_k64 68.19 µs 0.395 ms 2.81 ms 21.52 ms 167.10 ms int4_ptx_manual_pack 87.84 µs 0.611 ms 4.61 ms 35.55 ms 277.13 ms int4_ptx_3stage 56.16 µs 0.382 ms 2.91 ms 24.26 ms 197.63 ms TABLE X: RUN 3 — INT4 SPEEDUP VS. INT4_WMMA Kernel 512 1024 2048 4096 8192 int4_ptx_mma_k32 2.25× 2.51× 2.64× 2.65× 2.65× int4_ptx_mma_k64 2.87× 3.68× 4.07× 4.18× 4.27× int4_ptx_manual_pack 2.23× 2.38× 2.48× 2.53× 2.57× int4_ptx_3stage 3.48× 3.80× 3.92× 3.72× 3.61× TABLE XI: RUN 3 — KEY COUNTERS AT N = 8192 Metric int4_wmma int4_ptx_mma_k64 int4_ptx_3stage Duration (ms) 712.96 167.10 197.63 Avg. active threads / warp 16.89 32 32 Avg. divergent branches 4,610,118 0 0 Avg. executed instr. / scheduler ~295 M ~45 M ~46 M L1/TEX hit rate (%) 80.34 61.56 14.49 L2 throughput (%) lower moderate 95.6 IV-D Run 4 — INT4 k64 family: a local optimum Run 4 fixes the winning m16n8k64 MMA strategy and sweeps three knobs: the A-loader split (x1/x2/x4 ldmatrix), the cp.async cache-eviction policy (.ca vs .cg), and the B operand layout (non-transposed vs transposed). Table XII and Fig. 6 show that the chosen configuration—x4 A-loader, x2 B-loader, non-transposed B, .ca policy—sits in a genuine local-optimum basin. The non-transposed .ca variants differ by only a few percent because they do not change the bottleneck class; they merely redistribute pressure between L1, L2, and instruction issue. The two off-basin variants fail in instructive ways. The .cg policy pushes traffic from L1 to L2 (L1 hit rate ~64% → ~20%, L2 throughput ~26% → ~50%), producing a mild but consistent latency regression. The transposed-B variant collapses: with only ~2.2 of every 32 bytes per global-load sector used, eligible warps per scheduler fall from ~0.57 to ~0.20 and warp cycles per instruction more than double, making it up to ~3.1× slower at large N. This confirms that coalescing quality, not MMA arithmetic, governs this kernel family. Figure 6: INT4 k64 ablation (Run 4): effective memory throughput (GB/s) for different kernels. The transposed-B variant loses most of its usable bandwidth to uncoalesced loads, confirming that coalescing—not MMA arithmetic—governs this kernel family. TABLE XII: RUN 4 — INT4 K64 FAMILY SPEEDUP VS. THE X4_X2NONTRANS_CA BASELINE Variant 512 1024 2048 4096 8192 x1_x2nontrans_ca 0.97× 0.97× 0.95× 0.94× 0.92× x2_x2nontrans_ca 0.96× 0.97× 0.96× 0.95× 0.95× x4_x1nontrans_ca 0.94× 0.95× 0.94× 0.94× 0.93× x4_x2nontrans_cg 0.94× 0.94× 0.92× 0.90× 0.88× x4_x2trans_ca 0.93× 0.84× 0.60× 0.39× 0.32× ⬇ // int4_ptx_3stage -- triple-buffered main loop // (one group kept in flight) for (int k = 2 * WMMA_K; k < K; k += WMMA_K) const int next = (buf + 1) % 3; const int prefetch_buf = (buf + 2) % 3; // Prefetch tile two stages ahead for (...) cp_async16(&As[prefetch_buf][warp_id][row][byte_col], src); for (...) cp_async16(&Bs[prefetch_buf][warp_id][n][byte_col], src); asm volatile("cp.async.commit_group;"); mma_int4_k64(rc0, ra, rb0); mma_int4_k64(rc1, ra, rb1); asm volatile("cp.async.wait_group 1;"); // keep 1 group outstanding buf = next; ldmatrix_a_k64(ra, As[buf][warp_id], lane_id); ldmatrix_b_k64(rb0, Bs[buf][warp_id], lane_id, 0); ldmatrix_b_k64(rb1, Bs[buf][warp_id], lane_id, 8); Listing 3: Three-stage pipeline using wait_group 1. Deeper overlap helps at small N but becomes L2-bound as N grows. V Discussion The results support a single organizing principle: PTX is worth its complexity exactly when it removes instruction overhead that the WMMA path cannot avoid. For FP16 the compiler’s WMMA lowering is already efficient and the workload becomes bandwidth-bound at scale, so hand-written PTX has nothing to recover and never wins. For INT8 the win is moderate and mechanistic—fewer executed instructions and better coalescing from the k32 decomposition. For INT4 the win is large and structural, because WMMA emulates the s4 operation in software while the PTX kernel issues a single native m16n8k64.s4 MMA. A second principle is that at inference scale the binding constraint is memory, not arithmetic. Quantization helps not only by halving or quartering the bytes per operand but by shrinking the working set enough to stay resident in L2, which is why INT4 escapes the DRAM cliff that caps FP16 at N=8192. The near-1:1 correspondence between Average DRAM Active Cycles and wall time at large N, and the transposed-B collapse in Run 4, both point to coalescing and cache residency as the true levers. For practitioners serving quantized open-weight LLMs (e.g., Llama-, Mistral-, Qwen-, or Nemotron-class models) on a single L4-class GPU, the practical takeaways are: prefer native INT4/INT8 MMA over WMMA emulation; tune the K-tile and loader split for coalescing before reaching for deeper pipelines; and treat occupancy as a diagnostic, not a target. Because only the SM target passed to the compiler changes, the same kernels are expected to transfer to A100 (SM80) and H100 (SM90), with gains that remain qualitatively similar though quantitatively hardware-dependent. VI Scope and Generality The present study focuses on a single GPU (NVIDIA L4, SM89), square matrices, and a fixed block/warp tiling; absolute numbers and some rankings may shift on other architectures, non-square shapes, or alternative tilings. Timings for the very largest FP16 case reflect a workload that overflows L2, so small measurement variance is amplified there. Quantized kernels compute INT32 accumulations without in-kernel dequantization or scaling, matching a typical inference pipeline but not end-to-end accuracy; numerical-quality evaluation of the quantized paths is out of scope and left to future work. Finally, the profiling relies on Nsight Compute counter definitions, whose semantics can differ subtly across driver and tool versions. VII Conclusion We presented a controlled, Nsight-Compute-instrumented comparison of hand-written PTX Tensor-Core GEMM kernels against WMMA baselines across FP16, INT8, and INT4 on an NVIDIA L4. PTX offers no benefit for FP16, a 1.4×–1.8× benefit for INT8, and a 2.9×–4.3× benefit for INT4 over the same-precision WMMA baseline; relative to FP16 WMMA the best quantized kernels reach 34.4× (INT8) and 98.7× (INT4) at N=8192. The gains are explained not by tensor-core utilization but by instruction count, coalescing quality, and cache residency, and occupancy is shown to be a poor predictor of throughput. The work provides both a reusable, precision-matched kernel family and a clear decision rule for when descending to PTX is justified in quantized LLM-inference GEMM. Artifact availability. All kernels, build scripts, Nsight Compute reports, and the profiling charts referenced in this paper are available in the accompanying repository, https://github.com/MattJBorowski1991/TensorCorePTX, and are permanently archived on Zenodo: https://doi.org/10.5281/zenodo.21815137. References [1] NVIDIA, “Parallel Thread Execution ISA (PTX),” NVIDIA CUDA Toolkit Documentation. [2] NVIDIA, “CUDA C++ Programming Guide,” NVIDIA CUDA Toolkit Documentation. [3] NVIDIA, “Nsight Compute Profiling Guide,” NVIDIA Developer Documentation. [4] NVIDIA, “NVIDIA Ada GPU Architecture,” Technical Whitepaper, 2022. [5] S. Markidis, S. W. D. Chien, E. Laure, I. B. Peng, and J. S. Vetter, “NVIDIA Tensor Core Programmability, Performance & Precision,” IEEE IPDPSW, 2018. [6] NVIDIA, “CUTLASS: CUDA Templates for Linear Algebra Subroutines,” open-source library. [7] T. Dettmers, M. Lewis, Y. Belkada, and L. Zettlemoyer, “LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale,” NeurIPS, 2022. [8] E. Frantar, S. Ashkboos, T. Hoefler, and D. Alistarh, “GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers,” ICLR, 2023. [9] J. Lin et al., “AWQ: Activation-aware Weight Quantization for LLM Compression and Acceleration,” MLSys, 2024. [10] Z. Jia, M. Maggioni, B. Staiger, and D. P. Scarpazza, “Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking,” arXiv:1804.06826, 2018. [11] V. Volkov, “Understanding Latency Hiding on GPUs,” Ph.D. dissertation, UC Berkeley, 2016. [12] NVIDIA, “cuBLAS Library User Guide,” NVIDIA CUDA Toolkit Documentation. [13] NVIDIA, “Inline PTX Assembly in CUDA,” NVIDIA CUDA Toolkit Documentation. [14] D. Yan, W. Wang, and X. Chu, “Demystifying Tensor Cores to Optimize Half-Precision Matrix Multiply,” in IEEE Int. Parallel and Distributed Processing Symp. (IPDPS), 2020. [15] W. Sun, A. Li, T. Geng, S. Stuijk, and H. Corporaal, “Dissecting Tensor Cores via Microbenchmarks: Latency, Throughput and Numerical Behaviors,” IEEE Trans. Parallel Distrib. Syst., vol. 34, no. 1, 2023.