CUDA Convolutional Neural Network Kernels

Hand-written CUDA kernels for the two convolutional layers and two fully-connected (classifier) layers of a VGG-style CNN, plus GPU profiling, a CuDNN comparison, and a mechanistic roofline performance model. Graduate coursework for CS 259 — Learning Machines (MQST, Spring 2023).

📊 Results & docs site: https://the-iron-ryan.github.io/cuda-conv-kernels/

---

Overview

Four kernels are implemented from scratch in src/convolution.cu:

Kernel	Type	Dimensions
Conv1	Convolution	Ny=Nx=224, Ni=Nn=64, Ky=Kx=3
Conv2	Convolution	Ny=Nx=14, Ni=Nn=512, Ky=Kx=3
Classifier1	Fully-connected	Ni=25088 → Nn=4096
Classifier2	Fully-connected	Ni=4096 → Nn=1024

Parallelization strategy. Each kernel uses a block-strided launch: threads stride across the problem in blockDim.x * gridDim.x steps so any problem size maps onto a fixed grid. Reuse is squeezed out with register tiling (Tn, Ti, Tx, Ty) sized to fit working sets into the L1/L2 caches, and each output goes through a ReLU activation.

• Conv1 strides over the spatial (y, x) dimensions.
• Conv2 strides over the channel (n, i) dimensions instead — this gives better data reuse when the number of feature maps is large (the Conv2 layer shape), which is why it lands at a much higher arithmetic intensity on the roofline.

The main limits are the number of resident blocks/SMs and the amount of memory reuse per block; the kernels turn out to be memory-bandwidth bound (see Results).

Build & run

Requires the CUDA Toolkit (nvcc). Built and profiled on an NVIDIA Titan V (Volta, sm_70); the Makefile defaults to SMS=70.

cd src
make                       # builds ./convolution
./convolution -k Conv1 -b 16   # kernel: Conv1|Conv2|Classifier1|Classifier2, batches: 16

Profiling helpers (require nvprof):

make profilePerformance    # nvprof timing
make profileMemory         # nvprof -m dram_read_throughput

Results

Measured at batch size 1 on the Titan V. Our naive implementation vs. CuDNN:

Kernel	Ours (ms)	Ours BW (GB/s)	CuDNN (ms)	CuDNN BW (GB/s)
Conv1	121.34	52.34	0.342	406.34
Conv2	43.677	0.097	0.186	N/A (Winograd)
Class1	134.51	3.43	0.746	542.32
Class2	17.268	0.77	0.040	416.23

Takeaways

• The kernels sit well below both the DRAM and L2 roofline ceilings → memory-bound, with too little data reuse (no shared-memory scratchpad was used).
• We trail CuDNN by ~100×. CuDNN uses the Winograd algorithm backed by tensor-core matrix multiply; our naive direct convolution can't compete on either raw compute or memory throughput.
• The large block size left only 49 of 80 SMs worth of blocks occupied, costing further performance.

Performance model (Mini-Project 2)

A mechanistic model (perf/roofline.py, from Performance-Model.xlsx) predicts operational intensity as a function of the tiling parameters at each level of the memory hierarchy. Choosing tiles so the working set fits the L2 (4.5 MB) and L1 (65 KB/SM) caches yields target intensities of 14.05 FLOP/byte for Conv1 (L1 set) and 34.37 FLOP/byte for Conv2 (L2 set). Run python perf/roofline.py to reproduce these numbers and regenerate the roofline chart.

Repository layout

src/          convolution.cu (the kernels), vectorAdd_original.cu (NVIDIA starter), Makefile
perf/         roofline.py, nvprof output (profileOutput/profileResult), spreadsheets
reports/      Mini-Project 1 & 2 write-ups (PDF), roofline graph, LaTeX source
docs/         GitHub Pages site (index.html + assets/roofline.png)

Reports

• Mini-Project 1 write-up — strategy, timings, roofline, and the CuDNN comparison (Q4).
• Mini-Project 2 report — the mechanistic performance model.

Credits

CUDA kernels and Mini-Project 1 by Ryan Dougherty.

The performance analysis and mechanistic performance model (Mini-Project 2) were a team effort. Thanks to my teammates for their contributions to the analysis:

• Arvind Gangadhar
• Payal Kaushik
• Will Yzaguirre

The build system derives from NVIDIA's CUDA samples vectorAdd template.

License

MIT © 2023 Ryan Dougherty.