Tensor cores and matrix operations: Exercises — Hardware deep dive

Figure — Tensor cores and matrix operations

The picture above is our anchor: to fill one output cell C[i][j], you slide along row i of A and column j of B, multiply the pairs you meet, and add them up. Keep it in mind — every exercise is a variation on this slide-and-sum.

Level 1 — Recognition

L1.1 — Naming the operation

Problem. A Tensor Core executes D = A × B + C on small square tiles. In plain words, what are the two arithmetic actions packed into that single instruction, and what is the accumulator?

Recall Solution

The two actions are a matrix multiply (A × B) followed by a matrix add (+ C). This combined action is a matrix multiply-accumulate (MMA). The accumulator is C (and the result overwrites into D): it is the running total we add fresh partial products into. Keeping a separate accumulator is what lets us sum many tiles without losing anything.

L1.2 — Reading tile shapes

Problem. An Ampere Tensor Core advertises a 16×8×16 tile in M×N×K notation. Which of M, N, K is the shared inner dimension that gets summed over, and what are the dimensions of the output tile it produces?

Recall Solution

In C = A × B with A being M×K and B being K×N, the letter K is the shared dimension that appears in both A and B — it is the one we sum over (the length of the row/column slide). So here K = 16 is the inner dimension. The output tile C is M×N = 16×8.

Level 2 — Application

L2.1 — Counting FMAs and FLOPs

Problem. You multiply two 512 × 512 matrices. (a) How many FMAs does the naïve element-by-element algorithm perform? (b) How many FLOPs is that?

Recall Solution

Each output cell C[i][j] needs K FMAs (one per term in the sum). Here the matrices are M×K times K×N with M = N = K = 512.

Number of output cells: M·N = 512·512.
FMAs per cell: K = 512.
(a) Total FMAs = M·N·K = 512³ = 134{,}217{,}728 ≈ 1.34 × 10⁸ FMAs.
(b) Each FMA = 2 FLOPs, so 2 × 512³ = 268{,}435{,}456 ≈ 2.68 × 10⁸ FLOPs.

L2.2 — Counting Tensor Core instructions (4×4 tiles)

Problem. Multiply 256 × 256 matrices using Volta-style 4×4 hardware tiles. How many mma.sync tile-instructions are issued, and by what factor is that fewer than scalar FMAs?

Recall Solution

We chop each dimension into tiles of size TILE = 4. Along each axis there are 256 / 4 = 64 tiles.

The output grid is 64 × 64 tiles.
For each output tile we slide over the inner dimension in 64 tile-steps.
Tile-instructions = 64 × 64 × 64 = 64³ = 262{,}144.

Scalar FMAs would be 256³ = 16{,}777{,}216. Ratio: $\frac{25 6 ^{3}}{6 4 ^{3}} = 4^{3} = 64.$ Each tile op internally does 4×4×4 = 64 FMAs, so the 64× fewer instructions is exactly the work packed inside one Tensor Core op. 262,144 instructions, a 64× reduction.

L2.3 — Peak throughput timing

Problem. A GPU sustains 100 TFLOP/s on Tensor Cores. How long (in milliseconds) does the pure compute for a 1024 × 1024 × 1024 matrix multiply take, ignoring memory?

Recall Solution

FLOPs needed = 2 × 1024³ = 2 × 1{,}073{,}741{,}824 = 2{,}147{,}483{,}648 ≈ 2.147 × 10⁹.
Rate = 100 TFLOP/s = 100 × 10¹² = 10¹⁴ FLOP/s.
Time = FLOPs ÷ rate = 2.147 × 10⁹ / 10¹⁴ = 2.147 × 10⁻⁵ s. $t \approx 2.147 \times 1 0^{- 5} s = 0.0215 ms .$ So about 0.021 ms of pure arithmetic. Real time is longer because memory must feed this — see 9.1.5-Roofline-model.

Level 3 — Analysis

L3.1 — A tile multiply by hand

Problem. With TILE = 2 (a mini Tensor Core), compute D = A × B + C where $A = [1324], B = [5768], C = [1001] .$

Recall Solution

Use the slide-and-sum rule D[i][j] = Σₖ A[i][k]·B[k][j] + C[i][j], with k running 0,1.

D[0][0] = (1·5 + 2·7) + 1 = (5+14)+1 = 20
D[0][1] = (1·6 + 2·8) + 0 = (6+16)+0 = 22
D[1][0] = (3·5 + 4·7) + 0 = (15+28)+0 = 43
D[1][1] = (3·6 + 4·8) + 1 = (18+32)+1 = 51 $D = [20432251] .$ Notice the accumulator C only nudges the diagonal here — that is exactly what an identity matrix does.

L3.2 — Why accumulation over K stays correct

Problem. You split a K = 8 multiply into two 4×4 tile-steps (k = 0..3, then k = 4..7), adding the second result on top of the first. Explain the algebraic property that makes this equal to doing all 8 at once, and confirm numerically for one cell: the row A[0] = [1,1,1,1,1,1,1,1] times column B[:,0] = [2,2,2,2,2,2,2,2].

Recall Solution

The rule is that a sum can be split into partial sums and re-added — that is the associativity of addition: $\sum_{k = 0}^{7} x_{k} = (\sum_{k = 0}^{3} x_{k}) + (\sum_{k = 4}^{7} x_{k}) .$ The accumulator C holds the first partial sum, then the second tile op adds its own partial sum on top. Nothing is lost because addition doesn't care how we group the terms. Numerically each term is 1·2 = 2.

