⚡ GPU-Resident Shallow Neural Network Training

A fully GPU-centric CUDA training pipeline that eliminates host↔device round-trips — up to 47× faster than a naive per-operation CUDA baseline.

ENSI (École nationale Supérieure d'Informatique) — High Performance Computing, 2CS. This repository packages the code, benchmarks, and report behind "Performance Evaluation of a CUDA-Based Alternative Strategy for Shallow Neural Network Training."

---

TL;DR

We took a shallow neural network (32 → 256 → 1, ReLU, MSE, SGD) whose reference CUDA implementation only offloaded matrix multiplication to the GPU — leaving activation functions, weight updates, and every batch transfer on the CPU — and rebuilt the training loop as a single GPU-resident pipeline: the dataset, weights, gradients, and every intermediate activation live in device memory for the entire run, and ten custom kernels replace every host-side operation.

Dataset	Reference (naive CUDA)	GPU-resident (ours)	Speedup
Small (256 samples)	0.484 s	0.029 s	16.6×
Medium (2,560 samples)	3.105 s	0.195 s	15.9×
Large (25,600 samples)	31.359 s	0.661 s	47.4×

Best configuration per dataset, batch size 256, measured with omp_get_wtime() on an NVIDIA Tesla T4.

---

Table of Contents

• Why this exists
• Architecture
• Repository structure
• Quickstart
• Implementations at a glance
• Results
• Key findings
• Limitations & future work
• Reports & citation

---

Why this exists

Most "GPU-accelerated" neural network demos parallelize the obvious bottleneck — matrix multiplication — and stop there. That's exactly what our reference implementation does: one CUDA kernel per matmul, launched from a training loop that still runs on the CPU. It works, but every batch pays for:

• fresh cudaMalloc / cudaFree calls on every single operation,
• redundant cudaMemcpy of the (constant!) weight matrices on every matmul,
• CPU-side ReLU, subtraction, and weight-update loops that stall the pipeline waiting on the GPU,
• and non-coalesced memory access on the second matrix operand, with no shared-memory reuse.

None of this shows up in a single kernel's profile — it shows up as a training loop that plateaus around 31 s on the large dataset no matter how many threads per block you throw at it (Table below). The bottleneck isn't compute, it's orchestration.

So we asked: what if the entire training loop — forward pass, loss, backward pass, weight updates, even batch extraction — never left the GPU?

Architecture

┌─────────────────────────────────────────────────────────────────────┐
│                      HOST (one-time setup)                          │
│   load_csv() → single H2D copy of X, Y, W1, W2                      │
└───────────────────────────────┬───────────────────────────────────-─┘
                                 │  (all training-run matrices are
                                 │   pre-allocated on the GPU ONCE)
┌────────────────────────────────▼──────────────────────────────────┐
│                         DEVICE (GPU-resident)                      │
│                                                                     │
│   for epoch in 0..EPOCHS:                                          │
│     for batch in dataset:                                          │
│       ┌── copy_batch_kernel ───────────────┐  1D grid, GPU→GPU     │
│       │      X_batch, Y_batch extraction   │  batch extraction     │
│       └─────────────────────────────────────┘                     │
│       ┌── Forward ─────────────────────────┐                       │
│       │  mat_mult_kernel  (X·W1 → Z1)       │  2D grid, 16×16      │
│       │  relu_kernel      (Z1 → H)          │  1D grid, 256        │
│       │  mat_mult_kernel  (H·W2 → Y_pred)   │  2D grid, 16×16      │
│       └─────────────────────────────────────┘                     │
│       ┌── Loss ────────────────────────────┐                       │
│       │  mse_kernel (tree reduction,        │  shared memory       │
│       │              shared mem)            │  block reduction     │
│       └─────────────────────────────────────┘                     │
│       ┌── Backward ────────────────────────┐                       │
│       │  mat_sub_kernel, mat_scalar_mult    │  1D grid             │
│       │  transpose_kernel, mat_mult_kernel  │  2D + 1D grids       │
│       │  relu_derivative, elementwise_mult  │  1D grid             │
│       │  update_weights_kernel              │  1D grid             │
│       └─────────────────────────────────────┘                     │
│                                                                     │
│   ⟶ zero CPU involvement inside the loop, zero repeated cudaMalloc │
└──────────────────────────────────────────────────────────────────-┘

Two thread-mapping strategies, chosen per operation type:

• Matrix multiplication → 2D grid, one thread per output element: row = blockIdx.y*blockDim.y + threadIdx.y, col = blockIdx.x*blockDim.x + threadIdx.x
• Element-wise ops (subtraction, ReLU, scalar mult, weight update, transpose, batch copy) → 1D flat grid: idx = blockIdx.x*blockDim.x + threadIdx.x

Repository structure

.
├── src/
│   ├── sequential/
│   │   └── nn_sequential.c          # CPU baseline (OpenMP-timed)
│   ├── reference/
│   │   └── nn_reference.cu          # naive CUDA: GPU matmul only, CPU orchestration
│   └── alternative/
│       └── nn_gpu_resident.cu       # ours: fully GPU-resident pipeline, 10 kernels
├── scripts/
│   ├── generate_dataset.py          # synthetic dataset generator (convex / regression)
│   ├── build.sh                     # compiles all three implementations
│   └── run_benchmarks.sh            # sweeps threads-per-block × batch size, logs CSV
├── benchmarks/
│   └── results/
│       ├── reference_vs_alternative.csv
│       ├── alternative_sweep_no_tiling.csv
│       └── alternative_sweep_with_tiling.csv
├── docs/
│   ├── figures/                     # generated comparison charts
│   └── report/                      # full written reports (reference + alternative)
├── notebooks/
│   └── benchmarking.ipynb           # original Colab notebook (T4, all variants)
└── data/                            # (generated locally, not versioned — see below)

Quickstart

Requires the CUDA toolkit (nvcc) and a CUDA-capable GPU. Tested on Compute Capability 7.5 (Tesla T4, GTX 1650).

# 1. Generate the three benchmark datasets
python scripts/generate_dataset.py --mode convex --samples 256   --out data/synthetic_convex_small.csv
python scripts/generate_dataset.py --mode convex --samples 2560  --out data/synthetic_convex_medium.csv
python scripts/generate_dataset.py --mode convex --samples 25600 --out data/synthetic_convex_large.csv

# 2. Build everything
chmod +x scripts/*.sh
./scripts/build.sh sm_75          # pass your GPU's arch, e.g. sm_86 for Ampere

# 3. Run
./bin/nn_sequential   data/synthetic_convex_large.csv
./bin/nn_reference    data/synthetic_convex_large.csv
./bin/nn_gpu_resident data/synthetic_convex_large.csv

# 4. (optional) Reproduce the full threads×batch sweep
./scripts/run_benchmarks.sh data/synthetic_convex_large.csv large sm_75

Both CUDA binaries print Epoch N, MSE: ... progress lines and a final Training time: X.XXXX seconds.

Implementations at a glance

	nn_sequential.c	nn_reference.cu	nn_gpu_resident.cu
Where matmul runs	CPU (triple loop)	GPU (1 kernel/call)	GPU (1 kernel/call)
Where ReLU / subtraction / weight update run	CPU	CPU	GPU
Batch extraction	CPU	CPU	GPU (copy_batch_kernel)
Memory allocation pattern	per-batch malloc	cudaMalloc/cudaFree on every op	pre-allocated once, reused all epochs
Weight transfer per batch	n/a	re-uploaded every matmul	uploaded once, stays resident
Loss reduction	serial CPU loop	serial CPU loop	GPU tree reduction in shared memory
Thread mapping	—	2D (matmul only)	2D (matmul) + 1D (element-wise)
Optional tiling	—	—	k-dimension tiling, TILE_K compile flag

Results

1. Reference vs. GPU-resident, across dataset sizes and thread-block configs

The reference implementation's execution time is essentially flat regardless of threads-per-block (31–33 s on the large dataset for every configuration from 4×4 to 32×32) — a clear signature of a memory/orchestration bound problem, not a compute-bound one. The GPU-resident version, by contrast, actually benefits from larger thread blocks once the workload is big enough to use them.

2. Batch size × threads-per-block sweep (GPU-resident, no tiling)

Dataset	Best config	Time
Small	batch 256, 16×16	0.0291 s
Medium	batch 1024, 16×16	0.1318 s
Large	batch 1024, 32×32	0.2476 s

Small workloads barely move with more threads or a bigger batch — the GPU is already underutilized and kernel-launch overhead dominates. Medium and large workloads scale cleanly: 8× improvement between the worst and best large-dataset configuration.

3. Does k-dimension tiling help?

We also tried tiling the inner product loop of mat_mult_kernel (TILE_K-sized chunks instead of a single pass over all 32 input features). Short answer: no, not here.

	Best (no tiling)	Best (with tiling, TILE_K=32)
Large, batch 1024, 32×32	0.2476 s	0.3078 s

Out of 27 tested configurations, only 4 improved (1.6–5.3%) and 23 got worse (up to 24.6% slower). Since INPUT_SIZE = TILE_K = 32, the tiling loop degenerates to a single iteration that adds bookkeeping without improving cache locality — the whole 32-float row/column pair (256 B) already fits in L1. Tiling would start paying off with larger input dimensions where the per-thread working set actually exceeds cache capacity.

Key findings

1. Orchestration overhead, not raw compute, was the reference implementation's bottleneck. Flat execution time across thread-block sizes is the tell — you can't out-parallelize a bottleneck that lives in cudaMalloc/cudaMemcpy calls between kernel launches.
2. Keeping data GPU-resident for the whole training run is the single highest-leverage change. One H2D transfer at the start, zero mid-training transfers, and the speedup scales with dataset size (16× → 16× → 47×) rather than degrading.
3. Bigger batches help more as the dataset grows — small workloads are dominated by fixed kernel-launch overhead, so more parallelism has nothing to amortize it against.
4. Optimizations aren't free lunches. Tiling is a textbook cache optimization, but it needs a working set that doesn't already fit in L1 — applying it blindly can regress performance, as our numbers show.

Limitations & future work

• GPU memory ceiling. The whole dataset + every intermediate matrix must fit in device memory; very large datasets would need a chunked/streaming variant (CUDA streams overlapping transfer and compute).
• Small-workload underutilization. Sub-millisecond kernels are dominated by launch overhead; kernel fusion would help here.
• Shallow-only. The kernels are hand-written for one hidden layer; deeper architectures, other activations (sigmoid/tanh/softmax), and other losses are natural extensions.
• Single GPU. Multi-GPU data/model parallelism is unexplored.

See the full report for the detailed discussion.

Reports & citation

• 📄 docs/report/HPC_Project_Report.pdf — full write-up of the GPU-resident alternative strategy (this repo).
• 📄 docs/report/Reference_Strategy_Report.pdf — original reference-strategy report this work builds on and compares against.
• 📓 notebooks/benchmarking.ipynb — original Google Colab notebook (Tesla T4) used to produce the benchmark numbers.

Authors: Bouderbala Amira, Saidi Selma, Djoghlal Romaisa, Chouider Ikram, Allouche Reda Supervisor: Haichour Amina Selma Institution: École nationale Supérieure d'Informatique (ESI) — 2CS, 2025/2026

---

_{Hardware: NVIDIA Tesla T4 (Turing, CC 7.5, 40 SMs, 2560 CUDA cores) on Google Colab · CUDA 11.x · timings via omp_get_wtime().}