Verified Deep Learning with Lean 4

3 MNIST: 2D CNN

The MLP wastes information, and we can do better

Chapter 2’s MLP recognized digits by flattening every 28\(\times \)28 image into a 784-dim vector and feeding it through three dense layers. It worked. But it worked by ignoring something the data was telling us. Pixel \((5, 5)\) and pixel \((5, 6)\) are neighbors: they were next to each other on the page where the digit was drawn. Pixel \((5, 5)\) and pixel \((23, 14)\) are not. The dense layer has no idea. It treats every input coordinate as an independent feature. Its weight matrix has no notion that some pairs of pixels are geometrically close. A trained MLP eventually learns to compensate, but it has to learn it from the data, with the data, by paying the full price for each unrelated feature pair.

A handwritten “5” drawn in the top-left of the image and the same “5” drawn in the bottom-right are the same digit. The MLP sees them as completely unrelated vectors. It can be trained to classify both, but each is its own learning problem, and nothing it learned about one transfers to the other for free.

Convolutional networks fix this by building three priors directly into the architecture, priors that the data has but the MLP threw away.

3.1 Run it first

Before any of the math, train the thing. Four commands and about a minute of GPU time:

lake exe cache get                            # Mathlib oleans, ~30 s
./scripts/datasets/download_mnist.sh          # ~11 MB
lake build mnist-cnn-verified
./.lake/build/bin/mnist-cnn-verified data

On one RTX 4060 Ti (CUDA 12.9), from runs/2026-09-01-cnn-verified-xla-cuda/. XLA’s startup banner is removed and one long line is wrapped:

[pjrt_ffi] XLA backend: PJRT 0.114, 1 device(s)
[pjrt_ffi] command buffers: enabled (CUDA default)
[pjrt_ffi] compiled verified_mlir/cnn_train_step.mlir
             (@cnn_train_step, 11 outputs, 1 replica) in 2566 ms
[pjrt_ffi] compiled verified_mlir/cnn_fwd.mlir
             (@cnn_fwd, 1 outputs, 1 replica) in 210 ms
MNIST-CNN via the VERIFIED renderer (conv→conv→pool→512→512→10) → XLA/PJRT → GPU
  xla/pjrt verified_mlir/cnn_train_step.mlir
  xla/pjrt verified_mlir/cnn_fwd.mlir
  train 60000, test 10000; bs 128, MNIST-CNN (10 params, 3489130 floats),
    mean-loss SGD lr=0.100000, He init
  epoch 1: loss = 0.352314, test_acc = 9295/10000 = 92.950000%
             [95% CI 92.43–93.44] (5219ms)
  epoch 2: loss = 0.094411, test_acc = 9749/10000 = 97.490000%
             [95% CI 97.16–97.78] (5072ms)
  epoch 3: loss = 0.056773, test_acc = 9800/10000 = 98.000000%
             [95% CI 97.71–98.26] (4888ms)
  epoch 4: loss = 0.038316, test_acc = 9819/10000 = 98.190000%
             [95% CI 97.91–98.43] (4672ms)
  epoch 5: loss = 0.026202, test_acc = 9838/10000 = 98.380000%
             [95% CI 98.11–98.61] (4357ms)
  epoch 6: loss = 0.018312, test_acc = 9853/10000 = 98.530000%
             [95% CI 98.27–98.75] (5069ms)
  epoch 7: loss = 0.022045, test_acc = 9857/10000 = 98.570000%
             [95% CI 98.32–98.78] (5537ms)
  epoch 8: loss = 0.011130, test_acc = 9859/10000 = 98.590000%
             [95% CI 98.34–98.80] (5390ms)
  epoch 9: loss = 0.006994, test_acc = 9869/10000 = 98.690000%
             [95% CI 98.45–98.89] (4808ms)
  epoch 10: loss = 0.006176, test_acc = 9861/10000 = 98.610000%
             [95% CI 98.36–98.82] (4351ms)
  epoch 11: loss = 0.004395, test_acc = 9867/10000 = 98.670000%
             [95% CI 98.43–98.88] (4647ms)
  epoch 12: loss = 0.033374, test_acc = 9877/10000 = 98.770000%
             [95% CI 98.53–98.97] (4981ms)
done (trained MNIST-CNN via the proof-rendered StableHLO).

Twelve epochs at about 4.9 seconds each, and 98.77% on the full 10,000-image test set. Two convolutions in front of Chapter 2’s classifier buy about a point over the MLP’s 97.81% (runs/2026-09-01-mlp-verified-xla-cuda/), and they do it with 3.5M parameters against the MLP’s 670K.

This is the noisiest of the three MNIST runs, and the transcript shows it. Epoch 1 landed at 92.95% here; repeated, the same twelve epochs from the same seed put it anywhere between 92.5 and 94.8%, because a convolution’s weight gradient accumulates in an order the GPU is free to vary. By epoch 12 that has largely washed out — but not in the shape the usual summary assumes, and this is the chapter to say what that means.

Run this net thirty times and twenty-nine of them land between 98.48 and 98.94%. The thirtieth lands at 97.39%. Figure 3-1 is those thirty beside the MLP’s thirty.

\begin{tikzpicture} 
\begin{axis}[
    width=0.92\linewidth, height=3.4cm,
    xlabel={Final test accuracy (\%), 30 runs each},
    xmin=97.2, xmax=99.05, ymin=-0.6, ymax=1.6,
    ytick={0,1}, yticklabels={CNN, MLP},
    xtick={97.25,97.5,97.75,98,98.25,98.5,98.75,99},
    grid=major, grid style={gray!18},
    tick label style={font=\small}, label style={font=\small},
]
\addplot[only marks, mark=*, mark size=1.4pt, orange, opacity=0.75]
coordinates {(97.39,0) (98.48,0) (98.59,0) (98.60,0) (98.67,0) (98.67,0) (98.71,0) (98.72,0) (98.72,0) (98.73,0) (98.74,0) (98.75,0) (98.76,0) (98.78,0) (98.78,0) (98.79,0) (98.81,0) (98.81,0) (98.81,0) (98.84,0) (98.85,0) (98.85,0) (98.85,0) (98.85,0) (98.87,0) (98.87,0) (98.88,0) (98.90,0) (98.91,0) (98.94,0)};
\addplot[only marks, mark=*, mark size=1.4pt, blue, opacity=0.75]
coordinates {(97.49,1) (97.55,1) (97.63,1) (97.63,1) (97.64,1) (97.65,1) (97.65,1) (97.69,1) (97.70,1) (97.72,1) (97.75,1) (97.76,1) (97.76,1) (97.76,1) (97.77,1) (97.78,1) (97.78,1) (97.79,1) (97.79,1) (97.80,1) (97.81,1) (97.83,1) (97.84,1) (97.85,1) (97.86,1) (97.86,1) (97.87,1) (97.90,1) (97.92,1) (97.93,1)};
\end{axis}
\end{tikzpicture}
Figure 3-1 Final test accuracy of thirty runs each of the MNIST MLP and CNN, one dot per run (runs/2026-09-01-mnist-n30/).

The MLP’s runs sit in a band; the CNN’s do too, apart from one.

Summarised as a mean and a standard deviation the CNN is \(98.73 \pm 0.27\). Summarised as a median and quartiles it is \(98.78\), with the middle half between \(98.72\) and \(98.85\). The second describes the picture above; the first does not, and the whole difference is that one run — trim the three highest and three lowest and the standard deviation falls from \(0.27\) to \(0.07\). A standard deviation assumes the spread is symmetric about the middle, and this one is not: the distribution fails a normality test outright (\(p {\lt} 10^{-4}\)), with a long left tail and nothing answering it on the right.

