PMPP Reading Notes

Chapter1 - Introduction

• 异构计算：CPU 的设计哲学是 latency oriented，为低延时处理顺序任务而设计，GPU 的设计哲学是 throughput oriented，为高吞吐处理大量数据而设计。

• 为什么我们需要并行：现实应用对速度的需求不断增长，硬件发展的摩尔定律失效使得我们不能再仅依赖 CPU 硬件的进步，需要利用并行能力强的 GPU / 其他加速器来保持软件应用速度和能力的增长。

• 并行计算的挑战：
• 并行算法的设计是困难的（相比于同等算法复杂程度的顺序算法设计）
• 有些程序的执行速度受限与内存：受限于内存速度的程序称为 Memory Bounded，与之相对的是受限于计算速度的程序称为 Compute Bounded.
• 并行算法更容易受到输入数据分布的影响，对于不同分布特征的输入数据性能影响可能很大。
• 并行算法需要考虑线程之间的协作。

• 本书的目标：
• 学习并行编程以提高性能的方式。
• 学习如何编写正确可靠的并行程序。
• 学习如何编写并行程序以对未来硬件发展维持 scalability.

Chapter2 - Heterogeneous data parallel computing

• 数据并行：数据并行来源于那些我们可以独立地处理数据集的不同的部分的场景，一个简单的例子是将彩色图像转换为灰度表示，这个过程的计算中每个像素是独立的。数据并行是使得并行程序获得 scalability 的一个重要手段。

• 任务并行：任务并行来源于一个大任务的不同小任务可以被独立完成的场景。不相关的 I/O 和计算就是一个任务并行的例子。

• CUDA C 程序结构：
• CUDA C 是一个 language extension，通过基于 ANSI C 及 CUDA 相关的扩展语法用于并行程序设计。
• 一个 CUDA C 程序包括 host code(CPU 上执行的代码) 和 device code（GPU 上执行的代码），通常是从 host code 开始，某个时刻 launch device 上的 kernel(在 device 上运行的函数)，计算完成之后再拷贝结果回到 host，这个过程可能持续几次。
• 上图是一个简化的流程，在实际应用中 CPU 和 GPU 可以是异步的。
• thread 描述了一个处理单元上的顺序执行过程，一个 CUDA 程序的并行计算过程是在调用 kernel 的时候发起的，通过 CUDA 的运行时机制创建了一个 thread grid，其中有若干 thread，以数据并行的方式进行计算。

• Kernel Functions & Thread:
• 当 Kernel 被调用时，通过 CUDA 运行时发起一个 thread grid，由若干 thread blocks 组成，每个 block 由相同数量的 thread 组成。在 kernel 中通过运行时环境设置的内置变量来访问数据的不同部分。
• block 和 thread 都可以有 1, 2, 3 维的组织方式，有 x，y，z 三个维度的坐标，组织方式通常和数据的组织方式对应。通常来说，考虑到硬件效率因素，block 的每个维度最好是 32 的倍数。
• 每个 thread 可以通过三个内置变量 blockDim, blockIdx 和 threadIdx 来访问数据的不同部分，这对应了串行版本中的每次循环，在 CUDA 程序中它们被并行执行。
• CUDA extension 语法提供了不同设备上的函数声明的方式，__global__ 定义了一个 kernel function, __device __ 定义了一个 device 上可以调用，但是不发起新 thread 的函数，__host__ 定义了 host 上运行的函数。对于一个函数可以同时使用 __host __ 和 __device __，生成两个版本的代码。
• 调用 kernel function 需要在函数 identifier 和调用参数之间添加 <<< … >>> 包裹的 configuration parameters. 下图给出了 host 上调用 kernel 的代码，在实际应用场景中可能不会这么写，因为频繁内存分配释放和拷贝的代价已经超过了计算的代价。

• 编译过程：
• CUDA C 在 ANSI C 上做了若干拓展语法，因此一个使用了 CUDA C extension 语法的程序就不是一个 C 编译器能够接受的合法程序，需要使用支持 CUDA C extension 的编译器，比如 NVCC (NVIDIA C Compiler) 来编译程序。
• 对于一个 CUDA C 程序，其中的 host 代码会被 host 的 C 编译器编译为目标代码，和 device 相关的部分被 NVCC 编译成 PTX 文件，PTX 文件会被运行时组建进一步编译成目标代码并在 CUDA-enabled 的 GPU 上执行。

• 总结：
• CUDA C 在 ANSI C 的基础上提供了若干语言扩展，以支持并行计算。
• 函数声明：__host __, __device __, __global __
• Kernel 调用：f<<< … >>>(args)
• 内置变量：threadIdx, blockIdx, blockDim
• 运行时 API：cudaMalloc, cudaFree, cudaMemcpy …

Chapter3 - Multidimensional grids and data

• 多维 grid 组织：grid 中的 block 和 block 中的 thread 都可以按照 3 个维度组织，在调用 kernel 的执行参数中指出了其维度 <<< gridDim, blockDim >>>，其中两者都是一个 dim3 对象，也就是一个 3 维向量，在上一章中我们直接使用了两个整形表达式，这利用了 C++ 构造函数的机制，隐式地构造了一个 dim3 对象，只是其 y 和 z 维度为 1.

