什么是并行

计算机中程序的运行方式主要有三种：

串行、并发、并行

• 串行：同一时刻只有一个任务在执行，必须等前一个任务彻底结束，后一个任务才能开始。
• 并发：多个任务都已经开始，但同一时刻只有一个任务在执行，往往通过轮流转换的方式来实现。
• 并行：多个任务在同一时刻同时执行。

计算机如何并行

计算机中的运算部件可以并行的主要有两种：CPU 和 GPU。

CPU与GPU对比

CPU 的并行主要通过多核来执行。一台现代的 CPU 通常有多个核心，每个核心都能独立执行指令，这样可以把不同的任务分配到不同的核心上同时运行。如操作系统、Web 服务等。

GPU 的并行是通过将一个大任务拆分成海量相似的小任务，让许多计算单元同时做同样模式的工作。GPU 的内部拥有大量更轻量的执行单元，通过牺牲单个线程的灵活性换来极强的批量吞吐能力。

CUDA 编程

• C/C++语法：宿主语言为 C/C++，使用范围广，应用领域对口，新语法的学习门槛较低。
• 与 CPU 端协作：CPU 负责整理程序执行和处理逻辑等。
• SIMT 模式：Single Instruction, Multiple Threads，只需要编写一个指令，一个指令可同时被多个线程执行。
• 自动调度：根据设定的执行参数自动调度资源（优化）。

GPU 硬件单元

GPU 的硬件单元指的是流式多处理器（Streaming Multiprocessor， SM），是 GPU 内部的核心工作单元，GPU 会将所有的线程分配到一个个 SM 上执行。其中单个 SM 中又主要有 CUDA Core 和 Tensor Core：

SM

CUDA Core 通常处理的是一个线程的一条普通算术指令，如整数运算、浮点运算、加减乘除、逻辑判断、地址计算以及一些通用指令等，而 Tensor Core 主要用于专门加速矩阵乘加操作的，其最典型的操作即：

也就是矩阵乘法再加累加。这类操作在深度学习里极其常见，因为全连接层本质上是矩阵乘法，卷积在很多实现里也会转成矩阵乘法，Transformer 里的 attention、MLP 也高度依赖 GEMM。

CUDA 编程运行流程

CUDA代码执行流程

CPU 首先进行程序初始化、数据准备等工作，而后将数据从主机内存拷贝至设备全局内存。而后 SM 读取全局内存中的数据，并行执行计算任务，再将结果写回全局内存，最后将结果从设备全局内存拷贝回主机内存。

CPU 加法

CPU 加法的过程主要如下：准备和初始化数据，定义加法函数，依靠循环来进行所有的元素加法，调用函数后验证结果。

CPU 加法→GPU 加法

CPU先准备和初始化数据：

然后将数据传输到 GPU，类似于 C 语言中的申请堆内存，GPU 中的内存也需要先申请再使用：

内存申请成功后传输数据：

而后 GPU 从全局内存中读取并计算后写回。这里的计算是要写在 GPU 上运行的加法函数并进行调用，在 CPU 的加法运算中使用了循环，而在 GPU 中要使用的是 SIMT，SIMT 指挥每个线程，需要组织结构和编号。在 CUDA 中主要是 Grid→Block→Thread 的层级关系。

• Thread 是最小的执行单位，即一个线程，通常负责处理一小份数据，每个 Thread 中通常有自己的局部变量、寄存器上下文、程序计数状态和 threadIdx。
• Block 是一组 Thread 的集合，CUDA 不会让单个线程独自运行，而是将线程先组织成 Block，一个 Block 的线程中有几个比较重要的特性：它们可以通过 shared memory 共享数据，可以用 __syncthreads() 做块内同步，还会被分配到同一个 SM 上执行，SM 上可以同时驻留多个 Block。不同的 Block 之间一般不直接同步。Block 还有自己的索引：blockIdx.x,blockIdx.y,blockIdx.z，每个 Block 内部的线程也有自己的局部编号：threadIdx.x,threadIdx.y,threadIdx.z。Block 的尺寸在进行 kernel lauch 时指定，如 dim3 blockDim(256) 和 dim3 blockDim(16,16) 都是表示每个 block 有 256 个线程。
• Grid 是一次 kernel 启动时，所有 block 的总集合。即每调用一次 kernel<<gridDim, blockDim>>(...); 里面的 gridDim 表示 Block 的数量，blockDim 表示每个 Block 中有多少个 Thread。

因此就需要定义 Block 的数量和大小来指挥线程并行计算，还要定义 GPU 上的加法函数，结合定义的信息调用 GPU 加法函数。

其中核函数定义为：

其中的 dim3 是 CUDA 表示线程层级结构的类型，<<>> 中的内容传递线程层级信息给核函数，核函数（kernel）即设备侧的入口函数，__global__ 表示这是一个核函数，最后将结果拷回主机：

完整的代码如下：

最后通过以下命令编译运行即可：

nvcc 是 CUDA 的编译器，是 CUDA Toolkit 的一部分。其中编译的流程大致如下：

CUDA代码编译流程

nvcc 将同一份 .cu 源文件拆分成两条路径来处理：一条是 host 路径，用于生成 CPU 侧的目标代码，另一条是 device 路径，用于生成 GPU 侧的目标代码。具体而言，nvcc 会先对源文件中的 device 部分进行处理，生成 PTX 或 cubin 这类设备端中间结果，并将其封装进 fatbinary；随后再对 host 部分重新预处理，把 CUDA 的扩展语法转换成宿主编译器能够理解的标准 C/C++ 形式，同时把前面生成的 fatbinary 嵌入进去，最后再交给 g++、clang++ 或 MSVC 这样的 host compiler 生成 host object。

Nsight Systems（CLI）

英伟达提供了一个系统级性能分析工具 Nsight Systems，便于对编写的 CUDA 程序进行性能评估。

通过以上命令就能够对运行的程序进行性能评估。观察得到的结果是 HtoD 和 DtoH 的耗时远远大于核函数的时间。设定的指标为：

要进一步优化，就需要从 HtoD 和 DtoH 这两个时间入手。

设备信息

一般而言，启用越多的线程，就能获得越大的并行度，越大的并行度也就能获取越好的性能。一个设备最多能够启用的线程数量可以通过网络/官网查询获取，也可以通过代码获取，NVIDIA 提供了 cudaDeviceProp 来通过代码查询。