Three properties baked into a convolution

  1. Locality. Each output value depends on only a small neighborhood of inputs. A convolution with a \(3 \times 3\) kernel computes each output pixel as a weighted sum of its 9 neighbors. Faraway pixels do not enter directly. The network only learns to combine them later, by stacking convolutions and letting receptive fields grow.

  2. Translation equivariance. The same weights are applied at every spatial position, so a feature detector that recognizes the bottom-left curve of a “5” uses the same parameters whether the curve appears at row 5 or row 25, and the feature map shifts with the input. The architecture imposes the symmetry, and the data does not have to teach it. It is also why the layer is cheap: a \(3 \times 3 \times 1 \to 32\) convolution has \(3 \times 3 \times 1 \times 32 + 32 = 320\) parameters regardless of image size, where the matching fully-connected layer on a 28\(\times \)28 image would have \(28^2 \times 28^2 \times 32 \approx 20\) million.

  3. Hierarchy. Digits are made of strokes, and strokes of edges. Stacking convolutions lets each layer combine the previous one’s features over a growing neighborhood, and pooling between them discards where within a window a feature fired. That is where the invariance the “5” in the corner wanted comes from: not from any one layer, but from the stack.

Concretely, a convolution computes

\[ y_{c_{\text{out}},\, i,\, j} \; =\; \sum _{c_{\text{in}},\, k_h,\, k_w}\, W_{c_{\text{out}},\, c_{\text{in}},\, k_h,\, k_w} \cdot x_{c_{\text{in}},\, i + k_h,\, j + k_w} \; +\; b_{c_{\text{out}}}. \]

Each output cell is a small dot product between the kernel \(W\) and a window of the input. The sum is over a tiny \(3 \times 3\) spatial region and over input channels, usually a handful of terms rather than the 784 a dense layer touches.

The Jacobian is sparse and structured

The Jacobian of a convolution looks like a giant matrix, but it isn’t a generic giant matrix. It is mostly zeros. Specifically, \(\partial y_{c_{\text{out}}, i, j} / \partial x_{c_{\text{in}}, i', j'}\) is zero unless \((i', j')\) falls inside the kernel window centered at \((i, j)\). When it does, the entry is just the corresponding weight \(W_{c_{\text{out}}, c_{\text{in}}, k_h, k_w}\).

So a convolution is a linear map, like a dense layer, but its Jacobian is highly structured: each row has only \(k^2 \times C_{\text{in}}\) nonzero entries, and those entries are reused across many rows because the same kernel is used at every spatial position. The forward picture from Chapter 2 “the Jacobian of a linear map is the linear map” still holds. The Jacobian is just written in a different basis: instead of one big dense matrix it is a list of kernel weights plus a rule for where each one lives.

The backward of a convolution is another convolution

We want \(dx\), the way a small jiggle in each input pixel changes the loss, given an upstream \(dy\). Plugging into the Jacobian above and grinding through the algebra, the answer is:

\[ dx_{c_{\text{in}}, i', j'} \; =\; \sum _{c_{\text{out}},\, k_h,\, k_w}\, W_{c_{\text{out}},\, c_{\text{in}},\, k_h,\, k_w} \cdot dy_{c_{\text{out}},\, i' - k_h,\, j' - k_w}. \]

That is itself a convolution. Specifically, it is a convolution of the upstream gradient \(dy\) against the same kernel \(W\), but with the spatial axes reversed and the input/output channel axes swapped, which is a “transposed convolution”. The backward of a CNN block reuses the same kind of operation as the forward, and no new primitive is needed.

The weight gradient \(dW\) similarly works out to a convolution, this time of the saved input \(x\) against the upstream gradient \(dy\) (the “transpose trick”). The bias gradient \(db\) is just \(dy\) summed over spatial positions per channel. Theorems 35, 36, and 37 formalize each of these.

Max pooling: the winner gets the gradient

The other new primitive in this chapter is max pooling. A \(2 \times 2\) max pool collapses each \(2 \times 2\) block of the input into a single output value: the max of the four. It has no parameters, and its purpose is to halve the spatial resolution while keeping the strongest local response.

The Jacobian of \(\max \) is degenerate. At a generic point, only one of the four inputs is the maximum, and the output’s slope is \(1\) with respect to that one input and \(0\) with respect to the other three. “Jiggle the max” moves the output, and “jiggle a non-max” does nothing. So the backward of max pool routes the upstream gradient back to whichever input was the winner in each window, and zeros out the rest. At a tie the emitted select_and_scatter sends the whole cotangent to the first maximal cell, which is a different convention from ReLU’s mask, and the pdiv-derived backward is \(\operatorname {fderiv}\)’s zero map there, as at ReLU’s corner. Definition 39 states this formally.

The rest of the network is Chapter 2

After two convolutions and a max pool, the feature map is \(14 \times 14 \times 32\). At that point we flatten the \(14 \cdot 14 \cdot 32 = 6272\)-dim feature vector and feed it into the same three-dense-layer head from Chapter 2, ending in 10-class logits. The backward through the dense head and the softmax-CE loss is exactly the calculation of Chapters 1 and 2, and we do not re-prove it here. The composition rule from Chapter 1 guarantees that the full network’s VJP is just the per-layer VJPs strung together in reverse order: softmax-CE backward, dense backward \(\times 3\), flatten backward, max-pool backward, conv2d backward \(\times 2\). Each piece is its own theorem, and together they form a working trainer for MNIST that respects the image’s 2-D structure.

3.2 The theorems

Figure 3-2 is the dependency graph of this section.