• 将 threads 映射到高维数据：由于 CUDA C 是在 ANSI C 上的语言拓展，其在使用多维数组的时候需要在编译阶段知道列大小，因此在实际使用多维数组时需要线性化展平成一维数组保存，在 CUDA C 中一般采用 row major layout, 每行的元素连续存储。

• 例子：RGB2Grayscale, Image Blurring, Matrix Multiplication

• 总结：
• CUDA C 在组织 grid 和 block 时提供三维布局，这使得 thread 映射到高维数据更加方便，在每个 thread 中通过内置变量索引不同部分的数据。
• 在使用高维数组时需要线性化展平为一维来使用。

Chapter 4 - Compute Architecture and Scheduling

• 现代 GPU 架构：
• 由多个 Steaming Multiprocessor (SMs) 组成，每个 SM 包含控制单元，以及多个运算单元（CUDA cores）
• 每个 SM 都有片上存储，包括寄存器和共享内存；片外存储 HBM

• Block 调度：
• 调用 kernel 时，运行时会启动大量 threads，按照 block-by-block 的方式分配到 SM 上
• 每个 SM 能够容纳的 block 数量有限，通常一个 grid 包含非常多 block，为了让所有 block 都能够最终被执行，运行时需要维护一个待执行的 block 列表
• 这种 block-by-block 调度要求代码不管按照什么相对顺序执行 block 都能得到正确结果

• 块内同步: 使用 __syncthreads() 来作为 barrier synchronization 同步一个 block 内的 threads

• Warps & SIMD Hardware:
• 在一个 block 内 threads 被进一步分为多个 warp；warp 是一组相邻的 threads，是 block 内 threads 调度的基本单位
• 在目前的实现中 warp size 多为 32
• 一个 warp 内的 threads 按照 SIMD 的方式执行，共享控制单元（这使得更多空间可以给运算单元）
• block 大小如果不是 warp 大小的倍数需要 padding，padding 的 threads 为 inactive 状态

• Control Divergence:
• 一个 warp 内的 threads 如果执行了不同的控制流，那么控制流的每个 path 会被执行，当在执行某个 path 时，不走这个 path 的线程处于 inactive 状态
• divergence 造成硬件利用率下降，需要避免
• divergence 的一个 implication 是不应该默认 warp 内的 threads 步调都一致，如果需要同步点需要使用 __syncwarp() 来同步。

• Warp scheduling & latency hiding:
• 一个 SM 分配到的 threads 数量往往需要比它拥有的 CUDA core 的数量多很多（over-subscription）
• 当某个 warp 在等待一些高延迟操作的时候可以先上下文切换到可执行的 warp，来隐藏这些操作的延迟，over-subscription 保障大部分时候能有这种待执行的 warp
• 和 cpu 的线程不同，在 gpu 的场景下，上下文切换开销几乎可以忽略，因为并不需要保存线程状态到更高层级的存储，而是让线程状态停留在原来的地方，因此可以几乎无开销频繁进行上下文切换。

• Resource partitioning & occupancy:
• occupancy: 分配给一个 SM 的 warp 数量 / SM 支持的最大 warp 数量
• 通常无法达到满 occupancy 是因为资源限制，比如寄存器，共享内存，threads slot，block slot

Chapter 5 - Memory Architecture and Data Locality

• Pytorch 程序运行时间分解：
• python processing
• data administrative overhead：allocating tensor data structure etc
• data acquisition (I/O): between device and host, check this before diving to GPU optimization.
• GPU computation: fix cost (kernel launch, etc), memort access, “real” computation
• rule of thumb:
• 如果 gpu 利用率没到快满，检查 I/O
• 如果 tensor 不是太小，python processing 和 data administrative 开销基本在 10% 以下

• memory access as a bottleneck:
• Eager Pytorch 对于每一个 operation 都会：load 数据，进行计算，store 结果。
• 不同层次的 memory access 开销差距很大，register < shared memory << HBM
• 对于一个连续的 operation 序列，显然我们可以通过把多个计算一起执行完再 store 结果来减少访问 HBM 的开销，这就是 kernel fusion.

• roofline model:
• 关键指标：arithmetic intensity = #ops / #bytes
• 高 AI → 更可能是 compute bound; 低 AI → 更可能是 memory bound

• memory types:
• register：per threads ，访问速度最快
• shared memory：per block，高速 SRAM 存储
• local memory：per thread，但是 HBM 的一部分
• constant memory：per grid，支持快速只读访问
• HBM: per grid 全局内存，最慢

• tiling: 通过把一个操作变成分块运算，每个块计算的时候用 shared memory 缓存输入以复用，减少 HBM 访问

Key Takeaway from Ch.4 and Ch.5

• GPU 如何组织 threads, warp, block, grid

• 尽量提高 occupancy

• 避免 thread 发散

• roofline model

• 尽量减少 HBM 访问