为什么不把 Shader 实现成线程呢？

__device__ int *map;

// row = 1, 2, 3, 4, ..., 1080
// col = 1, 2, 3, 4, ..., 1920
void kernel(int row, int col) {
    t = ...;  // 任意 C/C++ 代码
    map[row * 1920 + col] = t;
}

• 在 GPU 上启动 2,073,600 个线程
  • 并行执行，允许访问 shared memory
  • 恭喜！你发明了 CUDA (的 programming model)

唯一的麻烦

Shared-memory CPU 很难实现

• GeForce 8800 (首个 CUDA GPU, G80 2006, 65nm)
  • 128 cores, v.s. Intel 桌面 CPU: 2C4T
  • The first GPU to support C
  • The first GPU to utilize a scalar thread processor, eliminating the need for programmers to manually manage vector registers
    • Fermi whitepaper
• RTX 5080: 10,752 CUDA Cores
  • 这是怎么做到的？
  • 它们也 (可以) 是 shared memory 啊！

同样的面积，更多的线程

Single instruction, multiple threads

• 一个 PC (program counter) 管多个线程
  • 假设 $(3, 0)$, $(3, 1)$, ... $(3, 31)$ 都被分配到了同一个 core 上
    • map[3 * 1920 + 0] = t;
    • map[3 * 1920 + 1] = t;
    • ...
    • map[3 * 1920 + 31] = t;
      • 魔法发生了！Memory coalescing

void kernel(int row, int col) {
    t = ...;  // 任意 C/C++ 代码
    map[row * 1920 + col] = t;
}

当然，这也意味着 CUDA 程序很难写

void kernel(int row, int col) {
    t = ...;  // 任意 C/C++ 代码

    map[row * 1920 + col] = t;
    map[row + col * 1080] = t;
    // 等价的写法，可能导致性能下降
    // SIMT 处理分支也并不擅长
}

难写，意味着有收益

渲染 $25600 \times 25600$ 的 Mandelbrot Set

• Ryzen 5 9600X: 25.1s @ 65Watt
• RTX 4060Ti (16GB): 6.1s @ 42Watt
  • 这已经是对 CUDA “很不友好” 的程序了
    • 不确定的 while 循环

我们甚至可以直接阅读 GPU 上的代码

• 经过编译器优化，比大家想象得要短
• (GPU 并不擅长执行数据中心中的业务逻辑)

CUDA 擅长什么呢……？

固定模式、海量计算的，它最擅长

• cuBLAS: Basic Linear Algebra
  • DeepGEMM 就是其中重要的 matrix multiplication

看看 gpt.c?

• 代码怎么都这么像矩阵乘法？
  • 没错！固定模式、海量的计算！

AlexNet 一战成名

• AlexNet “卷积神经网络” (~60M 参数) 开启了深度学习的时代
  • 1D: $(f * g)[n] = \sum_{m} f[m] \cdot g[n-m]$
  • Richard Sutton: The Bitter Lesson
    • (在正确的时间，心无旁骛地去做一件事)