100 DAYS OF GPU Course preview
Day 98Module 10
in-technical-review

Quantization and sparsity

Take a weight matrix whose columns do not share a dynamic range, which is every weight matrix in a real network, and quantize it to INT8 the way the first tutorial you find tells you to: one scale for the whole tensor, max|B| / 127. Now look at the column whose largest value is 64 times smaller than the matrix maximum. Its step size was set by somebody else's numbers, so all of its values land on two codes, plus zero. Three distinct values where there were thousands. The layer still runs, the test still passes if it uses one tolerance for the whole output, and one output channel of your model is now noise.

By the end of this page you can write an INT8 matmul with the scale arithmetic in the right place, bound its error against a double-precision reference, and say what the 8 bits bought in time, which is far less than four times.

What a scale actually promises

A quantized tensor is two objects: a buffer of small integers and the arithmetic that turns them back into numbers. Symmetric quantization uses one scale and nothing else, x = q * s, with s = max|x| / 127 and the code 0 meaning the value 0. Asymmetric quantization adds a zero point, x = (q - z) * s, where z is the code that dequantizes to zero and s covers the observed range rather than twice the largest magnitude. On a tensor that never goes negative, which is what a ReLU hands you, symmetric spends 127 of its 255 codes on values that cannot occur, so asymmetric is worth exactly one bit.

The second decision is what a scale is a promise about. Per-tensor means one number for the whole matrix and a step size set by the single largest element anywhere in it. Per-channel means one scale per column of B, one per output channel of the layer, so a quiet column keeps all 255 of its codes, at the cost of an array of N floats and one extra load per output element.

Now the part that surprises people: the multiplication by those scales happens once, at the very end. Inside the dot product there are no scales and no floats. Products of two 8-bit integers accumulated in INT32 round nothing at all, so a 2048-term INT8 dot product is exact integer arithmetic. That is the opposite of day 71, where the accumulator decided the error. Here |acc| is at most K * 127 * 128 and cannot lose a bit, so every bit of error in the answer was spent by the quantizer before the kernel launched.

Diagram: where a scale is a promise. One column of a weight matrix, drawn three times as a row of value ticks against the INT8 code grid it lands on, with the code count under each row. One scale for the whole matrix: the column's values are 64x smaller than the matrix maximum, so they land on two codes plus zero. Caption "3 of 255 codes used". One scale per column: the same values against a step of their own column's maximum over 127. Caption "255 of 255 codes used". One scale per 32 values: the same column with one outlier 8x its neighbours, so the outlier's block pays and the other blocks do not. Caption "255 codes per block of 32". Alt text: "The same weight column under three scaling policies. One scale for the matrix leaves it three codes out of 255. One scale per column gives it all 255. One scale per block of 32 stops an outlier spending its neighbours' range."

Why four times the bit width is not four times the speed

Eight bits instead of thirty-two reads as a factor of four, and for the weights on disk it is one. For the kernel it is not, and the reason is already in this course. Day 44 spent its whole ladder on one ratio, shared loads per fused multiply-add, because that is what a tiled matmul is limited by once the tiles are in shared memory. Narrowing the element from four bytes to one changes how many bytes each of those loads moves. It does not change how many loads there are, how many multiply-adds there are, or how many instructions the inner loop issues.

Three things stay exactly where they were. The output C is still FP32, so a quarter of the traffic never shrank. The epilogue still does its floating-point work per output element, and with a zero point it also reads a column sum. And the tile loop still issues one multiply-add per element pair, now an integer one, which on a Turing card is not cheaper than the float it replaced.

So the honest claim for this kernel is not a speedup at all. It is that the error is bounded, checkable and much smaller than the memory saving suggests it should be, and that whatever time it buys comes from DRAM traffic on the inputs alone. The 4x lives in the tensor-core path this program deliberately does not take.

The program, and the gate that has to be per column

Full program in quant_matmul.cu, under code/day98-quantization/. It runs day 16's tiled kernel in FP32 as the baseline, then the same geometry in INT8 three times, at 256, 1024 and 2048.

The three INT8 rows change one decision each. One scale for all of B, then a scale per column of B, then a zero point on A as well. The kernel is one template with those two decisions as template parameters, so a row of the table differs from the row above it in exactly one thing and never in a tile shape, a launch configuration or an access pattern.

