Skip to content
← CUDA, Triton & GPU Performance / 09
hardNewNVIDIARed HatTogether AI

In a W4A16 GEMM, where does the INT4 unpacking happen, and how does the kernel keep that work hidden under the memory loads?

A 4-bit weight is only cheaper if turning it back into FP16 costs less than the bytes it saved. The answer counts the instruction budget per weight at full HBM bandwidth and shows the bit trick that makes the conversion fit.

Updated Oct 2026 · Worked answers on inference mechanisms, with calculations, implementation details, and follow-up questions.

A 4-bit weight is only cheaper if turning it back into FP16 costs less than the bytes it saved. The answer counts the instruction budget per weight at full HBM bandwidth and shows the bit trick that makes the conversion fit.

an account raises the free count per topic
UP NEXT ON YOUR JOURNEY
Next in this trackHere is an Nsight Systems timeline of one decode step. Find the launch gaps, synchronizations and CPU stalls, and price each one.Next in this trackHopper replaced Ampere's cp.async and mma.sync with TMA and WGMMA. What changed in each, and how does it change the way you design a kernel?Next in this trackBlackwell added tensor memory and 2-CTA MMA. What problem does each one solve, and what changes in a GEMM kernel written for it?
DISCUSSION · 0

Nobody has commented on this answer. Add a correction or another way to work it.