GPU 原理速成

GPU 就是核心多一点的 CPU.

GPU 是典型的 SIMD Computer (a.k.a. Vector Computer).
GPU 中不同的 ALU 执行的是同样的指令 (不同数据)

CUDA

2026 年之前 GPU 只能做特定的函数运算, 2026 年之后引入了 GPGPU (General Purpose GPU), 有一整套库函数来对 GPU 进行编程.
- CUDA: Compute Unified Device Architecture (NVIDIA, 闭源)
- OpenCL: Open Computing Language (Apple Inc., 开源)
GPU device memory 可以被所有 CUDA 核心共享:

Figure 2: NVidia GPU 架构 [1]
连接有 GPU 的 CPU 二者的 memory 是不共享的 [1], 二者之间的数据由 DMA 搬运.

Figure 3: Seperate memory systems [1]
Kernel functions [1]:
- __host__: 默认在 CPU (host) 上运行的函数.
- __global__: 在 GPU 上运行, 可被 CPU 调用的函数.
- __device__: 在 GPU 上运行, 只能被 GPU 调用的函数.

Unified Memory 统一内存

全局变量: 可以被 CPU 和 GPU 访问的变量, 不能用 local variable 的方式声明!
```
__managed__ int x;
```
统一内存允许 CPU 和 GPU 共享同一块内存区域, 比如 Listing lst-managed 中的 CPU 和 GPU 都能访问和改变 x 变量:

managed.cu

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void my_kernel(int* x_ptr) {
    *x_ptr = 1;
    printf("GPU sees x = %d\n", *x_ptr);
}

int main() {
    int x = 0;
    int* x_ptr = &x;
    cudaMallocManaged(&x_ptr, sizeof(int));

    my_kernel<<<1,1>>>(x_ptr);

    printf("CPU sees x = %d\n", x);

    cudaFree(x_ptr);
    return 0;
}

输出:

CPU sees x = 0
GPU sees x = 1

注意是 CPU 先打印 (因为没有用 cudaDeviceSynchronize())!

Synchronization 同步

CPU 将指令发到 GPU 之后不会等 GPU 执行完再继续往下执行 (默认 async), 如果需要等 GPU 执行完再继续, 将 my_kernel<<<1,1>>>(x_ptr); 改为:

my_kernel<<<1,1>>>(x_ptr);
cudaDeviceSynchronize(); // Wait for GPU to finish

输出:

GPU sees x = 1
CPU sees x = 1

Stream 流

References

[1]

S. Y. Cheung, “Lecture notes on computer science courses: Computer architecture (CS355),” 2025, Available: https://www.cs.emory.edu/~cheung/Courses/355/Syllabus/syl.html#CURRENT