1
00:00:01,000 --> 00:01:03,524
[Hal Turing] Alrighty! Thanks for tuning in! Hello AI world! I am your host, Hal Turing, and my co-host is Dr. Ada Shannon.

And Ada, today we're going deep into the guts of the GPU. We're covering a paper called Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4. The lead author is Matt J. Borowski, with exactly one co-author, Blazej Osinski — so a tight two-person team. There's no academic affiliation listed here; both are described as independent kernel and machine learning engineers, which tells you something about who's actually doing this kind of ultra-low-level GPU work these days. It went up on arXiv on August 10th, 2026. And Ada, the thing that got me about this paper isn't one headline number — it's the shape of the result. For one precision, hand-written assembly-level GPU code buys you basically nothing over the compiler. For another precision, it changes everything. Same hardware, same authors, same kernel family.

2
00:01:03,524 --> 00:01:44,350
[Dr. Ada Shannon] Right, and that's really the question worth sitting with. Anybody who's ever hand-tuned a GPU kernel carries this strong prior that writing raw PTX instead of using NVIDIA's WMMA API is just better — faster, full stop, always. This paper picks that assumption apart by holding almost everything constant, same GPU, same tiling strategy, same benchmark harness, and only changing the numeric precision the kernel runs in. What falls out is that 'does PTX help' isn't a yes-or-no question. It's a question with a precision-dependent answer. That's a far more useful finding than 'PTX good' or 'PTX bad,' because it tells you exactly when the extra engineering effort is worth it and when it's a waste of your afternoon.

3
00:01:44,350 --> 00:02:00,925
[Hal Turing] Okay, before we get into that, let's actually back up, because I think a lot of listeners know Tensor Cores exist and make matmuls fast, but maybe not what a Tensor Core instruction physically is. Ada, what are we even talking about at the hardware level?

4
00:02:00,925 --> 00:02:49,225
[Dr. Ada Shannon] So a regular CUDA core is basically a tiny scalar ALU — one thread, one multiply-add per cycle, exactly like a CPU core. A Tensor Core is different in kind, not just degree. It's a warp-collective instruction, meaning all 32 threads in a warp jointly execute one instruction that multiplies small matrix tiles, think 16 by 16, in a single shot, with each thread holding just a fragment of the operands in its own registers. That's been true since the Volta generation back in 2017, and it's the actual silicon reason a transformer's attention and MLP projections run an order of magnitude faster than the same FLOPs would on plain CUDA cores. There are two doors into that hardware. WMMA, warp-matrix-multiply-accumulate, is the C++ API NVIDIA ships — you call load_matrix_sync, mma_sync, store_matrix_sync, and the compiler picks the tile shape and layout for you. Convenient, portable, the default since CUDA 9.

5
00:02:49,225 --> 00:02:56,050
[Hal Turing] And PTX is the other door, right? Walk me through it — when do you actually drop down that far?

6
00:02:56,050 --> 00:03:39,375
[Dr. Ada Shannon] That's the low-level path. Instead of trusting the compiler's fragment abstraction, you write PTX directly: cp.async to kick off an asynchronous copy from global memory into shared memory without stalling the warp, ldmatrix to cooperatively pull a properly-swizzled fragment out of shared memory into registers, and then mma.sync to actually issue the multiply-accumulate, at shapes like m16n8k16 or m16n8k64 that WMMA simply doesn't expose. You choose narrower, more numerous instructions instead of one wide WMMA call, control how deeply your memory pipeline is staged, and hand-place data in registers. Same hardware underneath, full stop — you're just picking the instruction stream yourself instead of letting—

7
00:03:39,375 --> 00:03:52,225
[Hal Turing] —Oh wait, wait, hold on, that's actually the part that gets me. So it's not that PTX unlocks new hardware capability, it's that WMMA is making a choice for you that you might not want?

8
00:03:52,225 --> 00:04:29,625
[Dr. Ada Shannon] Exactly. And whether overriding that choice pays off turns out to depend enormously on precision, which is the entire paper. There's good prior work behind this too. Markidis, Chien, Laure, Peng, and Vetter published one of the first systematic studies of WMMA performance and precision back in 2018 at IPDPS, when Volta was brand new. And separately, Jia, Maggioni, and Scarpazza reverse-engineered the actual SASS instructions WMMA compiles down to, that same year, which is honestly what makes hand-written PTX work possible at all, since NVIDIA never publishes that mapping itself.