The reference is double precision over the original FP32 inputs, never over the dequantized copies. A reference rebuilt from the quantized values would forgive the quantization error, and on this page the quantization error is the only error there is.

The inputs are shaped like the thing being measured. A is non-negative and uniform in [0, 1), which is what a ReLU produces and what makes the zero point worth a bit. B is signed, with a per-column gain spanning 64x and an outlier every 37th element at 8x its neighbours. Both come from a small linear congruential generator rather than a repeating pattern: the quantization errors of a periodic sequence are periodic too, and correlated errors would grow with K instead of its square root, which would make the tolerance below a lie.

The accumulate and the epilogue are the whole kernel worth reading:

    int acc = 0;
    for (size_t tileIdx = 0; tileIdx < tiles; ++tileIdx) {
        const size_t aCol = tileIdx * kTileDim + threadIdx.x;
        const size_t bRow = tileIdx * kTileDim + threadIdx.y;

        tileA[threadIdx.y][threadIdx.x] =
            (row < m && aCol < k) ? a[row * k + aCol] : aPad;
        tileB[threadIdx.y][threadIdx.x] =
            (bRow < k && col < n) ? b[bRow * n + col] : int8_t{0};
        __syncthreads();

        for (int p = 0; p < kTileDim; ++p) {
            acc += static_cast<int>(tileA[threadIdx.y][p]) *
                   static_cast<int>(tileB[p][threadIdx.x]);
        }
        __syncthreads();
    }

    if (row < m && col < n) {
        int corrected = acc;
        if constexpr (kAsymA) {
            corrected -= aZero * bColSums[col];
        }
        const float scale = aScale * (kPerChannelB ? bScales[col] : bScale);
        c[row * n + col] = static_cast<float>(corrected) * scale;
    }

The correction term is what the zero point costs. Expanding sum_k (a_q - z) * b_q gives the integer dot product minus z * sum_k b_q[k][col], and that column sum depends only on the weights, so it is computed once when they are quantized and read once per output element. A zero point is O(N^2) work bolted onto an O(N^3) loop.

The tolerance is derived rather than tuned, and it is per column because the columns are not comparable:

// The absolute tolerance for one output column, derived rather than tuned.
// Write a = a_q_dequantized + ea with |ea| <= sa/2 and likewise for b. One
// term of the dot product then carries at most |a|*sb/2 + |b|*sa/2, and with
// sa = amax/127 and sb = bColMax/127 that is amax*bColMax/127. A
// correctness gate cannot assume those K errors are independent, so its
// deterministic bound adds their magnitudes and scales with K. The sqrt(K)
// estimate remains useful as a statistical diagnostic, but it does not decide
// whether a kernel is correct. kGateSlack is the same factor 4 day 66 uses
// for a valid-but-different summation order and the omitted second-order term.
//
// The tolerance is per column because the columns of B differ by 64x here.
// One tolerance built from the matrix maximum would be 64x too loose on the
// smallest column, which is exactly the mistake the per-tensor row makes.
static double columnGateAtol(double amax, double bColMax, size_t kDim) {
    return kGateSlack * static_cast<double>(kDim) * amax * bColMax /
           static_cast<double>(kInt8Max);
}

static double columnStatAtol(double amax, double bColMax, size_t kDim) {
    return kGateSlack * std::sqrt(static_cast<double>(kDim)) * amax * bColMax /
           static_cast<double>(kInt8Max);
}

The per-tensor row is measured against that gate and never gated on it. It is expected to fail, its failure is the lesson, and a program that stopped there would not print the rows that fix it. Day 62 treats its racy kernels the same way. Timing is CUDA events, three warm-ups and the mean of ten runs, copies excluded.

The stretch tiers, and what a T4 can honestly say

Each of the stretch formats splits cleanly into a half that needs hardware the floor card does not have and a half that is arithmetic.

FP8 e4m3 and e5m2 need compute capability 8.9 for the mma instruction (".e4m3 and .e5m2 alternate floating point type mma operation requires sm_89 or higher", https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma , checked 2026-08-29). The format's rounding needs no hardware, so the program quantizes the same weight matrix to e4m3 on the host and reports what it costs. Three mantissa bits is a relative step of one part in sixteen everywhere, where INT8's step is one part in 254 of the matrix maximum and nothing else. Integers are uniform in absolute terms and floats in relative terms, which is the entire trade.

