README.md · tianbuaa/getting-started-guide-and-introduction-to-MXMACA

# getting-started-guide-and-summary-of-MXMACA

## CPU VS GPU

CPU，即中央处理器，由数百万个晶体管构成，可以具有多个处理核心，是计算机系统的运算和控制核心。CPU涉及到通用计算，适合少量的复杂计算。CPU虽然处理核心远没有GPU多，但是可以将核心集中在单个任务上并快速完成工作。

GPU，即图形处理器，由许多更小、更专业的核心组成的处理器。适合大量的简单运算。GPU最初是用来加速3D渲染任务，但是随着时间的推移，这些固定功能的引擎变得更加可编程、更加灵活。虽然图形和日益逼真的视觉效果仍然是GPU的主要功能，但GPU也已发展成为更通用的并行处理器，可以处理越来越多的应用程序。

| CPU                                | GPU                              |
| ---------------------------------- | -------------------------------- |
| 通用组件，负责计算机的主要处理功能 | 专用组件，主要负责图形和视频渲染 |
| 核心数：2-64                       | 核心数：数千                     |
| 串行运行进程                       | 并行运行进程                     |
| 更适合处理一项大任务               | 更适合处理多个较小的任务         |

### 加速深度学习和人工智能

GPU或其他加速器非常适合用神经网络或大量特定数据（e.g. 2D图像）进行深度学习训练。

GPU加速方法已经适用于深度学习算法，可以显著提升算法性能。

## 基本概念的解释

内存部分的解释详见MXMACA内存模型和管理。

### 主机端（host）

CPU所在的位置称为主机端。

可以简单理解为CPU。

### 设备端（device）

GPU所在的位置称为设备端。

可以简单理解为GPU。

主机和设备之间通过PCIe总线连接，用于传递指令和数据，让CPU和GPU一起来协同工作。

### 加速处理器（Accelerated Processors，AP）

每个AP都能支持数千个GPU线程并发执行。

执行具体的指令和指令和任务。

### 核函数（kernel）

核函数在设备端执行，需要为一个线程规定所进行的计算和访问的数据。当核函数被调用时，许多不同的MXMACA线程并行执行同一计算任务。

在设备侧（GPU）执行，可以在设备侧（GPU）和主机侧（CPU）被调用。

### 线程（thread）

一般通过GPU的一个核进行处理。

每个线程是Kernel的单个执行实例。在一个block中的所有线程可以共享一些资源，并能够相互通信。

### 线程束（wave）

GPU执行程序时的调度单位。

64个线程组成一个线程束，线程束中每个线程在不同数据集上同时执行相同的指令。

### 线程块（thread block）

由多个线程组成。可以是一维、二维或三维的。

各block是并行执行的。

同一个线程块内的线程可以相互协作，不同线程块内的线程不能协作。

当启动一个核函数网格时，它的GPU线程会被分配到可用的AP上执行。一旦线程块被调度到一个AP上，其中的线程将只在该指定的AP上并发执行。

多个线程块根据AP资源的可用性进行调度，可能会被分配到同一个AP上或不同的AP上。

### 线程网格（grid）

多个线程块可以构成线程网格。

和核函数（kernel）的关系：启动核函数（kernel）时，会定义一个线程网格（grid）。

网格可以是一维的、二维的或三维的。

### 流（stream）

相当于是GPU上的任务队列。

同一个stream的任务是严格保证顺序的，上一个命令执行完成才会执行下一个命令。

不同stream的命令不保证任何执行顺序。部分优化技巧需要用到多个stream才能实现。如在执行kernel的同时进行数据拷贝，需要一个stream执行kernel，另一个stream进行数据拷贝。

## 基本编程模型

1. 用户可以通过调用动态运行时库，申请、释放显存，并在内存和显存间进行数据拷贝。

2. 典型的MXMACA程序实现流程遵循以下模式：

1. 把数据从CPU内存拷贝到GPU内存；
   2. 调用核函数对GPU内存的数据进行处理；
   3. 将数据从GPU内存传送回CPU内存。

3. 用户可以编写kernel函数，在主机侧调用kernel函数，调用将创建GPU线程。

1. 用户可以在Kernel Launch时分别指定网格中的线程块数量、线程块中包含的线程数量。当用户指定的线程数量超过64，这些线程会被拆分成多个线程束，并在同一个AP上执行，这些线程束可能并发执行，也可能串行执行。
   2. 每个GPU线程都会完整执行一次kernel函数，kernel函数可以对显存进行读、写等操作，也可以调用设备侧函数对显存进行读、写等操作。不同的GPU线程可以通过内置变量进行区分，只需要通过读取内置变量，分别找到线程块的位置、线程的位置，就可以给每一个线程唯一地标识ThreadIdx（可以参考后文，相关的几个内置变量）。

