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 中通过内置变量索引不同部分的数据。
• 在使用高维数组时需要线性化展平为一维来使用。