Day 100Module 10

in-technical-review

Capstone 5: train a network with your CUDA kernels

Train a 784-256-10 network on Fashion-MNIST. You will write its forward pass, loss, backward pass, and momentum update in CUDA, then compare one epoch with PyTorch on the same GPU.

Your work passes only when cases 1 through 7 pass. You must also submit an Nsight Systems report and explain one measured limit in the training step.

Retrieve what you need

Day 99 gave you the gradient equations for one linear layer. Day 16 gave you a tiled matrix multiply, and day 48 showed how kernel fusion cuts launch overhead when each kernel does little work.

Before you continue, answer this without running code: which earlier kernel shape computes a weight gradient such as X^T dH? The answer is a matrix multiply with one transposed input.

Predict the profile

The network has 203,530 parameters and trains with batches of 128. The fused path launches 10 kernels per step; the unfused path launches 17.

Write down three predictions:

Which path will have the lower epoch time?
Will CUDA kernel time fill most of the epoch, or will launch work dominate?
Which kernel row will own the most GPU time?

Give each prediction a reason. You will compare it with the checked run only after you have built the model and the test plan.

The minimum model

For a batch X, the first layer computes H = ReLU(XW1 + b1). The second layer computes logits Z = HW2 + b2, and stable softmax cross entropy turns those logits into a loss and dZ.

Use the fixed recipe: batch 128, learning rate 0.1, momentum 0.9, 20 epochs, and seed 20260829. The repository ships a fixed 6,000-image training subset and the full 10,000-image Fashion-MNIST test split.

Fashion-MNIST uses the same 28 by 28 byte format as MNIST and has an explicit MIT license. The checked-in data, hashes, and rebuild command are in SOURCE.md, so the run needs no network.

Work through the backward pass

The calculus is not the task. Map each given equation to a CUDA operation:

dW2 = H^T dZ             db2 = sum_rows(dZ)
dH  = (dZ W2^T) * (H>0)
dW1 = X^T dH             db1 = sum_rows(dH)
dX  = dH W1^T

dW2, dW1, and dX are matrix multiplies with a transposed input where shown. Each db is a column reduction, while (H>0) is an elementwise ReLU mask.

This mapping tells you which earlier kernel to reuse. It also tells you what to test alone before you place the operations in a training loop.

Check your mapping

Why must db1 reduce across the batch dimension? Each image contributes one gradient for every hidden unit, but the model stores one bias per hidden unit.

Why does dH apply the ReLU mask? A hidden value at or below zero produced no forward output, so its gradient must also be zero.

Read the worked reference

The reference program is mlp_train.cu. The learner signatures are in submission.h.

The forward matrix multiply keeps bias and ReLU in its epilogue. Template parameters resolve both choices at compile time, so the code writes no separate bias or ReLU array:

    if (row < m && col < n) {
        if (kBias) {
            acc += bias[col];
        }
        if (kRelu) {
            acc = fmaxf(acc, 0.0f);
        }
        c[row * n + col] = acc;
    }

The four parameter tensors use one flat buffer. One grid-stride kernel can update all weights and velocity values in one launch:

__global__ void sgdMomentum(float* __restrict__ w, const float* __restrict__ g,
                            float* __restrict__ v, float lr, float momentum,
                            int n) {
    const int stride = gridDim.x * blockDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride) {
        const float vel = momentum * v[i] - lr * g[i];
        v[i] = vel;
        w[i] += vel;
    }
}

The loss kernel must subtract each row's maximum before taking exponentials. Case 3 uses logits at plus and minus 60 so an unstable softmax fails before it can spoil a full training run.

Build the reference

Compile for the compute capability of your GPU. This command targets compute capability 7.5:

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo \
    -I"$CONDA_PREFIX/include" -L"$CONDA_PREFIX/lib" \
    -o mlp_train mlp_train.cu -lz
./mlp_train --data ../../data/day100
python3 pytorch_baseline.py --data ../../data/day100
mkdir -p profile
nsys profile --trace cuda,nvtx --capture-range=nvtx \
    --nvtx-capture=profile-epoch --force-overwrite true \
    --output profile/day100-mlp ./mlp_train --profile \
    --data ../../data/day100

The full correctness run trains a float64 CPU reference, so it takes much longer than the GPU epochs under test. Use --profile for a short capture that warms the kernels and records one fused epoch.

Run the three nsys stats exports from the README after the capture. They produce the kernel, NVTX, and CUDA API CSV files used in the analysis.

Practice with less help each time

Stage 1: forward and loss

Implement solve_forward and solve_loss. Start with batch 1, then batch 128, and finish with the extreme-logit loss case.

Pass cases 1 through 3 before writing a gradient kernel. These checks separate matrix indexing errors from softmax stability errors.

Stage 2: backward and update

Implement the equations above through solve_backward, then implement solve_step. Compare every gradient with the float64 result before using it to change a weight.

Case 4 checks 200 seeded coordinates against both the analytic float64 backward pass and central finite differences. Case 5 checks one momentum update with |got-ref| <= 1e-6 + 1e-5*|ref|.

Stage 3: prove that it learns

Overfit 32 images for 200 steps. Case 6 requires a loss below 0.01, which catches paired errors that may agree at one gradient-check point.

Only then run all 20 epochs. Case 7 passes when your held-out accuracy is within 1.0 percentage point of the float64 CPU reference trained in the same run.

Stage 4: measure one epoch

