给25级xpy 伟大思想的pre，一天时间搓出来的，过程艰辛，但并不是学习的艰辛，具体遭到了什么问题就不说了，但已经解决了。自我感觉做的不错，遂放到blog上🤓 也是本人第一次用vscode-marp 做ppt，终于摆脱word的阴间排版了。

---

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 (中央处理器)
  • 拥有少量强大的核心 (Core)。
  • 具有复杂的控制单元 (Control) 和多级缓存 (L1, L2, L3 Cache)。
  • 擅长串行任务和复杂的控制逻辑。

---

What is a GPU?

• GPU (图形处理器)
  • 拥有大量简单的核心 (Core)。
  • 适合处理并行任务和大规模数据。
  • 常用于图形渲染和机器学习等领域。

---

---

What is RAM?

• 随机存取存储器 (RAM) 是计算机的主要内存，用于存储正在使用的数据和程序。
  • 动态随机存取存储器 (DRAM)：容量大，速度较慢，常用于主内存。
  • 静态随机存取存储器 (SRAM)：速度快，容量小，常用于缓存(Cache)。
• CPU 可以直接、快速地访问内存中任何位置的数据，无需按顺序查找。内存只能在电脑通电时存储数据。一旦你关机、重启或断电，内存中的所有数据都会立刻消失。

---

---

What is Cache?

• 缓存 (Cache) 是一种高速存储器，位于 CPU/GPU 和主内存(DRAM)之间。
• CPU 的运行速度极快，但访问主内存（DRAM，即“内存条”）的速度相对慢得多。这就形成了一个巨大的“速度鸿沟”。缓存就是为了弥补这个速度差而生的。它只存储 CPU/GPU 最可能立即需要的数据和指令。
• 多级缓存结构是现代 CPU 高性能的基础。

---

缓存级别	所属设备	速度/容量	数据共享范围
L1 缓存（SRAM）	CPU/GPU 核心内部	最快 / 容量最小	核心独享（每个核心都有自己的 L1）
L2 缓存	CPU/GPU 核心附近	较快 / 适中	通常核心独享或小组共享
L3 缓存	CPU	较慢 / 容量最大	所有核心共享
共享内存	GPU	极快（用户可控）	线程块内部协作共享

---

特性	CPU 缓存 (L1, L2, L3)	GPU 缓存 (L1, L2, 共享内存)
设计核心目标	低延迟 (Low Latency)：尽快完成单个复杂任务。	高吞吐量 (High Throughput)：同时处理大量简单任务。
层级数	L1、L2、L3 三级。	通常只有 L1、L2 两级。
L3 缓存	有。容量大，供所有核心共享。	基本无。架构上更依赖 L2 和共享内存。
最特殊结构	L3 缓存（复杂数据共享）。	共享内存 (Shared Memory)（程序员可控的高速存储）。
数据共享方式	核心间通过 L3 缓存或总线进行数据同步。	线程块内部通过共享内存直接快速协作。

---

GPU Architecture

现代 GPU（以 NVIDIA 的 CUDA 架构为例，如 Volta, Turing, Ampere, Hopper）从上到下可以分成三个主要层级：

1. 芯片级别：GPC (Graphics Processing Clusters) 整个 GPU 芯片由多个被称为 GPC (图形处理集群) 的大型单元组成。
2. 集群级别：SM (Streaming Multiprocessors) 每个 GPC 内部包含多个 SM (流式多处理器)。
   SM 是 GPU 的真正核心！ 它是执行并行计算的基本控制单元。

---

GPU Architecture

3. SM 内部：核心、Warp 和内存

• CUDAs 核心 (CUDA Cores)：大量（每个 SM 包含数十到数百个）。负责实际的数学运算（如加法、乘法）。
• Warp Scheduler： SM 的真正指挥官，负责将接收到的任务（线程块）拆解成 Warp（线程束，32 个线程一组），并向所有 32 个线程同时广播指令（SIMT 模型）
• 共享内存/L1 缓存：供 SM 内的线程块协作共享数据。
• 寄存器堆：每个线程独享的高速存储单元。

---

GPU Programming mode：SIMT

Single instruction multiple threads (SIMT) 单指令多线程

• 所有线程执行相同的代码，但它们可以根据程序中的分支或条件语句采取不同的执行路径（例如 if/else 结构）。
• Thread (线程)：执行计算的基本单位。
• Thread Block (线程块)：线程被分组到块中。同一块内的线程具有共享内存。
• Grid (网格)：线程块又被分组到网格中。

---

GPU Programming mode: SIMT

---

SIMT on GPU Hardware

• 一个 Grid 中的 Thread Blocks（线程块）会被分配到 GPU 的各个 SM 上执行。
• 当一个线程块被分配到 SM 上时，SM 内的调度器会首先将这个线程块中的所有线程（通常是 128、 256 或 1024 个）分成多个 Warp（每个 Warp 32 个线程）。
• 调度器的工作就是从就绪的 Warp 中选择一个，然后向该 Warp 中的所有 32 个线程同时发出相同的指令。
• GPU 不会调度单个线程，也不会调度整个线程块，它只以 Warp 为单位进行操作。

---

---

CUDA cores

当一个 Warp 被调度时，它会同时被送往 SM 内部的 CUDA 核心。在理想情况下（没有分支冲突），Warp 中的 32 个线程会同时在 32 个不同的 CUDA 核心上执行同一条指令。

---

We keep talking about CUDA,

then you guys may ask:

---

What is CUDA?

• CUDA (Compute Unified Device Architecture) 是 NVIDIA 提供的并行计算平台和编程模型。
• 它允许开发者使用 C/C++ 等高级语言编写程序，以充分利用 GPU 的强大计算能力。
• CUDA 提供了一套丰富的 API 和库，使得在 GPU 上进行高性能计算变得更加简单和高效。

---

CUDA Programming Example: Vector Addition

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

__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)

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) 是线性代数中最基本的操作之一，广泛应用于科学计算、图形处理和机器学习等领域。
• Compute C = dot(A.T, B)

---

Thread-level: register tiling

大规模线程并行执行（多个线程），每个线程又极高效率地利用了寄存器（平铺）。

类似于缓存+分块的思想

---

__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

整个 Thread Block (线程块) 作为一个协作团队工作

---

Block-level: shared memory tiling

• 平铺 (Tiling): 矩阵 A^T 和 B 被分解成大小为 L×S 或 S×L 的大块。
• 协作 (Cooperation): 线程块中的所有线程首先协同工作，将 A^T 和 B 的当前大块从慢速的全局内存（Global Memory）一次性搬运到快速的共享内存（Shared Memory）中。
• 重用 (Reuse): 一旦数据进入共享内存，块内的所有线程都可以多次、快速地访问这些数据，极大地提高了计算效率。

---

__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.