First tile (k=0..3): 2+2+2+2 = 8.
Second tile (k=4..7): 2+2+2+2 = 8.
Accumulated total: 8 + 8 = 16, identical to 2·8 = 16 done in one go. ✓

Level 4 — Synthesis

L4.1 — Mixed-precision error budget

Problem. TF32 keeps a 10-bit mantissa; FP32 keeps 23 bits. The relative rounding error of a format with an m-bit mantissa is about 2⁻ᵐ. (a) Give the TF32 and FP32 per-number relative errors. (b) By what factor is TF32 coarser?

Recall Solution

The mantissa is the "significant digits" part of a floating-point number; more mantissa bits = finer resolution. The relative error is roughly 2⁻ᵐ.

(a) TF32: 2⁻¹⁰ = 1/1024 ≈ 9.77 × 10⁻⁴ (about 0.098%). FP32: 2⁻²³ ≈ 1.19 × 10⁻⁷.
(b) Ratio = 2⁻¹⁰ / 2⁻²³ = 2¹³ = 8192. So each stored number in TF32 is about 8192× coarser, yet because Tensor Cores accumulate in FP32, the summed result stays accurate enough for training. See 8.3.4-Mixed-precision-training.

L4.2 — Choosing a precision, with justification

Problem. You must run inference on a pre-trained, frozen image classifier and want maximum throughput. Turing Tensor Cores support FP16 and INT8. If INT8 gives 4× the throughput of FP16 for the same matmul, and the accuracy drop from INT8 quantization is under 1%, which do you pick and why?

Recall Solution

Pick INT8. Reasoning chain:

The weights are frozen — no gradients to protect, so the wide dynamic range of floating point is unnecessary.
INT8 is 4× the throughput of FP16 here, and 4× FP16's already-large lead over FP32.
The accuracy cost (<1%) is acceptable for classification inference. The accumulator is INT32 to avoid overflow when summing many int8 × int8 products. This is exactly the 8.4.2-Quantization-techniques tradeoff: give up precision you no longer need to buy speed. If this were training, you'd instead reach for BF16 (see the L4.1 note).

Level 5 — Mastery

L5.1 — Padding to a full tile

Problem. A layer produces a 1000 × 1000 output but Tensor Cores want dimensions that are multiples of 16. (a) What is the smallest padded size? (b) How many extra "wasted" output cells does padding create, and what fraction of the padded work is wasted?

Recall Solution

(a) Round each dimension up to the next multiple of 16. 1000 / 16 = 62.5, so we need 63 tiles per axis → 63 × 16 = 1008. Padded size 1008 × 1008.
(b) Padded cells = 1008² = 1{,}016{,}064. Real cells = 1000² = 1{,}000{,}000. Wasted cells = 1{,}016{,}064 − 1{,}000{,}000 = 16{,}064.
Fraction wasted = 16{,}064 / 1{,}016{,}064 ≈ 0.01581 = 1.58%. A tiny padding tax buys full Tensor Core utilization — worth it because a partial tile falls off the fast path entirely (the "performance cliff").

L5.2 — End-to-end throughput comparison

Problem. Two 2048 × 2048 × 2048 matmuls run on the same GPU. Path A uses CUDA cores at 12 TFLOP/s. Path B uses Tensor Cores at 120 TFLOP/s. (a) Compute-only time for each. (b) Speedup. (c) If Path B is actually memory-bound and only reaches 40% of its peak, what is the realistic speedup over Path A?

Recall Solution

FLOPs = 2 × 2048³ = 2 × 8{,}589{,}934{,}592 = 17{,}179{,}869{,}184 ≈ 1.718 × 10¹⁰.
(a) Path A: 1.718×10¹⁰ / (12×10¹²) = 1.432×10⁻³ s ≈ 1.432 ms. Path B (peak): 1.718×10¹⁰ / (120×10¹²) = 1.432×10⁻⁴ s ≈ 0.1432 ms.
(b) Ideal speedup = 120 / 12 = 10× (also 1.432ms / 0.1432ms = 10).
(c) At 40% of peak, Path B effectively runs at 0.40 × 120 = 48 TFLOP/s. Time = 1.718×10¹⁰ / (48×10¹²) = 3.579×10⁻⁴ s ≈ 0.358 ms. Realistic speedup = 1.432 / 0.358 = 4×. Lesson: peak FLOP/s is a ceiling, not a promise. Whether you reach it depends on feeding the cores — see 6.2.8-Memory-hierarchy-in-GPUs and 9.1.5-Roofline-model.

Recall Quick self-test (cloze)

One Tensor Core instruction computes D = ==A × B + C== on a tile. Splitting a K-multiply into partial sums is valid because of the associativity of addition. A single FMA counts as 2 FLOPs. Tiling a matmul across all three axes reduces instruction count by (N/TILE)³. Mixed precision uses low-precision inputs but FP32 (high-precision) accumulation.

Related: 6.2.1-GPU-vs-CPU-architecture · 6.2.10-Warp-scheduling-and-execution · 7.3.6-cuBLAS-and-cuDNN · 6.2.12 Tensor cores and matrix operations (Hinglish)