Time fused and unfused epochs after epoch zero. Keep the initial weights, data order, batch size, and recipe fixed so the timing comparison changes only the kernel packaging.

Profile the profile-epoch NVTX range with Nsight Systems. Do not profile the whole program, because the CPU reference would hide the GPU step you need to inspect.

The independent submission gate

Submit your source, a .nsys-rep, the three CSV exports, and RESULTS.md. Your RESULTS.md must name the GPU, driver, toolkit, build command, run command, and training recipe.

The current repository still lacks the external stage harness and shipped reference weights listed in GAPS.toml. Until those gaps close, include the harness you used to call the five submission.h entry points.

Gate	Required result
1, `onebatch-forward`	Batch 1 logits match the float64 forward reference within the printed tolerance
2, `batch-forward`	Batch 128 logits match the same reference within the depth-scaled tolerance
3, `loss-extreme`	Logits at plus and minus 60 produce finite loss and gradients within tolerance
4, `grad-check`	200 seeded parameter coordinates pass the analytic and finite-difference checks
5, `step`	One momentum update passes `1e-6 + 1e-5*
6, `overfit-32`	Loss after 200 steps is below 0.01
7, `train`	Test accuracy is within 1.0 percentage point of the same-run float64 reference
8, `perf`	Fused and unfused epoch times are reported after one warm-up epoch

Cases 1 through 7 are the correctness gate. Exit code zero alone does not prove the network is correct.

For performance, bronze beats the shipped unfused path. Silver takes at most 0.90 times PyTorch's epoch time on compute capability 7.5, while gold takes at most 0.60 times PyTorch and keeps epoch-time CV below 5 percent.

Your analysis must quote one kernel row and one API or NVTX row from your own exports. State the change you would try next, the metric it should move, and the direction you expect it to move.

Checked results, after the prediction

The CUDA 13.0 run on 2026-09-02 passed all eight cases. These numbers describe that run only; use the relative gates above on your GPU.

What the run reports	Value
eight cases	passed
overfit-32 loss after 200 steps	0.000210, gate 0.010
fused epoch, mean of 19 timed epochs	16.571 ms
unfused epoch, mean of 19 timed epochs	9.342 ms
unfused / fused	0.564
PyTorch epoch median, same workload	48.687 ms, CV 2.36%
test accuracy, fused	82.89%
test accuracy, float64 CPU reference	83.18%
test accuracy, PyTorch	81.79%
captured `profile-epoch` NVTX range	21.759297 ms
widest kernel-summary row	`gemmTiled<true,false,false,false>`, 9.401344 ms, 44.7%
summed kernel time / captured range	21.054574 / 21.759297 ms
largest CUDA API row	`cudaDeviceSynchronize`, 18.599802 ms
`cudaLaunchKernel`	2.510467 ms over 470 calls

The fused path lost to the unfused path: 16.571 ms against 9.342 ms. Removing seven launches did not offset the slower fused matrix-multiply path in this implementation.

Kernel rows filled 21.054574 ms of the 21.759297 ms captured range. That result rejects the prediction that launch work would dominate the epoch, even though 470 cudaLaunchKernel calls still took 2.510467 ms of API time.

The largest row combined dW1 and dW2 because both launches used the same template specialization. The summary therefore cannot prove which of those two operations took more time.

The fused CUDA time is below 0.60 times the PyTorch median. It does not earn gold because the CUDA harness did not report epoch-time CV, and it misses bronze because it did not beat the shipped unfused path.

The checked transcripts and profiler files are listed in the page metadata. The failed case 2 and case 5 transcripts record bad early tolerances, not failed final kernels.

Explain a different result

Compare your three predictions with your tables. For each mismatch, name the row that changed your view instead of explaining it from memory.

If kernel time fills the epoch, start with the largest kernel row. If a large part of the epoch has no kernel, inspect cuda_api_sum and the gaps between kernels before changing arithmetic.

A CUDA graph is a testable option when launch work is large. Predict a lower cudaLaunchKernel API total and unchanged kernel durations, then capture a second report to test both claims.

Common failures

If case 4 passes but case 6 fails, check the ReLU mask and the signs in the weight update. A point gradient check can miss two errors that cancel, while the 32-image training path follows many updates.

If accuracy changes after an epilogue fusion, restart both runs from the same weights. A larger change than normal floating-point drift often means a race; day 68 shows how to test that.

Do not time epoch zero. It includes first-use costs for each kernel and gives a poor estimate of steady epoch time.

Keep the PyTorch comparison on the same GPU, architecture, batch size, data order, and 20-epoch recipe. A faster result on a different workload does not answer this capstone's question.

Ungraded extension: transformer forward

After the training capstone passes, you may replace the MLP with the two-layer transformer forward path from days 95 through 97. Check layer one against a float64 reference, then time the full forward pass.

This extension has no effect on the Day 100 gate. It omits the backward pass, so it cannot replace the required Fashion-MNIST training submission.

Sources

Fashion-MNIST source and license, checked 2026-08-29
llm.c by Andrej Karpathy, checked 2026-09-01
CUDA C++ Best Practices Guide, checked 2026-08-29
Nsight Systems User Guide, checked 2026-08-30
Programming Massively Parallel Processors, fourth edition, checked 2026-08-29

Finish

The course ends with the same order used in this capstone: predict, check correctness, measure, and explain the gap. Keep the report beside the code so the next change starts from evidence rather than memory.