9
00:04:29,625 --> 00:04:38,000
[Hal Turing] So let's talk about why anyone would bother with any of this. Where does it actually show up in something people run in production?

10
00:04:38,000 --> 00:05:28,500
[Dr. Ada Shannon] Quantized LLM serving. If you're running Llama, Mistral, Qwen, or Nemotron-class open-weight models at INT8 or INT4 instead of FP16, you've cut the bytes moved per operation by two or four times and shrunk your working set proportionally. That matters because past a certain problem size, a kernel stops being limited by how fast the Tensor Cores can multiply, that's compute-bound, and starts being limited by how fast you can pull data off DRAM, that's memory-bound. Quantization is partly an arithmetic trick and partly a memory trick — get the operands small enough and your tiles stay resident in cache instead of thrashing out to DRAM. That's the whole ballgame for serving models cheaply on something like an L4, NVIDIA's Ada-generation inference card, which is the only GPU this paper tests on.

11
00:05:28,500 --> 00:05:41,850
[Hal Turing] Okay, but I have to push back a little here, Ada — isn't testing on exactly one GPU kind of a narrow foundation to build a general claim on? How much of this even transfers off the L4?

12
00:05:41,850 --> 00:06:04,025
[Dr. Ada Shannon] I actually disagree with you there, Hal. A single-GPU study isn't a weakness here, it's the point. If you want to isolate whether a speedup comes from precision and instruction choice rather than GPU-to-GPU noise, you hold the hardware fixed and profile everything, full Nsight Compute metric set, on one well-understood chip. That's how you get to say 'this specific mechanism is what's driving the number' instead of eyeballing a table across three different architectures at once.

13
00:06:04,025 --> 00:06:15,350
[Hal Turing] Sure, but a result on one Ada-class card isn't the same as a result across Ampere and Hopper too — you could get a genuinely different picture on an A100.

14
00:06:15,350 --> 00:06:35,075
[Dr. Ada Shannon] That's fair, and it's a real limitation, but it's a different kind of limitation than 'the methodology is sloppy.' Controlled and narrow buys you causal confidence on this GPU. It doesn't automatically buy you generality across GPUs. Those are two separate questions, and I'd rather a paper nail the first one cleanly than hand-wave both at once.

15
00:06:35,075 --> 00:06:48,800
[Hal Turing] Okay, that's a fair distinction, controlled versus general. I'll take it. So with all that background in place, let's actually get into what they found when they ran the four precision families through this thing.

16
00:06:48,800 --> 00:06:59,150
[Hal Turing] So walk me through the actual experiment design, because I know this isn't just 'we ran one kernel and it was fast.' There are four separate runs here, right?

17
00:06:59,150 --> 00:07:41,075
[Dr. Ada Shannon] Four runs, each isolating a different precision family, plus a fourth that's an ablation on the winner. Run 1 is FP16 — six kernels, all PTX variants against the WMMA baseline. Run 2 is INT8, same structure, six kernels. Run 3 is INT4, the base comparison. And Run 4 takes the best INT4 kernel from Run 3 and sweeps three knobs on it — loader split, cache-eviction policy, operand layout — to check whether that winner is a real optimum or just got lucky. And the FP16 result is almost anticlimactic: every PTX variant stays within about five percent of the WMMA baseline, sometimes faster, sometimes slower, no consistent win either direction.

18
00:07:41,075 --> 00:07:51,625
[Hal Turing] Which is kind of a great result to lead with, honestly, because it immediately kills the 'PTX is always better' myth before they even get to the good stuff.

19
00:07:51,625 --> 00:08:30,525
[Dr. Ada Shannon] Right, and then INT8 is where it starts to move. The winning kernel, int8_ptx_mma_k32, is consistently faster than int8_wmma — 1.4x at the small end up to 1.8x at N equals 8192. And it's not vibes, they show the mechanism directly: it decomposes the K-tile differently and ends up executing 25 to 43 percent fewer instructions than WMMA, and its global-load coalescing is nearly perfect — 0.4 percent wasted sectors at the largest size versus roughly 50 percent for every other kernel in the family, including the WMMA baseline itself.

20
00:08:30,525 --> 00:08:41,325
[Hal Turing] Wait, hold on — fifty percent wasted sectors for basically everything else? That's not a small tuning gap, that's like half your memory bus doing nothing useful.

