CUDA Convolutional Neural Network Kernels

Overview

Four kernels implemented from scratch in a single CUDA translation unit.

Each kernel launches with a block-strided pattern — threads stride across the problem in blockDim.x * gridDim.x steps so any problem size maps onto a fixed grid. Data reuse is squeezed out with register tiling (Tn, Ti, Tx, Ty) sized to keep working sets inside the L1/L2 caches, and every output passes through a ReLU. The two convolution kernels differ in which axis they stride over — the key idea behind their very different roofline placement.

The kernels

Two convolutional layers and two fully-connected classifiers.

Conv1 3×3, 224², 64→64

Strides over the spatial (y, x) dimensions. Large feature maps, modest channel count. Non-contiguous access hurts reuse → low arithmetic intensity, memory-bound.

Conv2 3×3, 14², 512→512

Strides over the channel (n, i) dimensions instead. With many feature maps this yields far better reuse → much higher arithmetic intensity on the roofline.

Classifier1 25088 → 4096

Fully-connected layer as a tiled matrix–vector product with ReLU. Threads strided over the output dimension, input tiled by Ti.

Classifier2 4096 → 1024

The smaller fully-connected layer, same structure — the cheapest of the four kernels to evaluate.

Results

Batch size 1 on the NVIDIA Titan V — our naive kernels vs. CuDNN.

Execution time and DRAM read throughput. CuDNN's Conv2 uses Winograd, for which a direct throughput comparison isn't meaningful.
Kernel	Ours (ms)	Ours BW (GB/s)	CuDNN (ms)	CuDNN BW (GB/s)	Slowdown
Conv1	121.34	52.34	0.342	406.34	~355×
Conv2	43.677	0.097	0.186	N/A	~235×
Classifier1	134.51	3.43	0.746	542.32	~180×
Classifier2	17.268	0.77	0.040	416.23	~430×

What the numbers say

All four kernels are memory-bandwidth bound — they sit well below both roofline ceilings, with too little data reuse (no shared-memory scratchpad).
We trail CuDNN by roughly 100×: it uses the Winograd algorithm on tensor-core matrix multiply, which a naive direct convolution can't match.
The large block size left only 49 of 80 SMs worth of blocks occupied, leaving performance on the table.

Roofline model

Where each kernel lands relative to the Titan V's DRAM and L2 ceilings.

Log-log roofline plot for the Titan V. DRAM and L2 ceilings rise to a 14.9 TFLOP/s FP32 peak; the four kernel operating points sit far below the ceilings, indicating a memory-bound implementation. — FP32 peak ≈ 14.9 TFLOP/s · DRAM BW 651.3 GB/s · L2 BW 2155 GB/s. Regenerate with `python perf/roofline.py`.

Mechanistic performance model

A separate model (Mini-Project 2) predicts operational intensity from the tiling parameters at each level of the memory hierarchy. Sizing tiles so the working set fits the L2 (4.5 MB) and L1 (65 KB/SM) caches gives target intensities of 14.05 FLOP/byte for Conv1 (L1 set) and 34.37 FLOP/byte for Conv2 (L2 set).

Build & run

Needs the CUDA Toolkit (nvcc). Makefile defaults to SMS=70.

cd src
make                             # builds ./convolution
./convolution -k Conv1 -b 16     # kernel + batch count

# kernels: Conv1 | Conv2 | Classifier1 | Classifier2
make profilePerformance          # nvprof timing
make profileMemory               # nvprof -m dram_read_throughput

Reports & artifacts

Full write-ups and raw profiling data.

📄Mini-Project 1 write-up Strategy, timings, roofline placement, optimizations. 📊CuDNN comparison (Q4) Side-by-side timing and throughput vs. CuDNN. 📐Mini-Project 2 report Mechanistic performance model of the convolution kernels. 🗂️Profiling data & spreadsheets nvprof output plus the performance-model workbooks.