This was a presentation for the great ideas of 2025 xpy, made in one day and archived on the blog as a memory. It was also my first time using vscode-marp to make slides, finally escaping Office’s painful layout system.

marp: true

theme: gaia

footer: ‘JaneZ 2025-10-14’

paginate: true

html: true

style: |

section a {

font-size: 30px;

}

GPU Acceleration

Yihan Zhu @JaneZ ACM Class 2024

2025.10.14

Overview

  • Basic Architecture of CPU, GPU, RAM, Cache
  • GPU Architecture
  • GPU Programming
  • Case study: Matrix Multiplication on GPU

What is a GPU?

  • CPU (Central Processing Unit)
    • Has a small number of powerful cores.
    • Has a complex control unit and multi-level cache (L1, L2, L3 Cache).
    • Good at serial tasks and complex control logic.

What is a GPU?

  • GPU (Graphics Processing Unit)
    • Has a large number of simple cores.
    • Suitable for parallel tasks and large-scale data processing.
    • Commonly used in graphics rendering, machine learning, and related fields.

What is RAM?

  • Random Access Memory (RAM) is the main memory of a computer, used to store data and programs currently in use.
    • Dynamic Random Access Memory (DRAM): large capacity, slower speed, commonly used as main memory.
    • Static Random Access Memory (SRAM): fast, small capacity, commonly used as cache.
  • The CPU can directly and quickly access data at any memory location without sequential search. Memory can store data only while the computer is powered on. Once you shut down, restart, or lose power, all data in memory disappears immediately.

What is Cache?

  • Cache is a high-speed memory located between the CPU/GPU and main memory (DRAM).
  • The CPU runs extremely fast, but accessing main memory, namely DRAM modules, is much slower. This creates a huge speed gap. Cache exists to bridge this gap. It stores only the data and instructions that the CPU/GPU is most likely to need immediately.
  • Multi-level cache hierarchy is the foundation of modern CPU performance.
Cache level Device Speed / Capacity Data sharing scope
L1 Cache (SRAM) Inside CPU / GPU core Fastest / smallest capacity Private to each core
L2 Cache Near CPU / GPU core Faster / medium capacity Usually private to a core or shared by a small group
L3 Cache CPU Slower / largest capacity Shared by all cores
Shared Memory GPU Extremely fast and user-controllable Shared cooperatively within a thread block
Feature CPU Cache (L1 / L2 / L3) GPU Cache (L1 / L2 / Shared Memory)
Core design goal Low latency: finish one complex task as quickly as possible High throughput: process many simple tasks simultaneously
Number of levels L1, L2, L3 Usually only L1 and L2
L3 Cache Exists. Large capacity, shared by all cores Generally absent. The architecture relies more on L2 and shared memory
Most special structure L3 cache for complex data sharing Shared memory, a programmer-controllable high-speed storage
Data sharing method Cores synchronize through L3 cache or buses Threads within a block cooperate directly through shared memory

GPU Architecture

Modern GPUs, taking NVIDIA CUDA architectures such as Volta, Turing, Ampere, and Hopper as examples, can be divided into three major levels from top to bottom:

  1. Chip level: GPC (Graphics Processing Clusters). The whole GPU chip consists of multiple large units called GPCs.
  2. Cluster level: SM (Streaming Multiprocessors). Each GPC contains multiple SMs.

The SM is the real core of the GPU. It is the basic control unit for parallel computation.

GPU Architecture

  1. Inside an SM: cores, warps, and memory
  • CUDA Cores: many cores, with each SM containing dozens to hundreds. They perform actual mathematical operations such as addition and multiplication.
  • Warp Scheduler: the real commander of the SM. It breaks received tasks, namely thread blocks, into warps, each containing 32 threads, and broadcasts instructions to all 32 threads at once under the SIMT model.
  • Shared Memory / L1 Cache: shared cooperatively by threads within the SM.
  • Register File: high-speed storage private to each thread.

GPU Programming Mode: SIMT

Single Instruction Multiple Threads (SIMT)

  • All threads execute the same code, but they can take different execution paths according to branch or condition statements, such as if/else.
  • Thread: the basic unit of computation.
  • Thread Block: threads are grouped into blocks. Threads in the same block have shared memory.
  • Grid: thread blocks are grouped into a grid.

SIMT on GPU Hardware

  • Thread blocks in a grid are assigned to different SMs on the GPU for execution.
  • When a thread block is assigned to an SM, the scheduler inside the SM first divides all threads in the block, usually 128, 256, or 1024 threads, into multiple warps, each with 32 threads.
  • The scheduler’s job is to select a ready warp and issue the same instruction simultaneously to all 32 threads in that warp.
  • The GPU does not schedule individual threads, nor does it schedule the entire thread block. It operates only at the warp level.

CUDA Cores

When a warp is scheduled, it is sent to the CUDA cores inside the SM. Ideally, with no branch divergence, the 32 threads in the warp execute the same instruction simultaneously on 32 different CUDA cores.