21
00:08:41,325 --> 00:09:18,925
[Dr. Ada Shannon] Exactly, and that's the tell that this isn't a Tensor Core story, it's a memory-access-pattern story. Which is why the occupancy angle matters — int8_wmma has the highest occupancy of the whole group at every single size, and it still loses to k32. Same thing shows up back in FP16: the accumulator variant that switches from FP32 to FP16 accumulators frees up registers and pushes occupancy up to 70 to 82 percent, and GFLOPS don't move at all. Two separate precisions, same conclusion — once you're not warp-count-limited, more occupancy buys you nothing.

22
00:09:18,925 --> 00:09:39,351
[Hal Turing] And then INT4 is where the numbers get almost silly. 2.9x to 4.3x over int4_wmma, and 98.7x relative to the FP16 WMMA baseline at N equals 8192. That's not a tuning win, that's a different category.

23
00:09:39,351 --> 00:10:18,276
[Dr. Ada Shannon] Because it's not the same kind of bottleneck. WMMA's INT4 path is software-emulated — there's no native WMMA fragment for s4 at that shape, so the compiler expands it into a divergent, lane-dependent instruction sequence. The PTX kernels just issue mma.sync.m16n8k64.s4 directly, one native instruction, full 32 active threads per warp, zero divergent branches. INT8 relative to FP16 WMMA lands at 34.4x for comparison — still huge, but INT4 nearly triples that because it's removing an entire emulation layer, not just trimming instruction count.

24
00:10:18,276 --> 00:10:36,051
[Hal Turing] Okay but here's where I want to push on something — inside the INT4 results, the best kernel actually changes depending on N. The three-stage pipeline wins small, and the k64 kernel takes over at scale. Isn't that a little bit of a shrug — 'it depends'?

25
00:10:36,051 --> 00:11:05,226
[Dr. Ada Shannon] No, I'd push back on that, Hal — that's not a shrug, that's the finding. Three-stage wins small because deeper prefetch overlap hides latency when you're compute-bound. But as N grows its L1 hit rate collapses to about 14.5 percent while k64 holds around 61.6 percent, and that gap is entirely responsible for the crossover around N equals 2048. Same MMA instruction in both kernels, identical arithmetic — the only thing that changes is cache residency.

26
00:11:05,226 --> 00:11:15,776
[Hal Turing] I hear you, but doesn't 'the answer depends on N' undercut the clean story they're telling? If I'm an engineer picking one kernel to ship, I don't get a single winner.

27
00:11:15,776 --> 00:11:56,151
[Dr. Ada Shannon] You get a single winner per regime, which is more useful than a fake universal answer. And Run 4 backs that up rather than muddying it — they fix the k64 MMA shape and just vary loader width, cache policy, and B-layout. Every non-transposed variant lands within a few percent of each other. But flip B to transposed and it collapses to 0.32x at N equals 8192 — only about 2.2 of every 32 bytes per load sector actually get used. Same MMA instruction, same precision, just uncoalesced loads, and that alone is worth over 3x. That's about as clean a confirmation as you'll get that coalescing, not the arithmetic, is what's governing this whole kernel family.

28
00:11:56,151 --> 00:12:35,751
[Hal Turing] Fair enough. But that discipline is exactly what makes me want to push on the baseline, Ada — every comparison in this paper is PTX versus the WMMA C++ API. Not CUTLASS, not cuBLASLt. And their own background section cites CUTLASS as reference six, calling it the production route to exactly this level of tuning. So when they report one-point-four to four-point-three x over WMMA, how much of that survives against a properly tuned CUTLASS kernel at the same precision, instead of against a comparatively naive double-buffered baseline nobody actually ships?

29
00:12:35,751 --> 00:13:12,326
[Dr. Ada Shannon] Honestly, that's the biggest hole in the paper. CUTLASS is built from the same mma-sync, ldmatrix, cp-async primitives they hand-wrote — it's just years of tuned templates wrapped around them. Cite it and never run it as a comparison, and we can't tell if this is a fundamental PTX advantage or just the gap between naive and competent. Same blind spot with Triton, which isn't mentioned at all — that's what most LLM-serving engineers actually reach for now instead of hand-rolling anything. PTX versus WMMA may not even be the dichotomy practitioners are choosing between anymore.