MX block scaling is what makes 8-bit and 4-bit weights usable: one shared scale per block of 32 values, so an outlier spends its own block's range and not the tensor's. The hardware form is Blackwell's; the idea is arithmetic, and the program measures it by scaling e4m3 per 32 values down a column.

2:4 structured sparsity is two of every four weights forced to zero. Forcing the pattern is free and costs accuracy, which the program measures by pruning B and rerunning the per-column kernel against the dense reference. Getting the 2x back needs the sparse tensor cores at 8.0 and cuSPARSELt (cuSPARSE and friends; https://docs.nvidia.com/cuda/cusparselt/index.html , checked 2026-08-29), which this program does not link, so nothing here is timed against a dense kernel reading zeros.

The card also has one capability this program declines to use. A T4 has INT8 tensor cores at 7.5, from Table 33 of the compute capability appendix (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-29). The inner loop above is integer multiply-adds on CUDA cores, so every time in the table is a CUDA-core time, and the transcript prints that sentence itself.

Results

Re-verified on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Every error, census and deterministic correctness value reproduced exactly; all three cases passed and the process exited 0. Timing drift was material at 1024, while the three 2048 INT8 rows remained close. Both transcripts are listed in front matter; this table is CUDA 13.0.

case kernel max err rms err K gate frac sqrtK frac time (ms) vs f32
256 matmulTiledF32 1.907e-05 1.139e-06 0.000 0.000 0.0832 1.00
256 int8 one scale 6.449e-01 1.611e-01 0.459 7.337 0.1012 0.82
256 int8 per-column 5.360e-01 6.792e-02 0.011 0.181 0.1014 0.82
256 int8 per-col + zp 5.334e-01 6.683e-02 0.010 0.161 0.1018 0.82
256 int8 2:4 pruned 8.765e+00 1.045e+00 n/a n/a not timed n/a
1024 matmulTiledF32 8.392e-05 4.364e-06 0.000 0.000 5.7960 1.00
1024 int8 one scale 2.287e+00 3.211e-01 0.186 5.944 7.1859 0.81
1024 int8 per-column 1.229e+00 1.320e-01 0.015 0.483 4.4374 1.31
1024 int8 per-col + zp 1.221e+00 1.294e-01 0.015 0.471 3.3166 1.75
1024 int8 2:4 pruned 1.961e+01 2.142e+00 n/a n/a not timed n/a
2048 matmulTiledF32 2.924e-04 8.771e-06 0.000 0.000 28.0189 1.00
2048 int8 one scale 7.333e+00 4.899e-01 0.190 8.588 23.5013 1.19
2048 int8 per-column 8.311e+00 2.468e-01 0.016 0.742 23.7347 1.18
2048 int8 per-col + zp 8.160e+00 2.441e-01 0.016 0.729 23.7970 1.18
2048 int8 2:4 pruned 4.120e+01 2.865e+00 n/a n/a not timed n/a

The 2:4 rows were accuracy-only. The five predictions resolved as follows.

  1. Refuted. The 2048 FP32 row improved from 32.2266 to 28.0189 ms but still did not reproduce day 44's 25.350 ms.
  2. Held. One-scale sqrt(K) fractions were 7.337, 5.944 and 8.588; per-column fractions stayed below 0.75. Improvements were 40.45x, 12.30x and 11.57x, while every deterministic K-scaled gate passed.
  3. The old speed band did not reproduce. The three 2048 INT8 ratios moved from roughly 1.35x to 1.18x or 1.19x. They remain tightly grouped, so the shared integer-loop cost still dominates the scale policy.
  4. Held. Symmetric activation RMS was 2.008x asymmetric. End-to-end RMS improved only 1.016x, 1.020x and 1.011x with the zero point.
  5. Held. The one-scale INT8 maximum storage error was 3.937e-03 of max|B|. FP8 one-scale max error was 9.07x larger while its RMS was lower; per-32 scaling cut FP8 max error 4.89x. At 2048, 2:4 RMS error was 11.6x the dense per-column INT8 row.

The support report also kept the mechanism honest: this T4 has INT8 tensor cores, but the program used integer MADs on CUDA cores; its FP8 and 2:4 rows were host arithmetic and never timed.

Run it yourself

Any card at sm_75 or above builds and runs this, a free Colab session included, and it needs no root, no profiler and no library beyond the toolkit:

nvcc -std=c++17 -O3 -arch=sm_75 -o quant_matmul quant_matmul.cu

No Compiler Explorer embed for this one: the double-precision reference at 2048 is 8.6 billion multiply-adds on one core, which does not fit a 20 second execution cap. On an sm_89 card the capability report's FP8 line flips to yes and nothing else changes, because the kernels use no tensor cores.

Exercise

The page fixed B's scales and left A on one scale for the whole tensor. Quantize A per row instead, one scale per token, and measure it. Then make one row of A carry a value 20 times the rest and measure again.

Time: 30 to 45 minutes. Submit: the changed quantizer and epilogue, the max, RMS and worst-column numbers for your per-row variant next to the shipped per-tensor row on both inputs, and one sentence on why the two measurements disagree.

Check: your variant must pass the same deterministic per-column gate the shipped per-column rows pass, |got-ref| <= 4*K*max|A|*max|B[:,c]|/127, with the first offending row and column printed on failure. Also report the sqrt(K) fraction as a statistical diagnostic. The interesting number is the ratio between your variant and the shipped one, on each of the two inputs.

Hint 1

A scale is a promise about a set of values. Ask what set each scale in this program covers, and what has to be true of that set for the promise to be cheap. Then ask which of the two inputs makes it true.

Hint 2

The epilogue already multiplies by a per-column array for B. A per-row scale for A is the same array indexed by the other coordinate. Before you time it, predict the ratio from the row maxima alone: what is max(row)/max(tensor) on the shipped input, and on the one with the outlier?

Solution

aScales[row] beside bScales[col], and aScale * bScales[col] in the epilogue becomes aScales[row] * bScales[col]. On the shipped input every row of A is drawn from the same distribution, so every row maximum is close to the tensor maximum and per-row scaling changes the error by a percent or two while adding an array and a load. On the input with one 20x row, the tensor scale is set by that row, every other row loses more than four bits, and per-row scaling gets them back.

Granularity is not a quality setting you turn up. It pays exactly where the dynamic range varies along the axis you split, and costs a load everywhere else, which is why production stacks quantize activations per token and weights per channel rather than everything everywhere.

Pitfalls

One tolerance for the whole output hides a dead channel. A test that compares against atol = 1e-2 * max|C| is dominated by the largest column, so a column that came back as three distinct values still passes. Scale the tolerance per column, from that column's own inputs, the way columnGateAtol does. The same trap in a different format is day 71's zero-error FP16 test.

Checking the quantized kernel against a dequantized reference. If the reference is built from the same rounded inputs the kernel got, the quantization error cancels on both sides and the test measures the indexing only. The reference has to be the original FP32 values in double, which is also day 66's rule.

Rounding to 128 and storing it in an int8_t. max|x| / 127 puts the largest element exactly on 127, but the next value up, or a scale computed from a stale calibration, rounds to 128, which wraps to -128. The symptom is one enormous negative weight in an otherwise fine matrix. Clamp after rounding, always, and keep the clamp even when you can prove it is never hit.

Dropping the zero-point correction term. The output is not noisy, it is offset: every element of a column is wrong by the same z * sum_k b_q multiple. It looks like a bias bug, it survives eyeballing a few values, and it goes away the moment you switch to symmetric quantization, which is how people conclude the zero point is broken.

Reporting an INT8 speedup as a tensor-core result. These kernels never touch a tensor core; any win is DRAM traffic on the inputs. The instruction that would use the INT8 tensor cores a 7.5 card already has is day 72's subject with a different input type. Day 47 is the standing lesson on attributing a number to the wrong mechanism.

Go deeper

Next

Day 99 returns to FP32 and assembles one forward and backward pass. Its new problem is not representing a value with fewer bits but preserving the shape and orientation of every gradient: a transposed GEMM, a column reduction and a ReLU mask all have to agree before day 100 can train. The mixed precision question then belongs to the whole training step rather than one kernel.