4. 相关的几个内置变量

1. `threadIdx`，获取线程`thread`的ID索引；如果线程是一维的那么就取`threadIdx.x`，二维的还可以多取到一个值`threadIdx.y`，以此类推到三维`threadIdx.z`。可以在一个线程块中唯一的标识线程。
   2. `blockIdx`，线程块的ID索引；同样有`blockIdx.x`，`blockIdx.y`，`blockIdx.z`。可以在一个网格中唯一标识线程块。
   3. `blockDim`，线程块的维度，同样有`blockDim.x`，`blockDim.y`，`blockDim.z`。可以代表每个维度下线程的最大数量。
      1. 对于一维的`block`，线程的`threadID=threadIdx.x`。
      2. 对于大小为`（blockDim.x, blockDim.y）`的 二维`block`，线程的`threadID=threadIdx.x+threadIdx.y*blockDim.x`。
      3. 对于大小为`（blockDim.x, blockDim.y, blockDim.z）`的 三维 `block`，线程的`threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y`。
   4. `gridDim`，线程格的维度，同样有`gridDim.x`，`gridDim.y`，`gridDim.z`。可以代表每个唯独下线程块的最大数量。

5. 常用的GPU函数

1. `mcMalloc()`

负责内存分配。类似与C语言中的`malloc`。不过mcMalloc是在GPU上分配内存，返回device指针。

2. `mcMemcpy()`

负责内存复制。

可以把数据从host搬到device，再从device搬回host。

3. `mcFree()`

释放显存的指针。

（可以参考示例代码）

## 基本硬件架构及其在Kernel执行中的作用

## MXMACA内存模型和管理

### MXMACA内存模型

MXMACA的内存是分层次的，每个不同类型的内存空间有不同的作用域、生命周期和缓存行为。一个内核函数中，每个线程有自己的私有内存，每个线程块有自己工作组的共享内存并对块内的所有线程可见，一个线程网格中的所有线程都可以访问全局内存和常量。可以参考下图：

书里提到了它们的初始化方式，这里主要介绍它们的用途、局限性。

#### 可编程存储器、不可编程存储器

根据存储器能否被程序员控制，可分为：可编程存储器、不可编程存储器。

可编程存储器：需要显示控制哪些数据放在可编程内存中。包括全局存储、常量存储、共享存储、本地存储和寄存器等。

不可编程存储器：不能决定哪些数据放在这些存储器中，也不能决定数据在存储器中的位置。包括一级缓存、二级缓存等。

#### GPU寄存器

寄存器延迟极低，对于每个线程是私有的，与核函数的生命周期相同。

寄存器是稀有资源，使用过多的寄存器也会影响到性能，可以添加辅助信息控制限定寄存器数量。

书中也提到了一些方式，可以让一个线程束内的两个线程相互访问对方的寄存器，而不需要访问全局内存或者共享内存，延迟很低且不消耗额外内存。

#### GPU私有内存

私有内存是每个线程私有的。

私有内存在物理上与全局内存在同一块储存区域，因此具有较高的延迟和低带宽。

#### GPU线程块共享内存

共享内存的地址空间被线程块中所有的线程共享。它的内容和创建时所在的线程块具有相同生命周期。

共享内存让同一个线程块中的线程能够相互协作，便于重用片上数据，可以降低核函数所需的全局内存带宽。

相较于全局内存，共享内存延迟更低，带宽更高。

适合在数据需要重复利用、全局内存合并或线程之间有共享数据时使用共享内存。

不能过度使用，否则会限制活跃线程束的数量。

书里也提到了共享内存的分配、共享内存的地址映射方式、bank冲突以及最小化bank冲突的方法。bank冲突时，多个访问操作会被序列化，降低内存带宽，就没有什么并行的意义了。

#### GPU常量内存

常量内存在设备内存中，并在每个AP专用的常量缓存中缓存。

如果线程束中所有线程都从相同内存读取数据，常量内存表现最好，因为每从一个常量内存中读取一次数据，都会广播给线程束里的所有线程。

#### GPU全局内存

GPU中内存最大、延迟最高、最常使用。

可以在任何AP上被访问，并且贯穿应用程序的整个生命周期。

优化时需要注意对齐内存访问与合并内存访问。

## MXMACA程序优化

### 性能优化的目标

1. 提高程序执行效率，减少运行时间，提高程序的处理能力和吞吐量。
2. 优化资源利用率，避免资源的浪费和滥用。
3. 改善程序的响应时间。