\begin{tikzpicture} [x=1pt,y=1pt, every node/.style={inner sep=1.50pt, font=\fontsize{7.00}{8.40}\selectfont\ttfamily}]
  \node[draw=green!50!black, fill=green!15, rectangle, align=center, minimum width=31.2pt, minimum height=12.7pt, double] at (17.7pt,319.9pt) {\hyperref[ax:conv2d]{conv2d}};
  \node[draw=green!50!black, fill=green!15, rectangle, align=center, minimum width=35.4pt, minimum height=12.7pt, double] at (17.7pt,170.9pt) {\hyperref[def:hasvjp3]{hasvjp3}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=60.6pt, minimum height=12.7pt, double] at (188.2pt,261.9pt) {\hyperref[ax:conv2dHasVJP3]{conv2dHasVJP3}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=98.4pt, minimum height=12.7pt, double] at (188.2pt,344.9pt) {\hyperref[ax:conv2dWeightGradHasVJP]{conv2dWeightGradHasVJP}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=90.0pt, minimum height=12.7pt, double] at (188.2pt,317.9pt) {\hyperref[ax:conv2dBiasGradHasVJP]{conv2dBiasGradHasVJP}};
  \node[draw=green!50!black, fill=green!15, rectangle, align=center, minimum width=39.6pt, minimum height=12.7pt] at (87.2pt,141.9pt) {\hyperref[ax:maxPool2]{maxPool2}};
  \node[draw=green!50!black, fill=green!15, rectangle, align=center, minimum width=69.0pt, minimum height=12.7pt] at (188.2pt,53.9pt) {\hyperref[ax:maxPool2HasVJP3]{maxPool2HasVJP3}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=90.0pt, minimum height=12.7pt] at (314.4pt,53.9pt) {\hyperref[thm:mnistCnnNoBnHasVJPAt]{mnistCnnNoBnHasVJPAt}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=39.6pt, minimum height=12.7pt] at (314.4pt,372.9pt) {\hyperref[thm:cnn_fold]{cnn\_fold}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=60.6pt, minimum height=12.7pt] at (314.4pt,303.9pt) {\hyperref[thm:cnn_loss_grad]{cnn\_loss\_grad}};
  \node[draw=green!50!black, fill=green!15, rectangle, rounded corners=4.0pt, align=center, minimum width=73.2pt, minimum height=12.7pt] at (314.4pt,233.9pt) {\hyperref[thm:cnn_sgd_descends]{cnn\_sgd\_descends}};
  \node[draw=gray!70, dashed, rounded corners=2pt, minimum width=39.0pt, minimum height=17.8pt, text=gray!60!black, align=center, font=\fontsize{5.50}{6.60}\selectfont\ttfamily] at (188.2pt,172.9pt) {\hyperref[ax:pdiv_dense]{pdiv\_dense}\\(ch 1)};
  \node[draw=gray!70, dashed, rounded corners=2pt, minimum width=45.6pt, minimum height=17.8pt, text=gray!60!black, align=center, font=\fontsize{5.50}{6.60}\selectfont\ttfamily] at (188.2pt,463.9pt) {\hyperref[ax:pdiv_dense_W]{pdiv\_dense\_W}\\(ch 1)};
  \draw[->, >=stealth, black!70, line width=0.40pt] (33.7pt,314.8pt) .. controls (61.2pt,305.3pt) and (120.7pt,284.8pt) .. (157.0pt,272.4pt) -- (167.4pt,268.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (33.6pt,323.3pt) .. controls (43.2pt,325.4pt) and (56.0pt,328.1pt) .. (67.4pt,329.9pt) .. controls (86.8pt,333.2pt) and (108.2pt,336.0pt) .. (127.4pt,338.4pt) -- (138.7pt,339.7pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (33.7pt,319.8pt) .. controls (55.4pt,319.5pt) and (97.0pt,319.0pt) .. (131.5pt,318.6pt) -- (142.8pt,318.5pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (18.8pt,313.5pt) .. controls (19.6pt,276.7pt) and (29.4pt,93.0pt) .. (139.0pt,40.9pt) .. controls (186.4pt,18.4pt) and (249.3pt,32.2pt) .. (285.1pt,43.6pt) -- (295.8pt,47.1pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (21.1pt,326.5pt) .. controls (30.9pt,352.2pt) and (70.6pt,445.8pt) .. (139.0pt,479.9pt) .. controls (178.1pt,499.5pt) and (198.5pt,500.0pt) .. (237.4pt,479.9pt) .. controls (274.1pt,461.0pt) and (297.0pt,415.2pt) .. (307.1pt,390.1pt) -- (311.1pt,379.7pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (30.3pt,326.8pt) .. controls (40.0pt,332.2pt) and (54.2pt,339.5pt) .. (67.4pt,343.9pt) .. controls (139.2pt,368.2pt) and (165.2pt,381.0pt) .. (237.4pt,357.9pt) .. controls (261.7pt,350.2pt) and (284.9pt,331.5pt) .. (299.1pt,318.3pt) -- (307.1pt,310.6pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (26.8pt,313.2pt) .. controls (45.3pt,298.4pt) and (92.6pt,262.9pt) .. (139.0pt,248.9pt) .. controls (180.9pt,236.3pt) and (230.8pt,233.2pt) .. (266.5pt,232.8pt) -- (277.5pt,232.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (31.4pt,177.8pt) .. controls (59.8pt,193.1pt) and (129.0pt,230.5pt) .. (164.9pt,249.9pt) -- (174.9pt,255.3pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (21.7pt,164.2pt) .. controls (28.2pt,150.4pt) and (44.6pt,118.9pt) .. (67.4pt,100.9pt) .. controls (90.2pt,82.9pt) and (120.8pt,71.1pt) .. (145.3pt,63.8pt) -- (156.1pt,60.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (21.0pt,164.1pt) .. controls (30.3pt,137.4pt) and (68.3pt,40.7pt) .. (139.0pt,14.9pt) .. controls (180.1pt,0.0pt) and (194.9pt,4.7pt) .. (237.4pt,14.9pt) .. controls (258.3pt,20.0pt) and (279.9pt,31.8pt) .. (294.7pt,41.1pt) -- (304.1pt,47.3pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (20.0pt,177.8pt) .. controls (26.4pt,208.7pt) and (57.1pt,335.2pt) .. (139.0pt,382.9pt) .. controls (184.6pt,409.5pt) and (248.9pt,395.3pt) .. (285.3pt,383.5pt) -- (296.1pt,379.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (24.0pt,177.8pt) .. controls (39.5pt,197.0pt) and (86.3pt,251.5pt) .. (139.0pt,274.9pt) .. controls (182.0pt,294.1pt) and (236.1pt,300.7pt) .. (272.5pt,302.9pt) -- (283.7pt,303.5pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (23.8pt,164.2pt) .. controls (31.8pt,154.3pt) and (48.4pt,136.0pt) .. (67.4pt,128.9pt) .. controls (101.8pt,116.2pt) and (176.6pt,104.1pt) .. (237.4pt,135.9pt) .. controls (271.2pt,153.6pt) and (294.6pt,193.8pt) .. (305.8pt,216.8pt) -- (310.6pt,227.1pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (193.9pt,255.2pt) .. controls (202.7pt,242.4pt) and (222.2pt,213.8pt) .. (237.4pt,188.9pt) .. controls (262.9pt,147.1pt) and (290.6pt,96.5pt) .. (304.4pt,70.7pt) -- (309.8pt,60.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (196.7pt,268.6pt) .. controls (206.2pt,277.1pt) and (223.0pt,292.1pt) .. (237.4pt,304.9pt) .. controls (258.2pt,323.5pt) and (282.1pt,344.9pt) .. (297.5pt,358.7pt) -- (305.9pt,366.2pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (209.4pt,268.8pt) .. controls (229.2pt,275.5pt) and (259.5pt,285.7pt) .. (282.4pt,293.4pt) -- (293.2pt,297.1pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (219.0pt,255.2pt) .. controls (234.8pt,251.7pt) and (254.7pt,247.2pt) .. (272.0pt,243.3pt) -- (283.0pt,240.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (219.6pt,351.8pt) .. controls (238.9pt,356.1pt) and (263.7pt,361.7pt) .. (283.0pt,366.1pt) -- (294.2pt,368.6pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (213.8pt,338.1pt) .. controls (221.4pt,335.9pt) and (229.8pt,333.4pt) .. (237.4pt,330.9pt) .. controls (253.1pt,325.9pt) and (270.4pt,319.8pt) .. (284.5pt,314.7pt) -- (295.1pt,310.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (222.8pt,338.1pt) .. controls (228.0pt,336.3pt) and (233.0pt,333.9pt) .. (237.4pt,330.9pt) .. controls (239.5pt,329.5pt) and (280.3pt,276.9pt) .. (301.2pt,249.8pt) -- (308.1pt,240.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (216.6pt,324.8pt) .. controls (223.5pt,326.8pt) and (230.8pt,329.2pt) .. (237.4pt,331.9pt) .. controls (256.9pt,339.9pt) and (278.0pt,351.5pt) .. (293.0pt,360.4pt) -- (302.7pt,366.2pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (233.7pt,312.9pt) .. controls (246.5pt,311.5pt) and (260.4pt,309.9pt) .. (272.9pt,308.5pt) -- (284.0pt,307.3pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (199.4pt,311.1pt) .. controls (220.0pt,297.1pt) and (267.2pt,265.2pt) .. (294.0pt,247.1pt) -- (303.4pt,240.7pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (91.1pt,135.4pt) .. controls (97.7pt,121.6pt) and (115.0pt,89.0pt) .. (139.0pt,70.9pt) .. controls (141.8pt,68.8pt) and (144.9pt,67.0pt) .. (148.2pt,65.3pt) -- (158.7pt,60.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (105.5pt,135.2pt) .. controls (144.2pt,120.1pt) and (238.2pt,83.3pt) .. (285.4pt,64.9pt) -- (296.0pt,60.7pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (88.7pt,148.5pt) .. controls (91.6pt,181.6pt) and (105.9pt,329.3pt) .. (139.0pt,357.9pt) .. controls (178.8pt,392.4pt) and (244.6pt,387.1pt) .. (283.0pt,380.1pt) -- (294.2pt,377.8pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (94.7pt,148.8pt) .. controls (104.0pt,158.4pt) and (121.9pt,176.0pt) .. (139.0pt,188.9pt) .. controls (179.9pt,219.8pt) and (194.8pt,220.6pt) .. (237.4pt,248.9pt) .. controls (257.9pt,262.6pt) and (280.6pt,279.2pt) .. (295.9pt,290.6pt) -- (304.7pt,297.3pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (107.2pt,139.9pt) .. controls (136.7pt,137.6pt) and (194.4pt,136.4pt) .. (237.4pt,156.9pt) .. controls (266.0pt,170.6pt) and (290.0pt,199.8pt) .. (302.9pt,218.0pt) -- (309.3pt,227.3pt);
  \draw[->, >=stealth, black!70, line width=0.40pt] (223.1pt,53.9pt) .. controls (233.9pt,53.9pt) and (246.1pt,53.9pt) .. (258.0pt,53.9pt) -- (269.2pt,53.9pt);
  \draw[->, >=stealth, dashed, gray!60, line width=0.40pt] (208.2pt,180.3pt) .. controls (217.1pt,183.9pt) and (227.9pt,188.4pt) .. (237.4pt,192.9pt) .. controls (256.1pt,201.9pt) and (276.8pt,213.1pt) .. (291.9pt,221.6pt) -- (301.7pt,227.2pt);
  \draw[->, >=stealth, dashed, gray!60, line width=0.40pt] (211.1pt,456.7pt) .. controls (219.7pt,453.5pt) and (229.3pt,449.1pt) .. (237.4pt,443.9pt) .. controls (262.0pt,428.1pt) and (286.0pt,403.8pt) .. (300.2pt,388.1pt) -- (307.7pt,379.7pt);
\end{tikzpicture}
Figure 3-2 Dependency graph of the convolution chapter’s theorems.
Definition 33 Conv2d forward
✓
#

Concrete \(\sum _{c,kh,kw}\) cross-correlation with SAME padding. Codegen emits stablehlo.convolution; proofs reason about it via the explicit definition.

Definition 34 3D VJP record
✓
#

The rank-3 analogue of Definition 10. For \(f : \mathrm{Tensor3} \to \mathrm{Tensor3}\), define \(\operatorname {pdiv}_3 f\, x\, (c_i, h_i, w_i)\, (c_o, h_o, w_o)\) as \(\operatorname {pdiv}\) of the flattened map \(v \mapsto \mathrm{flatten}\bigl(f(\mathrm{unflatten}\, v)\bigr)\) with both indices decoded through the \((c, h, w) \leftrightarrow \text{flat}\) bijection. Then \(\mathsf{HasVJP3}\, f\) bundles a backward function \(B\) with: for all \(x\), \(dy\), \((c_i, h_i, w_i)\),

\[ B(x, dy)_{c_i h_i w_i} = \sum _{c_o, h_o, w_o} \operatorname {pdiv}_3 f\, x\, (c_i, h_i, w_i)\, (c_o, h_o, w_o) \cdot dy_{c_o h_o w_o}. \]
Theorem 35 Conv2d input VJP
✓

\(\mathsf{HasVJP3}\, \bigl(\mathrm{conv2d}(W, b)\bigr)\), with the input backward

\[ dx_{c_i h_i w_i} = \sum _{c_o, h_o, w_o} [\text{pad-valid}]\; W_{c_o, c_i, \hat{k}_h, \hat{k}_w} \cdot dy_{c_o h_o w_o}, \qquad \hat{k}_h = h_i + p_H - h_o,\; \; \hat{k}_w = w_i + p_W - w_o . \]

No hypotheses on the input or the kernel: conv2d is affine in its input. For odd \(k_H\), \(k_W\) this sum is the reversed-kernel, channel-swapped convolution of \(dy\) against \(W\) that the emitted backward runs (conv_back_bridge, §3.4); at an even kernel the reversed-kernel convolution is the adjoint of a convolution shifted by one pixel; the emitter pads such a kernel asymmetrically, which is the kernel zero-extended to the next odd size with symmetric padding, and the ties are stated through that zero-extension (padOdd, Chapter 8).

Proof ▶

Sketch: the same affine-rule skeleton as the dense Jacobians (Theorems 17, 18): the layer is linear in its input plus the bias, the Jacobian entry is the convolution of a basis vector, and that entry is a pad-conditional indicator which the closing collapse resolves.

  1. Define \(B(x, dy)_{c_i h_i w_i} := \sum _{c_o, h_o, w_o} [\text{pad-valid}]\; W_{c_o, c_i, \hat{k}_h, \hat{k}_w} \cdot dy_{c_o h_o w_o}\), where \(\hat{k}_h = h_i + p_H - h_o\), \(\hat{k}_w = w_i + p_W - w_o\) are the reconstructed kernel offsets and \([\text{pad-valid}]\) requires them to lie in \([0, k_H) \times [0, k_W)\). For odd kernels, where \(k_H - 1 - p_H = p_H\), the substitution \(k_h \leftrightarrow k_H - 1 - k_h\) turns this into the reversed-kernel, channel-swapped convolution of \(dy\) against \(W\) that the MLIR backward emits (transpose + reverse + convolution).

  2. suffices: for all \(x\), \(dy\), \((c_i, h_i, w_i)\): \(B(x, dy)_{c_i h_i w_i}\) equals the \(\operatorname {pdiv}_3\) triple sum.
    proof: Definition 34, with \(B\) as the candidate backward function.

  3. The flattened forward \(v \mapsto \mathrm{flatten}\bigl(\mathrm{conv2d}(W, b)\, (\mathrm{unflatten}\, v)\bigr)\) is affine in \(v\): the bias-free convolution \(\sum _{c, k_h, k_w} W_{o(\mathrm{idx}), c, k_h, k_w} \cdot (\text{pad-conditional projection of } v)\), additive and homogeneous in \(v\), plus the broadcast bias.
    proof: Unfold \(\mathrm{conv2d}\) and \(\mathrm{flatten}\) (Definition 33); the two side conditions distribute through the \((c, k_h, k_w)\) sum and the pad case split (Finset.sum_add_distrib, Finset.mul_sum).

  4. Per-index Jacobian: for every flat \(\mathrm{idx_{in}}\), \(\mathrm{idx_{out}}\), \(\operatorname {pdiv}\)(flattened forward) \(= \sum _{c, k_h, k_w} W_{o(\mathrm{idx_{out}}), c, k_h, k_w} \cdot [\text{pad-valid} \wedge \mathrm{idx_{in}} = \text{encoded input position}]\).
    proof: Affine rule (Theorem 9) with 3: the entry is the bias-free convolution of \(\mathbf{e}_{\mathrm{idx_{in}}}\), and a basis vector read at the pad-conditional position is exactly the indicator.

  5. Contract 4 with \(dy\) and collapse: re-index the sum over flat \(\mathrm{idx_{out}}\) to \((c_o, h_o, w_o)\); for each fixed output position the triple sum over \((c, k_h, k_w)\) collapses at \(c = c_i\), \(k_h = \hat{k}_h\), \(k_w = \hat{k}_w\).
    proof: The indicator in 4 is nonzero for exactly that one summand (padTap_indicator, Finset.sum_ite_eq’).

  6. q.e.d.
    proof: What remains after 5 is exactly \(B(x, dy)_{c_i h_i w_i}\) from 1; with 2 this is the goal.

Theorem 36 Conv2d weight VJP
✓

The transpose-trick formula. As with the dense weight Jacobian (Theorem 18), differentiate in the flattened kernel: let \(F(v) := \mathrm{flatten}\bigl(\mathrm{conv2d} (\mathrm{unflatten}\, v,\, b)\, x\bigr)\) for \(v \in \mathbb {R}^{o_c \cdot i_c \cdot k_H \cdot k_W}\). prove: \(\mathsf{HasVJP}\, F\), with backward \(dW_{o, c, k_h, k_w} = \sum _{h_i, w_i} [\text{pad-valid}]\; x_{c,\, h_i + k_h - p_H,\, w_i + k_w - p_W} \cdot dy_{o, h_i, w_i}\) — a convolution of the saved input against the upstream gradient.

Proof ▶

Sketch: the dual of Theorem 35 — the same affine decomposition, but now \(v\) (the kernel) is the variable and the saved input \(x\) is the constant coefficient.

  1. Define \(B(v, dy)_{\mathrm{idx}}\), decoding \(\mathrm{idx} \leftrightarrow (o', c', k_h', k_w')\), as \(\sum _{h_i, w_i} (\text{pad-conditional } x \text{ value at } (c', h_i + k_h' - p_H, w_i + k_w' - p_W)) \cdot dy_{\mathrm{flat}(o', h_i, w_i)}\).

  2. suffices: for all \(v\), \(dy\), \(\mathrm{idx}\): \(B(v, dy)_{\mathrm{idx}} = \sum _{\mathrm{idx_{out}}} \operatorname {pdiv}F\, v\, \mathrm{idx}\, \mathrm{idx_{out}} \cdot dy_{\mathrm{idx_{out}}}\).
    proof: Definition 10, with \(B\) as the candidate backward function.

  3. \(F\) is affine in \(v\): the bias-free convolution \(\sum _{c, k_h, k_w} v_{\varphi (o(\mathrm{idx_{out}}), c, k_h, k_w)} \cdot (\text{pad-conditional } x \text{ term})\), additive and homogeneous in \(v\), plus the broadcast bias, where \(\varphi \) is the kernel-index bijection (Kernel4.flatten).
    proof: Unfold \(\mathrm{conv2d}\) and the flattens (Definition 33); add_mul, Finset.sum_add_distrib, Finset.mul_sum.

  4. Per-index Jacobian: \(\operatorname {pdiv}F\, v\, \mathrm{idx}\, \mathrm{idx_{out}} = \sum _{c, k_h, k_w} [\mathrm{idx} = \varphi (o(\mathrm{idx_{out}}), c, k_h, k_w)] \cdot (\text{pad-conditional } x \text{ term})\).
    proof: Affine rule (Theorem 9) with 3: the bias-free convolution with the basis kernel \(\mathbf{e}_{\mathrm{idx}}\) has one nonzero tap.

  5. Contract 4 with \(dy\) and collapse: the \((c, k_h, k_w)\) sum collapses on \(\mathrm{idx}\)’s decode \((c', k_h', k_w')\) by injectivity of \(\varphi \); re-index \(\mathrm{idx_{out}}\) to \((o, h_i, w_i)\) and collapse the channel sum at \(o = o'\).
    proof: The indicator forces \(o = o'\); the spatial sums survive.

  6. q.e.d.
    proof: What remains after 5 is exactly \(B(v, dy)_{\mathrm{idx}}\) from 1; with 2 this is the goal.

Theorem 37 Conv2d bias VJP
✓

Sum the cotangent over spatial dims per channel. prove: \(\mathsf{HasVJP}\, \bigl(b \mapsto \mathrm{flatten} (\mathrm{conv2d}(W, b)\, x)\bigr)\), with backward \(db_o = \sum _{h_i, w_i} dy_{\mathrm{flat}(o, h_i, w_i)}\).

Proof ▶
  1. Define \(B(b, dy)_o := \sum _{h_i, w_i} dy_{\mathrm{flat}(o, h_i, w_i)}\).

  2. suffices: for all \(b\), \(dy\), \(o\): \(B(b, dy)_o = \sum _{\mathrm{idx}} \operatorname {pdiv}\bigl(b' \mapsto \mathrm{flatten}(\mathrm{conv2d}(W, b')\, x)\bigr)\, b\, o\, \mathrm{idx} \cdot dy_{\mathrm{idx}}\).
    proof: Definition 10, with \(B\) as the candidate backward function.

  3. As a function of \(b\), the flattened forward is \((\text{channel reindex of } b) + (\text{$W\! ,x$ term constant in } b)\), so \(\operatorname {pdiv}= \delta _{o,\, \mathrm{chan}(\mathrm{idx})}\).
    proof: Affine rule (Theorem 9): the channel broadcast \(b' \mapsto (\mathrm{idx} \mapsto b'_{\mathrm{chan}(\mathrm{idx})})\) is linear, and \(\mathbf{e}_o\) broadcast is that indicator.

  4. q.e.d.
    proof: Contract 3 with \(dy\): re-index the flat sum to \((c, h_i, w_i)\); the Kronecker keeps exactly channel \(o\)’s spatial positions (Finset.sum_ite_eq’), leaving \(B(b, dy)_o\); with 2 this is the goal.

Definition 38 MaxPool 2x2 stride 2 forward
✓
#

Concrete four-way max over 2x2 windows.

Definition 39 MaxPool input VJP
✓

The max pool’s VJP record, with the \(\operatorname {pdiv}\)-derived backward: at a point where every window has a strict maximum, each window’s cotangent goes to that cell; at a point where any window ties at its maximum the map is not differentiable, so the backward is the zero tensor in every window, strict ones included. The emitted backward routes a tied window to its first maximal cell.

Those are the primitives. The rest of the section is the network: its VJP at a smooth input, then the three certificates about training it — the emitted step is the certified step, each gradient node is the loss’s derivative, and one SGD step descends.

assume:

  1. no coordinate of either convolution’s pre-activation, nor of either hidden dense layer’s, is zero [h1, h2, h3, h4]

  2. every \(2 \times 2\) window the max pool sees has a strict maximum [h_mp]

prove: the backward composed from the six per-layer backwards in reverse order — dense, dense+ReLU, dense+ReLU, max pool, conv+ReLU, conv+ReLU — satisfies, for all \(dy\), \(i\):

\[ \mathrm{backward}(dy)_i = \sum _j \operatorname {pdiv}(\texttt{mnistCnnNoBnForward})\, x\, i\, j \cdot dy_j . \]
Proof ▶

Chain rule at a point through the six stages; see Proofs.mnistCnnNoBnHasVJPAt_correct.

prove: all ten parameter updates cnn_train_step.mlir emits — the dense head’s \(W_5, b_5, W_4, b_4, W_3, b_3\) and the convolutions’ \(W_2, b_2, W_1, b_1\) — denote, at a batch of \(N\) examples, the certified SGD step at the real forward: each update is the parameter minus \(\alpha \) times the sum over the batch of its layer’s VJP at example \(k\), contracted with example \(k\)’s chain cotangent. That cotangent starts as the softmax cross-entropy gradient at the logits; the dense head carries it back to the max pool, the max pool to \(\mathrm{conv}_2\), and \(\mathrm{conv}_2\) to \(\mathrm{conv}_1\). At the batched index the chain is the emitted backward subgraph itself — the selects, the select_and_scatter and the reversed-kernel convolution — and its slices are the hand-named chain cotangents of the single-example reading (cnnChainCotW2 outright, cnnChainCotW1 for odd kernels, the artifact’s being \(3\times 3\)). For the output weight \(W_5\) the term is \(\partial L_k/\partial W_5\), the gradient of example \(k\)’s whole loss, as for the MLP. For every \(N\); the artifact is \(N = 128\).

Proof ▶

A fold over the ten updates: each denotes its parameter minus \(\alpha \) times the layer’s gradient — the dense ones of Theorems 20 and 21, the convolutions’ of Theorems 36 and 37 — at the cotangent the chain hands it; see Proofs.CnnFoldB.cnn_train_step_tiedB.

assume:

  1. odd kernels [hkH, hkW]

  2. at \(x\), every ReLU is off its corner, and every pooling window is entirely \(\le 0\) before its ReLU or has its maximum strictly above every other cell except its twins: cells equal at every value of both convolutions’ weights and biases, which two cells with identical receptive fields are — the tie a blank stretch of an MNIST image produces [hx]

  3. \(\sigma \) names a maximal cell of each window

  4. \(L\) is any loss of the logits, with gradient \(g\) there [hL]

prove: each of the ten gradient nodes, at the chain cotangent with each window’s cotangent routed to the cell \(\sigma \) names, is \(\partial L/\partial \theta \) of mnistCnnNoBnForward with that one parameter varied. Instantiated at the softmax cross-entropy with the emitted loss cotangent; and at \(\sigma \) the first maximal cell of each window — the routing the emitted select_and_scatter uses — the chain is Theorem 41’s, so the emitted updates step along these gradients, ties included.

Proof ▶

Along any parameter, near the given point, the pooled ReLU is the gather at \(\sigma \): a dead window stays negative, a strict maximum stays strict, and a twin stays tied. The gather is linear, so the chain rule runs on it, and the gradient moves to the real net with the germ; see Proofs.CnnFold.cnn_net_lossGrad.

assume:

  1. the second convolution’s input and the dense weights are bounded, \(|x_1| \le a\) and \(|W_k| \le w_k\) [hx, hW\(_3\), hW\(_4\), hW\(_5\)]

  2. a margin at every ReLU after \(\mathrm{conv}_2\) and in the dense head [hm2, hm3, hm4]

  3. in every live pooling window the maximum clears the other cells by a margin, up to twins: cells that read the same input patch, so are equal at every kernel — a blank corner of an MNIST image is one [hmq, hT]

  4. an \(\eta \)-accurate gradient oracle, and the small-step and dominance conditions in the shape of Theorem 25, with the chain’s own constants: the small step is \(2\, w_5 d_4 w_4 d_3 w_3\, (4hw)\, a D {\lt} 1\), the Lipschitz constant carries the same product squared with the class count, and the first convolution’s adds the kernel’s fan-in [hgh, hsmall, h1, h2]

  5. the bounds, \(\alpha \) and \(\eta \) are non-negative [ha, hw\(_3\), hw\(_4\), hw\(_5\), hlr, h\(\eta \)]

prove: one SGD step on the second convolution’s kernel, at one example with the other parameters fixed, decreases the cross-entropy loss:

\[ L(W_2 - \alpha g) \le L(W_2) - \alpha \, \| \nabla L\| _2^2/2 . \]

The same holds at the first convolution and at both convolutions’ biases, and with the binary32 gradient in place of \(g\), its \(\eta \) the rounding budget of §B.5 and its pool margin strict. Witness: a trained reduced-width network and a real MNIST test image meet every hypothesis, at both kernels and both biases.

Proof ▶

Within the step a tied window stays tied and a strict one stays strict, so the pooled ReLU is a fixed gather; the chain from the kernel to the logits is then linear up to its ReLU masks, and the descent argument of Theorem 25 runs with the margins supplying its Lipschitz constant; see Proofs.cnn_conv2_sgd_descends.

3.3 Example: MNIST 2D CNN

Chapter 2’s MLP flattened the image and threw dense layers at it. This chapter does MNIST again with the structure kept, the way the first book’s Chapter 2 did: two plain \(3 \times 3\) convs to pull out local features, a max-pool to collapse the spatial grid, then the same dense head you already know.

Architecture

Figure 3-3 is the network.

\begin{tikzpicture} [
  >={Stealth[length=1.8mm]},
  every node/.style={font=\sffamily\scriptsize},
  col/.style    = {align=center, rounded corners=2pt, inner sep=2pt, minimum height=0.62cm, minimum width=4.0cm},
  io/.style     = {col, draw=blue!55!black,   fill=blue!8},
  conv/.style   = {col, draw=orange!65!black, fill=orange!12},
  pool/.style   = {col, draw=teal!60!black,   fill=teal!10},
  flat/.style   = {col, draw=purple!60!black, fill=purple!8},
  dense/.style  = {col, draw=orange!65!black, fill=orange!12},
  head/.style   = {col, draw=red!60!black,    fill=red!8},
  logits/.style = {col, draw=green!50!black,  fill=green!14, very thick},
  arr/.style    = {->, thick, gray!60, shorten >=1pt, shorten <=1pt},
]
  % Vertical layer column: input at top -> logits at bottom.
  \node[io]                          (input) {Input \;\; $28\times28\times1$};
  \node[conv, below=0.24cm of input] (c1)    {\textbf{Conv2D} $1\to32$, $3\times3$, ReLU};
  \node[conv, below=0.24cm of c1]    (c2)    {\textbf{Conv2D} $32\to32$, $3\times3$, ReLU};
  \node[pool, below=0.24cm of c2]    (p1)    {\textbf{maxPool} $2\times2$ \;\; $28\to14$};
  \node[flat, below=0.24cm of p1]    (fl)    {flatten \;\; $14\times14\times32 \to 6272$};
  \node[dense, below=0.24cm of fl]   (d1)    {\textbf{Dense} $6272\to512$, ReLU};
  \node[dense, below=0.24cm of d1]   (d2)    {\textbf{Dense} $512\to512$, ReLU};
  \node[head, below=0.24cm of d2]    (d3)    {\textbf{Dense} $512\to10$ \;(identity)};
  \node[logits, below=0.24cm of d3]  (out)   {Logits \;\; 10 classes, softmax-CE};
  \foreach \a/\b in {input/c1, c1/c2, c2/p1, p1/fl, fl/d1, d1/d2, d2/d3, d3/out}
     \draw[arr] (\a) -- (\b);
\end{tikzpicture}
Figure 3-3 The MNIST CNN: two \(3\times 3\) convolutions and a max-pool in front of the MLP’s dense head.

Two convolutions lift the 1-channel grayscale input to 32 channels, maxPool halves the spatial resolution, then the flattened \(14 \times 14 \times 32 = 6272\) feature vector goes through three dense layers into 10-class logits. No batch norm yet, because BN shows up in Chapter 4.

Code: the full training program

Same format as the MLP chapter. The only thing that changes from Chapter 2’s training program is the layers list. This is the spec §3.1 actually ran.

-- 1
import LeanMlir.Verified.NetsCore
import LeanMlir.Verified.Train

-- 2
def cnnVerified : VerifiedNetSpec where
  name     := "MNIST-CNN"
  slug     := "cnn"
  inC      := 1
  imageH   := 28
  imageW   := 28
  data     := .mnist
  layers   := [.conv 1 32 3 1, .relu, .conv 32 32 3 1, .relu, .maxPool 2 2,
               .flatten,
               .dense 6272 512, .relu, .dense 512 512, .relu, .dense 512 10]
  blurb    := "MNIST-CNN via the VERIFIED renderer
                (conv→conv→pool→512→512→10) → %LOWERER% → GPU"
  lossSlot := true

-- 3
def cnnConfig : VerifiedConfig where
  epochs    := 12
  batchSize := 128

-- 4
def main (argv : List String) : IO Unit :=
  cnnVerified.toNet.train cnnConfig (argv.head?.getD "data")

Walking through the numbered sections:

1. The import. Same as Chapter 2.

2. The architecture. Eleven entries: two .conv layers at 32 channels, a .maxPool, a .flatten, three .dense layers down to 10, and four .relus between them. .conv ic oc k stride takes in-channels, out-channels, kernel size and stride — no activation argument, because activations are their own entries here rather than a field on the layer that precedes them. That is why Chapter 2’s five-entry list became eleven and not seven.

The .conv2d backward pass is the theorem from earlier in this chapter (Theorem 35): the input VJP is another convolution, using the kernel reversed along both spatial axes. The max-pool backward (Definition 39) routes the gradient to the argmax position in each \(2 \times 2\) window. The dense backward is Theorem 19. Adding convolution to the model introduces no new training-loop code, only new layer types, which the framework already knows how to compose.

Two fields carry over from Chapter 2 and do the same jobs: blurb is the banner the run prints, and lossSlot declares that this render returns a trailing report-only %loss, which is the one line of verified_mlir/cnn_train_step.mlir that is not pretty of a proven graph node — the render labels it as such in a comment of its own. Everything else in the file is: Proofs.StableHLO.cnnTrainStepText emits it, and CnnFold proves that each of the ten parameter updates denotes the certified softmax cross-entropy loss-descent step derived from the Mathlib \(\operatorname {fderiv}\) math. Both convolutions, the max pool, the three dense layers and the SGD update are proof-backed operations.

3. The training hyperparameters. Twelve epochs at batch 128, SGD, no regularization and no augmentation. Note what a VerifiedConfig does not carry: the learning rate. It is baked into verified_mlir/cnn_train_step.mlir, which is §1.6’s point in a second chapter — the rate is inside the reach of the proofs, so it cannot be set to one thing and trained at another.

4. The program entry point. One line, same as the MLP chapter. cnnVerified.toNet.train cnnConfig replaces mlpVerified.toNet.train mlpConfig. That’s the whole diff.

What .conv, .maxPool and .flatten contribute.

Three more arms on Chapter 1’s toSpecs, and only one of them carries parameters:

  | conv ic oc k _ => #[(#[oc,ic,k,k],0),(#[oc],2)]   -- W, bias
  | maxPool _ _    => #[]
  | flatten        => #[]

Pooling and flattening are pure reshapes of data, so they have nothing to initialise — which is the parameter-list statement of the same fact the backward passes above make: maxPool2’s VJP routes gradient without owning any, and flatten’s is a reshape. Fold the eleven-entry list through those arms and you get the 10 params, 3489130 floats the run printed: two conv kernels and their biases, three dense matrices and theirs.

Results

The run in §3.1 is the result. Figure 3-4 is its per-epoch training loss, against Chapter 2’s MLP on the same recipe.

\begin{tikzpicture} 
\begin{axis}[
    width=0.92\linewidth, height=6.5cm,
    xlabel={Epoch}, ylabel={Training loss},
    ymode=log,
    xmin=0, xmax=13, ymin=0.003, ymax=0.5,
    xtick={0,2,4,6,8,10,12},
    legend pos=north east,
    legend cell align={left},
    grid=major, grid style={gray!18},
    tick label style={font=\small},
    label style={font=\small},
    every axis plot/.append style={line width=1pt, mark size=1pt},
]
\addplot[blue, mark=*, mark options={fill=blue}] coordinates {
(1,0.413496) (2,0.191186) (3,0.139057) (4,0.109009) (5,0.088619) (6,0.073548) (7,0.061923) (8,0.052682) (9,0.045157) (10,0.038872) (11,0.033574) (12,0.029061)
};
\addlegendentry{MLP (12 epochs)}
\addplot[orange, mark=*, mark options={fill=orange}] coordinates {
(1,0.352314) (2,0.094411) (3,0.056773) (4,0.038316) (5,0.026202) (6,0.018312) (7,0.022045) (8,0.011130) (9,0.006994) (10,0.006176) (11,0.004395) (12,0.033374)
};
\addlegendentry{CNN (12 epochs)}
\end{axis}
\end{tikzpicture}
Figure 3-4 Mean training loss per epoch of the MNIST CNN against the MLP, both on SGD 0.1 for twelve epochs, on a log scale (runs/2026-09-01-{mlp,cnn}-verified-xla-cuda/).

The CNN sits below the MLP through epoch 11 and reaches by epoch 5 a loss the MLP has not reached by twelve. Test accuracy is 98.77% against the MLP’s 97.81%.

The exception is the last point, and it is the interesting one. The CNN’s training loss bumps at epoch 7 and then jumps almost eightfold at epoch 12, from \(0.0044\) to \(0.0334\) — far enough to cross back above the MLP’s \(0.0291\). That is ordinary SGD at \(\text{lr}=0.1\) rather than a measurement artifact. And the run’s test accuracy is at its highest on that same epoch, 98.77%. Training loss and held-out accuracy are different measurements, and a figure where the two disagree this loudly is the one to remember when a later chapter asks you to read a loss curve.

Both curves come off the binaries these chapters’ theorems describe.

Why this network is still 99% dense-head

At 3,489,130 total parameters, the no-BN CNN is actually \(\sim \)5.2\(\times \) bigger than the MLP, not smaller. You might have expected the opposite, because convolutions are supposed to be parameter-efficient. They are, but the head is what’s eating the budget: the flatten of a \(14 \times 14 \times 32\) feature map produces a 6272-dim vector, and the first dense layer (6272 \(\to \) 512) alone is 3,211,776 parameters. The entire conv+maxPool backbone is 9,568 parameters, which is 0.27% of the total.

This is the same “fat FC head” pattern you’ll see in AlexNet, VGG, and YOLOv1: a compact convolutional backbone spending its parameter budget on a large fully-connected classifier. ResNet (Chapter 5) is the first architecture in this book without one: its head is a globalAvgPool followed by a single small dense layer, which is why ResNet-18 (a smaller version of ResNet-34) has \(\sim \)11M parameters for 1000-class ImageNet but Chapter 3’s 10-class MNIST CNN has 3.5M.

The fat head is fat because it can be, not because it must be

If the dense head holds 99% of the parameters, the obvious question is how many of them earn their keep. The parametric renderer answers it directly: mnist-cnn-grid holds the convolutional backbone fixed at 32 channels and sweeps only the dense head (\(\ldots {\to }d{\to }d{\to }10\)), rendering each width from the faithful CNN emitter. Figure 3-5 overlays that sweep on the MLP’s (both trained on their own verified StableHLO), and shows how little the head needs once good convolutional features exist.

\begin{tikzpicture} 
\begin{axis}[
    width=0.92\linewidth, height=6.8cm,
    xlabel={Hidden width $d$ (dense head / MLP hidden layers)},
    ylabel={MNIST test accuracy (\%)},
    xmode=log, log basis x=2,
    xmin=6.5, xmax=5000, ymin=90.5, ymax=99.6,
    xtick={8,16,32,64,128,256,512,1024,2048,4096},
    xticklabels={8,16,32,64,128,256,512,1024,2048,4096},
    ytick={91,92,93,94,95,96,97,98,99},
    legend pos=south east, legend cell align={left},
    legend style={font=\footnotesize, draw=none, fill=none},
    grid=major, grid style={gray!18},
    tick label style={font=\small},
    label style={font=\small},
    every axis plot/.append style={line width=1pt, mark size=1.5pt},
]
\addplot[green!55!black, mark=*, mark options={fill=green!55!black},
  error bars/.cd, y dir=both, y explicit,
  error bar style={green!45!black, line width=0.6pt},
  error mark options={rotate=90, mark size=2.5pt, green!45!black, line width=0.6pt}]
coordinates {
(8,98.17) += (0,0.03) -= (0,0.18) (16,98.49) += (0,0.14) -= (0,0.02) (32,98.46) += (0,0.28) -= (0,0.09) (64,98.47) += (0,0.16) -= (0,0.15) (128,98.65) += (0,0.08) -= (0,0.05) (256,98.77) += (0,0.04) -= (0,0.15) (512,98.49) += (0,0.17) -= (0,1.83) (1024,98.76) += (0,0.08) -= (0,0.04) (2048,98.71) += (0,0.08) -= (0,0.06) (4096,98.75) += (0,0.06) -= (0,5.23)
};
\addlegendentry{CNN (conv@32 fixed, head $d{\to}d$)}
\addplot[blue, mark=square*, mark options={fill=blue}] coordinates {
(8,91.99) (16,94.01) (32,96.76) (64,97.02) (128,97.25) (256,97.67) (512,97.83) (1024,97.89) (2048,98.00) (4096,98.01)
};
\addlegendentry{MLP ($784{\to}d{\to}d{\to}10$)}
\end{axis}
\end{tikzpicture}
Figure 3-5 Test accuracy against dense-head width for the verified CNN, the convolutional backbone held at 32 channels, overlaid on the MLP width sweep of Chapter 2 (runs/cnn_grid_rep{1..5}.tsv, runs/mlp_grid_results_xla.tsv). Each CNN point is the median of five runs and its whiskers the observed range.

The whiskers are not an interval, just the highest and lowest the configuration actually produced. Eight of the ten widths repeat to within 0.4 and their whiskers are barely visible. The other two do not: \(d=512\) dropped one run to 96.66% and \(d=4096\) one to 93.52%, which is what the two long tails are.

Those tails are the reason this curve is a median of five. A symmetric error bar would have been the wrong shape for them — four runs cluster and one falls away, so there is no spread to put a \(\pm \) on — and the 95% interval on the test set, \(\pm 0.22\) at this accuracy, would have hidden a five-point failure behind a rounding error. The MLP points are single runs, which Chapter 2 reads with half a point of give.

The cost is that a handful of runs cannot see this. Draw five of the thirty runs in §3.1, build the ordinary 95% interval on their mean, and ask whether it contains the mean of the twenty-five you did not draw: it does 79% of the time, not 95%, because five draws usually miss the tail and report a spread as small as \(0.05\) against the true \(0.27\). Chapter 2’s MLP, whose thirty runs are unremarkable, gets 92% from the same procedure. How many runs a claim needs is a property of the network, and a convolution’s is higher than a dense stack’s. The early epochs of a conv net are not a measurement; the last one is, and even then not from a single run.

The whole CNN curve sits about a point above the MLP’s, and it is flat across the entire range: an 8-wide head already reaches 98.2% at 60K parameters, beating the widest 4096-wide MLP. So the 3.2M-parameter \(6272{\to }512\) layer is not buying accuracy, because shrinking the head to \(\ldots {\to }8{\to }8{\to }10\) keeps essentially all of it. The head is fat because the flatten hands it a 6272-vector and nobody trimmed it, not because the task needs the width.

3.4 MLIR: Convolution

What is already proven. mnistCnnNoBnForward is the conv \(\to \) relu \(\to \) conv \(\to \) relu \(\to \) maxPool \(\to \) flatten \(\to \) dense network, and mnistCnnNoBnHasVJPAt (Theorem 40) proves its reverse-mode derivative by chaining the per-layer VJPs: the dense backward (Theorem 19), the max-pool backward (maxPool2HasVJP3, Definition 39), and the convolution backward (conv2dHasVJP3, Theorem 35), all through the IR-level chain rule vjpCompAt. The convolution’s reverse-mode derivative is the chapter’s headline result: the input VJP of a conv2d is another conv2d, against the kernel with its in- and out-channels swapped and both spatial axes reversed. The parameter gradients are pinned by conv2dWeightGradHasVJP (the transpose trick) and the dense grads (Theorems 20, 21), and the loss gradient by softmaxCE_grad (Theorem 22). So what the convolution’s gradient is, is not in question.

The gap and how we close it. The emitted backward convolution is given a denotation \([\! [\cdot ]\! ]\) valued in the proofs’ own tensor type,

\[ [\! [\texttt{convBackDenote}\, W]\! ] = \texttt{conv2d}\, (\texttt{reverseSwap}\, W), \qquad (\texttt{reverseSwap}\, W)_{c_i\, c_o\, k_h\, k_w} = W_{c_o\, c_i\, \bar k_h\, \bar k_w}, \]

where \(\bar k = K - 1 - k\) is the spatial flip, and the bridge theorem

\[ \texttt{conv\_ back\_ bridge}\; (k_H, k_W \text{ odd}):\quad [\! [\texttt{convBackDenote}\, W]\! ](dy) = (\texttt{conv2dHasVJP3}\, W\, b).\mathrm{backward}(x, dy) \]

says the denotation of that graph is the proven conv input-VJP. It holds at every channel count, image size and odd kernel size at once: per output coordinate the reversed-kernel sum and the gradient formula of Theorem 35 are the same finite sum under the reindexing \(k \leftrightarrow K-1-k\) (reverseSlab_eq_gradSlab), and that reindexing is a bijection exactly when \(K - 1 - p = p\), which is oddness (hkH, hkW).

The committed step of the chapter’s network, verified_mlir/cnn_train_step.mlir, carries that graph at the trained shape, \(1{\to }32{\to }32\) channels at \(28\times 28\), batch 128. Here are the three operators the chapter added, as the backward pass meets them: the max-pool, the second ReLU, and the second convolution’s input gradient (SSA names the artifact’s, default precision dropped, lines wrapped, the select_and_scatter body elided):

%v50 = stablehlo.reshape %v17 : (tensor<128x25088xf32>)
      -> tensor<128x32x28x28xf32>
%v51 = stablehlo.reshape %v49 : (tensor<128x6272xf32>)
      -> tensor<128x32x14x14xf32>
%v52 = stablehlo.constant dense<0.0> : tensor<f32>
%v53 = "stablehlo.select_and_scatter"(%v50, %v51, %v52) ({
// ...
%v58 = stablehlo.compare GT, %v56, %v57 : (tensor<128x32x28x28xf32>,
      tensor<128x32x28x28xf32>) -> tensor<128x32x28x28xi1>
%v59 = stablehlo.select %v58, %v55, %v57 : tensor<128x32x28x28xi1>,
      tensor<128x32x28x28xf32>
// ...
%v62 = stablehlo.transpose %W2, dims = [1, 0, 2, 3]
      : (tensor<32x32x3x3xf32>) -> tensor<32x32x3x3xf32>
%v63 = stablehlo.reverse %v62, dims = [2, 3] : tensor<32x32x3x3xf32>
%v64 = stablehlo.convolution(%v61, %v63)
      dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1],
      window = {stride = [1, 1], pad = [[1, 1], [1, 1]],
      lhs_dilate = [1, 1], rhs_dilate = [1, 1]}
      {batch_group_count = 1 : i64, feature_group_count = 1 : i64}
      : (tensor<128x32x28x28xf32>, tensor<32x32x3x3xf32>)
      -> tensor<128x32x28x28xf32>

Read it against the graph, bottom up. %v62 is the transpose dims [1,0,2,3] that swaps the kernel’s out- and in-channel axes, and %v63 is the reverse dims [2,3] that flips both spatial axes (\(\bar k_h, \bar k_w\)). Together they build reverseSwap \(W_2\). The convolution at %v64 convolves the incoming cotangent %v61 against that kernel, exactly \(\texttt{conv2d}\, (\texttt{reverseSwap}\, W_2)\), the \(32 \leftarrow 32\) adjoint of the forward conv, under the same pad = [[1, 1], [1, 1]] the forward used. Above it, %v58/%v59 is the ReLU mask at the saved pre-activation %v56, the second convolution’s output, and %v53 is the max-pool backward: select_and_scatter walks each \(2\times 2\) window of the pre-pool activation %v50 with a GE select, which picks the window’s first maximal cell, and scatters the pooled cotangent %v51 onto it with add. That is the lowest-index tie-break of Definition 39 written as an operation.

Three theorems cover the three. conv_back_bridge is the convolution; maxPoolBack_faithful says the select_and_scatter denotes maxPoolFlatHasVJPAt’s backward wherever every window has a single maximum (MaxPool2Smooth); the ReLU mask is Chapter 2’s relu_back_bridge. Theorem 41 then composes them at batch 128. Its slice lemmas identify example \(k\)’s slice of %v61, the cotangent entering the second convolution’s backward, with the certified chain (dhc2B_slice), and the slice of %v70, the cotangent at the first convolution’s output once the first ReLU’s mask %v69 is applied to %v64, with the certified chain through conv_back_bridge (dhc1B_slice), which is where the odd-kernel hypothesis enters. The cotangent each convolution’s weight update is stated at is the slice of these emitted operations, not a hand-named value beside them. One detail of printing: this per-example renderer writes transpose then reverse, and the batched renderers of Chapters 5 onward write reverse then transpose. The two kernels are the same tensor, since swapping the channel axes and flipping the spatial ones commute, and the batched graph Theorem 41 reads this file at carries a descriptor (convBackP) for exactly this order, so that the printed text does not change.

Caveats.

  • The convolution bridge holds at every input, for odd kernels. A convolution is linear, so there is no hypothesis on \(x\); the one hypothesis is on the kernel shape, \(k_H\) and \(k_W\) odd, and every kernel in Chapters 3–7 is. Where a net has an even kernel (ConvNeXt’s \(4\times 4\) stem and \(2\times 2\) downsamples, Chapter 8) the tie passes the kernel zero-extended to the next odd size (padOdd), which is the asymmetric pad the emitter writes there.

  • ReLU is a smooth-point bridge (no pre-activation exactly zero), stated only off that set, which is null for generic weights.

  • Max-pool is a smooth-point bridge too. It holds where every \(2\times 2\) window attains its maximum at one cell (MaxPool2Smooth); the other cells may tie with each other. An all-zero post-ReLU window fails it, so for pooling the excluded set is not null the way ReLU’s corner is, and the whole-net theorem states it as a hypothesis (h_mp in Theorem 40).

  • Shape. The listing is the trained net; the bridge theorems are shape-general and the tie is stated at batch 128.

3.5 Recap

Two operations, and the first whole-network theorem with a hypothesis in it. The convolution is a linear map, so its backward needs no hypothesis on the input: for an odd kernel it is a convolution of the cotangent against the kernel with its channel axes swapped and its spatial axes flipped, which is what the emitted transpose and reverse compute. The max-pool is the second convention after ReLU’s, argmax ties broken by the lowest index, and its bridge holds where every window has one maximum. An all-zero window after ReLU fails that, so the excluded set is not null, and Theorem 40 carries it as a stated hypothesis rather than hiding it. Theorems 41 and 42 tie the four convolution updates and every gradient node to the loss’s derivative, ties included, and Theorem 43 is the descent through the convolutions.

The network is two convolutions deep, and that is as deep as this recipe goes: the MNIST head holds nearly every parameter, and the width sweep shows the head barely matters once the features are good. Going deeper, on a harder dataset, is where two things start to bite that nothing so far has needed, a normalization inside the stack and a choice of optimizer. The next chapter adds the one operator whose gradient needs real analysis to bound.