We Keep Talking About CUDA. What Is CUDA?

  • CUDA (Compute Unified Device Architecture) is a parallel-computing platform and programming model provided by NVIDIA.
  • It allows developers to write programs in high-level languages such as C/C++ to fully use GPU computing power.
  • CUDA provides a rich set of APIs and libraries, making high-performance GPU computing simpler and more efficient.

CUDA Programming Example: Vector Addition

1
2
3
4
5
void VecAddCPU(float* A, float *B, float* C, int n) {
 for (int i = 0; i < n; ++i) {
   C[i] = A[i] + B[i];
 }
}

CUDA Programming Example: Vector Addition

1
2
3
4
5
6
__global__ void VecAddGPU(float* A, float *B, float* C, int n) {
 int i = blockIdx.x * blockDim.x + threadIdx.x;
 if (i < n) {
   C[i] = A[i] + B[i];
 }
}

Vector Addition (Host Code)

1
2
3
4
5
6
7
8
9
10
11
12
13
void VecAddCUDA(float* Acpu, float *Bcpu, float* Ccpu, int n) {
 float *dA, *dB, *dC;
 cudaMalloc(&dA, n * sizeof(float));
 cudaMalloc(&dB, n * sizeof(float));
 cudaMalloc(&dC, n * sizeof(float));
 cudaMemcpy(dA, Acpu, n * sizeof(float), cudaMemcpyHostToDevice);
 cudaMemcpy(dB, Bcpu, n * sizeof(float), cudaMemcpyHostToDevice);
 int threads_per_block = 512;
 int nblocks = (n + threads_per_block - 1) / threads_per_block;
 VecAddKernel<<<nblocks, thread_per_block>>>(dA, dB, dC, n);
 cudaMemcpy(Ccpu, dC, n * sizeof(float), cudaMemcpyDeviceToHost);
 cudaFree(dA); cudaFree(dB); cudaFree(dC);
}

Case Study: Matrix Multiplication on GPU

  • Matrix Multiplication (GEMM) is one of the most fundamental operations in linear algebra and is widely used in scientific computing, graphics processing, machine learning, and many other fields.
  • Compute C = dot(A.T, B)

Thread-Level: Register Tiling

Massive threads execute in parallel, and each thread uses registers very efficiently through tiling.

This is similar to the idea of cache plus blocking.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
__global__ void mm(float A[N][N], float B[N][N], float C[N][N]) {
 int ybase = blockIdx.y * blockDim.y + threadIdx.y;
 int xbase = blockIdx.x * blockDim.x + threadIdx.x;
 float c[V][V] = {0};
 float a[V], b[V];
 for (int k = 0; k < N; ++k) {
 a[:] = A[k, ybase*V : ybase*V + V];
 b[:] = B[k, xbase*V : xbase*V + V];
 for (int y = 0; y < V; ++y) {
 for (int x = 0; x < V; ++x) {
 c[y][x] += a[y] * b[x];
 }
 }
 }
 C[ybase * V : ybase*V + V, xbase*V : xbase*V + V] = c[:];
}

Block-Level: Shared Memory Tiling

The entire Thread Block works as a cooperative team.

Block-Level: Shared Memory Tiling

  • Tiling: matrix A^T and B are decomposed into large tiles of size LxS or SxL.
  • Cooperation: all threads in the block first cooperate to move the current tiles of A^T and B from slow global memory into fast shared memory.
  • Reuse: once data enters shared memory, all threads in the block can access it many times and very quickly, greatly improving computation efficiency.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
__global__ void mm(float A[N][N], float B[N][N], float C[N][N]) {
 __shared__ float sA[S][L], sB[S][L];
 float c[V][V] = {0};
 float a[V], b[V];
 int yblock = blockIdx.y;
 int xblock = blockIdx.x;
 for (int ko = 0; ko < N; ko += S) {
 __syncthreads();
 // needs to be implemented by thread cooperative fetching
 sA[:, :] = A[k : k + S, yblock * L : yblock * L + L];
 sB[:, :] = B[k : k + S, xblock * L : xblock * L + L];
 __syncthreads();
 for (int ki = 0; ki < S; ++ ki) {
 a[:] = sA[ki, threadIdx.y * V : threadIdx.y * V + V];
 b[:] = sB[ki, threadIdx.x * V : threadIdx.x * V + V];
 for (int y = 0; y < V; ++y) {
 for (int x = 0; x < V; ++x) {
 c[y][x] += a[y] * b[x];
 }
 }
 }
 }
 int ybase = blockIdx.y * blockDim.y + threadIdx.y;
 int xbase = blockIdx.x * blockDim.x + threadIdx.x;
 C[ybase * V : ybase*V + V, xbase*V : xbase*V + V] = c[:];
}

Thanks for Listening!

Acknowledgements

CMU 10-414/714: Deep Learning Systems by Prof. Tianqi Chen and Zico Kolter.