### 程序性能评估

#### 精度

GPU 的单精度计算性能要远远超过双精度计算性能，需要在速度与精度之间选取合适的平衡。

#### 延迟

#### 计算量

如果计算量很小，或者串行部分占用时间较长，并行部分占用时间较短，都不适合用GPU进行并行计算。

### 优化的主要策略

#### 硬件性能优化

#### 并行性优化

可以通过设置线程块的大小、每个线程块的共享内存使用量、每个线程使用的寄存器数量，尽量提升occupancy。

#### 内存访问优化

##### 提高`Global Memory`访存效率

对齐内存访问：一个内存事务的首个访问地址尽量是缓存粒度（32或128字节）的偶数倍，减少带宽浪费。

合并内存访问：尽量让一个线程束的线程访问的内存都在一个线程块。

##### 提高`Shared Memory`访存效率

若`wave`中不同的线程访问相同的`bank`，则会发生bank冲突(bank conflict)，bank冲突时，`wave`的一条访存指令会被拆分为n条不冲突的访存请求，降低`shared memory`的有效带宽。所以需要尽量避免bank冲突。

#### 算法优化

1. 如何将问题分解成块、束、线程
2. 线程如何访问数据以及产生什么样的内存模式
3. 数据的重用性
4. 算法总共要执行多少工作，与串行化的方法之间的差异

#### 算数运算密度优化

1. 超越函数操作：可以查阅平方根等超越函数和加速函数，以及设备接口函数
2. 近似：可以在速度和精度之间进行折衷
3. 查找表：用空间换时间。适合GPU高占用率的情况，也要考虑到计算的复杂度，计算复杂度低时，计算速度可能大大快于低GPU占用下的内存查找方式。

#### 编译器优化

1. 展开循环
2. 常量折叠 e.g. 编译时直接计算常数，从而简化常数
3. 常量传播：将表达式中的变量替换为已知常数
4. 公共子表达式消除：将该类公共子表达式的值临时记录，并传播到子表达式使用的语句
5. 目标相关优化：用复杂指令取代简单通用的指令组合，使程序获得更高的性能

#### 其他

1. 用结构体数组（结构体的成员是数组），而不是数组结构体（数组的每个元素都是结构体）。
2. 尽量少用条件分支。CPU具有分支预测的功能，GPU没有这一功能，GPU执行if，else语句的效率非常低。因此只能让束内每一线程在每个分支都经过一遍（但不一定执行），当然如果所有线程都不用执行，就可以忽略这一分支。只要有一个线程需要执行某一个分支，其他线程即使不需要执行，也要等着一个线程执行完才能开始自己的计算任务。而且不同的分支是串行执行的，因此要减少分支的数目。
   1. 通过计算，去掉分支（可以参考书中8.3.4相关内容）。
   2. 通过查找表去掉分支。
   3. 尽量使`wave`块完美对齐，让一个`wave`里的所有线程都满足条件或者都不满足条件。
3. 引入一些指令级并行操作，尽可能终止最后的线程束以使整个线程块都闲置出来，并替换为另一个包含一组更活跃线程束的线程块。

### 优化性能需要考虑的指标

1. 最大化利用率
2. 最大化存储吞吐量
3. 最大化指令吞吐量
4. 最小化内存抖动
5. 时间消耗（整体运行所需时间、GPU和CPU之间的传输所需时间、核函数运行所需时间）

## MXMACA生态的人工智能和计算加速库

### mcBLAS

主要用于多种形式的计算。

`Level-1 Functions`定义了向量与向量、向量与标量之间的运算，还为多种数据类型（单精度浮点实数、单精度浮点复数、双精度浮点实数、双精度浮点复数）定义了专用的接口。

`Level-2 Functions`定义了矩阵与向量之间的运算。

`Level-3 Functions`定义了矩阵与矩阵之间的运算。是求解器和深度神经网络库的底层实现基础。

### mcDNN

提供常用深度学习算子。

### mcSPARSE

稀疏矩阵线性代数库。稀疏矩阵是指零元素数目远多于非零元素数目的矩阵。

可以用对应的接口完成稀疏矩阵线性代数运算。

### mcSOLVER

稠密矩阵线性方程组的求解函数库。

### mcFFT

快速傅里叶变换库。

tianbuaa/getting-started-guide-and-introduction-to-MXMACA

简介

发行版

贡献者

近期动态

tianbuaa/getting-started-guide-and-introduction-to-MXMACA .gitee-modal { width: 500px !important; }

简介

发行版

贡献者

近期动态

搜索帮助

tianbuaa/getting-started-guide-and-introduction-to-MXMACA