30
00:13:12,326 --> 00:13:48,126
[Hal Turing] Same worry applies to the shapes they tested. Every benchmark is a square dense GEMM, N from five-twelve to eight-thousand-one-ninety-two. But single-batch LLM decode looks nothing like that — it's skinny, M around one to thirty-two, with huge K and N. Completely different memory access pattern. So when they frame all of this around serving Llama, Mistral, Qwen, Nemotron-class models, I genuinely don't know if the k32 coalescing win or the k64 residency trick survives in that regime at all.

31
00:13:48,126 --> 00:14:09,951
[Dr. Ada Shannon] Can't defend that one for them either. Those coalescing and instruction-count advantages come directly from how a fixed sixteen-by-sixteen warp tile decomposes across a square matrix — shrink one dimension to single digits and the whole load pattern changes. The ranking could flip entirely. Square-GEMM results are honest; the 'this is how you serve quantized LLMs' framing reaches past what was actually measured.

32
00:14:09,951 --> 00:14:32,177
[Hal Turing] I'd almost let that slide, since square GEMM still matters for prefill. What bugs me more is they never measure accuracy — no dequant, no scale, straight INT32 accumulate. But it's a kernel paper, not a quantization paper. Isn't 'that's out of scope' a fair thing to say and just let someone else measure whether INT4 tanks quality?

33
00:14:32,177 --> 00:15:05,702
[Dr. Ada Shannon] No — wait, hold on, I actually disagree with you there, Hal. You can't build the entire motivation around serving quantized LLMs and then wave away the one number that decides whether that's usable. GPTQ and AWQ exist because naive INT4 without calibration can wreck real model outputs. A kernel that's four times faster but needs AWQ-style scaling bolted on isn't reporting speed in isolation — it's reporting half a two-variable tradeoff and only showing the flattering half.

34
00:15:05,702 --> 00:15:26,302
[Hal Turing] Isn't that disclosed, though? Section six says numerical-quality evaluation is explicitly out of scope and left to future work — that's honest scoping, not something they're hiding from the reader. Nobody in the paper is claiming raw INT32 accumulation is what you'd actually ship to production untouched.

35
00:15:26,302 --> 00:15:46,677
[Dr. Ada Shannon] Disclosed, sure, I'll give them that — it's not deceptive. But disclosed doesn't mean it stops mattering for the audience they're targeting. If you're the engineer deciding whether to adopt native INT4 mma-sync, you need to know how much of that four-x gets eaten by calibrated scaling. Scoping it out is fine for the paper's honesty; it's still a real gap for its stated audience.

36
00:15:46,677 --> 00:16:10,127
[Hal Turing] Fair, we're agreeing on the facts, just weighing them differently. Zooming out — this whole 'PTX wins exactly where WMMA can't avoid overhead' principle isn't new either, is it? That's basically Markidis, Chien, Laure, Peng and Vetter's Tensor Core programmability paper, IEEE IPDPSW 2018 — reference five here — just applied across more precisions.

37
00:16:10,127 --> 00:16:49,652
[Dr. Ada Shannon] Right, that's the earliest systematic look at this exact tradeoff, just at FP16 since INT8 and INT4 Tensor Cores didn't exist yet on that hardware. The occupancy finding is the same story — Vasily Volkov's 'Understanding Latency Hiding on GPUs,' his UC Berkeley dissertation from 2016, already showed occupancy doesn't predict throughput once you're not warp-count-limited. This paper reconfirms both nicely in the Tensor Core era, but neither is a new discovery. And the A100, H100 transfer claim is asserted, never tested — plus no run-count or variance reported anywhere, which matters when some Run 4 deltas are a few percent apart.

38
00:16:49,652 --> 00:17:19,627
[Hal Turing] So practically: real, useful decision rule — go native for INT4, don't bother for FP16, tune coalescing before pipeline depth — but only inside the box they actually tested. One GPU, square matrices, no accuracy check, no CUTLASS in the comparison. Use it as a heuristic, validate on your own shape and hardware, and don't treat the multipliers as universal. Good place to land, honestly — great mechanism, narrow scope. Thanks for digging into this one with me, Ada.

39
00:17:19,627 --> 00:17:24,102
[Dr. Ada Shannon] Always fun, Hal. Thanks for listening, everyone — see you next time.