# getting-started-guide-and-introduction-to-MXMACA **Repository Path**: clausade/getting-started-guide-and-introduction-to-mxmaca ## Basic Information - **Project Name**: getting-started-guide-and-introduction-to-MXMACA - **Description**: 克隆自https://github.com/bxttttt/getting-started-guide-and-introduction-to-MXMACA，方便学校网络访问 - **Primary Language**: C++ - **License**: MIT - **Default Branch**: master - **Homepage**: None - **GVP Project**: No ## Statistics - **Stars**: 0 - **Forks**: 2 - **Created**: 2024-06-14 - **Last Updated**: 2024-06-14 ## Categories & Tags **Categories**: Uncategorized **Tags**: None ## README # 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 快速傅里叶变换库。