#GPUの内存体系及其优化指南
要在冯·诺依曼架构的硬件中实现高性能计算，最重要的两点就是：访存和计算。这两点分别对应着 IO bound 和 compute bound ，硬件系统的内存体系深刻影响着这两点。因此要实现软件层面的高性能计算，必须要对内存体系有深刻的理解。本篇主要讨论GPU的内存体系，并在此基础上进行CUDA的编程实践。
一、GPU的内存体系
内存的访问和管理编程语言的重要组成部分，也是实现高性能计算的重要环节。CUDA 内存模型结合了主机和设备的内存系统，具有完整的层次结构，并可以显式地进行控制和优化。
1.1 各级内存及其特点
下图展示了CUDA 内存模型的层次结构，每一种都有不同的作用域、生命周期以及缓存行为，接下来将逐一介绍：
CUDA 内存模型的层次结构
- 全局内存(global memory)
全局内存是GPU中最大、延迟最高、最长使用的内存，通常说的“显存”中的大部分都是全局内存。全局内存的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。
全局内存的主要角色是为核函数提供数据，并在主机与设备及设备与设备之间传递数据。可以用 cudaMemcpy函数将主机的数据复制到全局内存，或者反过来。如将中 M 字节的数据从主机复制到设备，操作如下：
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);全局内存变量可以被静态声明和动态声明， 如 静态全局内存变量由以下方式在任何函数外部定义 ：
__device__ T x; // 单个变量
__device__ T y[N]; // 固定长度的数组后续将会重点研究如何优化全局内存访问，以及如何提高全局内存的数据吞吐率。
- 常量内存(constant memory)
常量内存是指存储在片下存储的设备内存上，但是通过特殊的常量内存缓存（constant cache）进行缓存读取，常量内存为只读内存。常量内存数量有限，一共仅有 64 KB，由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程(一个线程块中相邻的 32 个线程)要读取相同的常量内存数据。
一个使用常量内存的方法是在核函数外面用 __constant__ 定义变量，并用 API 函数 cudaMemcpyToSymbol 将数据从主机端复制到设备的常量内存后 供核函数使用。
- 纹理内存(texture memory)和表面内存(surface memory)
纹理内存和表面内存类似于常量内存，也是一 种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读(表面内存也可 写)。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。
- 寄存器（register）
寄存器是线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上（on chip）的存储，用来存储一些线程的暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。
在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器(register)中。各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中，以便高效访问。在上期求和的例子中：
const int n = blockDim.x * blockIdx.x + threadIdx.x;
c[n] = a[n] + b[n];中的 n 就是一个寄存器变量。寄存器变量仅仅被一个线程可见。也就是说，每一个线程都有一个变量 n 的副本。虽然在核函数的代码中用了这同一个变量名，但是不同的线程中该寄存器变量的值是可以不 同的。每个线程都只能对它的副本进行读写。寄存器的生命周期也与所属线程的生命周期 一致，从定义它开始，到线程消失时结束。
- 局部内存(local memory)
局部内存和寄存器几乎一 样，核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。
虽然局部内存在用法上类似于寄存器，但从硬件来看，局部内存只是全局内存的一部 分。所以，局部内存的延迟也很高。每个线程最多能使用高达 512 KB 的局部内存，但使用 过多会降低程序的性能。
- 共享内存(shared memory)
共享内存和寄存器类似，存在于芯片 上，具有仅次于寄存器的读写速度，数量也有限。一个使用共享内存的变量可以 __shared__ 修饰符来定义。
不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。也 就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中 可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能 访问其他线程块的共享内存变量副本。共享内存的主要作用是减少对全局内存的访问，或 者改善对全局内存的访问模式。
以上内存的主要特点如下表所示：
- L1和L2 缓存
每个 SM 都有一个 L1 缓存，所有 SM 共享一个 L2 缓存。L1 和 L2 缓存都被用来存储局部内存和全局内存中的数据，也包括寄存器中溢出的部分，以减少延时。
从物理结构上来说，在最新的GPU架构中，L1 缓存、纹理缓存及共享内存三者是统一的。但从编程的角度来看，共享 内存是可编程的缓存(共享内存的使用完全由用户操控)，而 L1 和 L2 缓存是不可编程的缓存(用户最多能引导编译器做一些选择)。
1.2 SM 构成及典型GPU的对比
一个 GPU 是由多个 SM 构成的。一个 SM 包含如下资源:
- 一定数量的寄存器。
- 一定数量的共享内存。
- 常量内存的缓存。
- 纹理和表面内存的缓存。
- L1缓存。
- 线程束调度器(warp scheduler) 。
- 执行核心，包括:
- 若干整型数运算的核心(INT32) 。
- 若干单精度浮点数运算的核心(FP32) 。
- 若干双精度浮点数运算的核心(FP64) 。
- 若干单精度浮点数超越函数(transcendental functions)的特殊函数单元(Special Function Units，SFUs)。
- 若干混合精度的张量核心(tensor cores)
下图是 H100 的 SM 结构图，可以按图索骥找到上述的对应部分
H100 的 SM 结构图
为了更好理解内存体系与性能的关系，下表列举了当前几款主流的GPU产品的数据
1.3 GPU 之外：近存计算与存算一体
在GPU的层次结构之外，为了降低访存成本，获得更高的性能，近存计算与存算一体逐渐成为热门的方向。
近存计算: 以 Graphcore IPU 为例
IPU芯片没有高速的片外存储，而是把存储放到了片上。整个芯片由1472个核心组成，称作Tile。每个Tile由独立的计算单元和存储单元组成，整个片上存储是分布式的。每个Tile中有624KB的SRAM，因此整个芯片的存储大小为624KB * 1472 = 900MB。
IPU芯片采用纯分布式的架构，每个Tile有自己的存储和计算资源，采用MIMD的计算架构（与NVIDIA CUDA的SIMT不同），每个Tile可以独立地执行不同的指令，可以独立地访存。Tile和Tile之间的memory不能共享访问，只能访问自己Tile内部的memory，叫做 local memory。因此整个芯片的访存带宽 = Tile 访存带宽 * Tile数量。
Graphcore IPU 结构图示
存算一体: 以 后摩智能 H30 为例
存算一体或者存内计算的核心思想是，通过对存储器单元本身进行算法嵌入，使得计算可以在存储器单元内完成。
鸿途™H30包含多个存算单元，既能存储数据也能处理数据，打破传统芯片性能瓶颈并提升能效比，物理算力可达到256TOPS，实现大算力、低功耗、低成本。
存内计算示意图
二、通过规约(Reduction)操作理解GPU内存体系
reduce 算法可以描述为
由于GPU没有针对global数据的同步操作，只能针对block的数据进行同步。所以，一般而言将reduce分为两个阶段，其示意图如下：
假设给定一个长度为N的数组，需要计算该数组的所有元素之和。首先需要将数组分为m个小份。而后，在第一阶段中，开启m个block计算出m个小份的reduce值。最后，在第二阶段中，使用一个block将m个小份再次进行reduce，得到最终的结果。
以下实验代码开源在：ifromeast/cuda_learning: learning how CUDA works (github.com)
2.1 仅使用全局内存实现规约
对于数组归约的并行计算问题，要从一个数组出发，最终得到一个数。所以，必须使用某种迭代方案。假如数组元素个数是 2 的整数次方，可以将数组后半部分的各个元素与前半部分对应的数组元素相加。如果重复此过程，最后得到的第 一个数组元素就是最初的数组中各个元素的和。这就是所谓的折半归约(binary reduction) 法。
其实现如以下代码所示：
void __global__ reduce_global(real *d_x, real *d_y)
{
const int tid = threadIdx.x;
real *x = d_x + blockDim.x * blockIdx.x;
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
{
if (tid < offset)
{
x[tid] += x[tid + offset];
}
__syncthreads();
}
if (tid == 0)
{
d_y[blockIdx.x] = x[0];
}
}其中有以下几点需要注意：
- 同步函数
__syncthreads是为了保证一个线程块中的所有线程(或者说所有线程束)在执行该语句后面的语句之 前都完全执行了该语句前面的语句。 - 将(动态)数组 d_x 中第
blockDimx.x * blockIdx.x个元素的地址赋给指针x，该句等价于real *x = &d_x[blockDim.x * blockIdx.x]; - for 循环内在各个线程块内对其中的数据独立地进行归约，同一个线程块内的线程按照代码出现的顺序执行指令 。不同线程块之间由于处理的数据不同，因此不需要同步。
- 在 offset 的计算过程中使用了位操作，对于2的幂而言，这样做更加高效。
- 该核函数仅仅将一个长度为的数组 d_x 归约到一个长度为/128的数组 d_y 。
- 因为 global 修饰符的限制，d_y 和 x 均是全局内存的变量。
2.2 使用共享内存实现规约
由于全局内存访问速度最低，因此性能较低，本节将使用 对整个线程块可见的共享内存来实现同样的规约操作。使用 __shared__ 修饰共享内存变量 s_y，同时其长度等于线程块大小。然后将全局内存中的数据复制到共享内存中，每个线程块都有一个共享内存变量的副本。
具体实现如下所示， 当 blockIdx.x 等于 0 时，将全局内存中第 0 到第 blockDim.x - 1 个数组元素复制给 第 0 个线程块的共享内存变量副本，以此类推。
const int tid = threadIdx.x;
const int idx = blockIdx.x * blockDim.x + threadIdx.x;
__shared__ real s_y[128];
s_y[tid] = (idx < N) ? d_x[idx] : 0.0;
__syncthreads();上述过程可表示如下，需要说明的是，GPU需要分配两种资源，一个是存储资源，一个是计算资源。计算资源其实是根据thread数量来确定的，一个block中分配128个thread线程，32个线程为一组（即一个 warp），绑定在一个SIMD单元。所以128个线程可以简单地理解为分配了4组SIMD单元。
在核函数中对共享内存访问的次数越多，则由使用共享内存带来的加速效果越明显。在我们的数组归约问题中，使用共享内存相对于仅使用全局内存还有两个好处: 一个是不再要求全局内存数组的长度 N 是线程块大小的整数倍，另一个是在规约的过程中不会改变全局内存数组中的数据(在仅使用全局内存时，数组 d_x 中的部分元素被改变)。这两点在实际的应用中往往都是很重要的。
2.3 使用动态共享内存实现规约
在上边使用共享内存数组时，指定了一个固定的长度(128) ，使这个长度与核函数的执行配置参数 block_size (也就是核函数中的 blockDim.x)是一样的。这种静态的方式可能会导致错误的发生，因此有必要使用动态操作。
将静态共享内存改成动态共享内存，只需要做以下两处修改:
- 在调用核函数的执行配置中写下第三个参数:
<<<grid_size, block_size, sizeof(real) * block_size>>>前两个参数分别是网格大小和线程块大小，第三个参数就是核函数中每个线程块需要 定义的动态共享内存的字节数， 其默认值为零。
- 要使用动态共享内存，还需要改变核函数中共享内存变量的声明方式
extern __shared__ real s_y[];它与之前静态共享内存的声明方式 有两点不同:第一，必须加上限定词 extern; 第二，不能指定数组大小。
- 种方式都完成后，我们可以编译运行代码
nvcc reduce_gpu.cu -o reduce通过计时函数可以看到，每种方法的完整计算总时间都在 7.5 ms 左右。
然后通过 nvprof 命令查看GPU各部分的耗时
nvprof ./reduce结果如下，由于该例子中不存在对内存的频繁读写，因此以上几种方式的性能差别不大
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 98.25% 30.8438s 300 102.81ms 80.686ms 259.15ms [CUDA memcpy HtoD]
0.44% 137.86ms 100 1.3786ms 1.3756ms 1.3822ms reduce_global(float*, float*)
0.44% 137.66ms 300 458.85us 343.93us 863.09us [CUDA memcpy DtoH]
0.43% 136.51ms 100 1.3651ms 1.3647ms 1.3666ms reduce_shared(float*, float*)
0.43% 136.50ms 100 1.3650ms 1.3647ms 1.3664ms reduce_dynamic(float*, float*)如果要放大访存速度的差别，可以使用双精度，编译方式如下：
nvcc reduce_gpu.cu -DUSE_DP -o reduce_dp此时可以看到全局内存的性能出现了明显下降：
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 98.65% 60.8587s 300 202.86ms 189.54ms 304.60ms [CUDA memcpy HtoD]
0.65% 398.54ms 300 1.3285ms 1.2460ms 2.9135ms [CUDA memcpy DtoH]
0.26% 157.32ms 100 1.5732ms 1.5677ms 1.5803ms reduce_global(double*, double*)
0.22% 137.32ms 100 1.3732ms 1.3716ms 1.3746ms reduce_shared(double*, double*)
0.22% 137.31ms 100 1.3731ms 1.3714ms 1.3750ms reduce_dynamic(double*, double*)2.4 规约的其他优化方法
2.4.1 使用原子函数
在前边几个版本的数组归约函数中，核函数并没有做全部的计算，而只是将一个长 一些的数组 d_x 变成了一个短一些的数组 d_y，后者中的每个元素为前者中若干元素的和。在调用核函数之后，将短一些的数组复制到主机，然后在主机中完成了余下的求和。所有这些操作所用时间约为 7.5 ms (单精度的情形)，而在 GPU 中实际计算的时间仅为 1.4ms 左右。
如果能在 GPU 中计算出 最终结果，则有望显著地减少整体的计算时间，提升程序性能。有两种方法能够在 GPU 中 得到最终结果，一是用另一个核函数将较短的数组进一步归约，得到最终的结果(一个数值); 二是在先前的核函数的末尾利用原子函数进行归约，直接得到最终结果。本节讨论原子函数的方法。
在前述的几种方式中，核函数最后执行的都是
if (tid == 0)
{
d_y[bid] = s_y[0];
}即 将每一个线程块中归约的结果从共享内存 s_y[0] 复制到全局内 存 d_y[bid]。为了将不同线程块的部分和 s_y[0] 累加起来，存放到一个全局内存地址，我 们尝试将上述代码改写如下:
if (tid == 0) {
d_y[0] += s_y[0];
}问题是该语句在每一个线程块的第 0 号线程都会被执行， 但是它们执行的次序是不确定的。在每一个线程中，该语句其实可以分解为两个操作:首 先从 d_y[0] 中取数据并与 s_y[0] 相加，然后将结果写入 d_y[0]。不管次序如何，只有当 一个线程的“读-写”操作不被其他线程干扰时，才能得到正确的结果。如果一个线程还未 将结果写入 d_y[0]，另一个线程就读取了 d_y[0]，那么这两个线程读取的 d_y[0] 就是一 样的，这必将导致错误的结果。
要得到所有线程块中 的 s_y[0] 的和，必须使用原子函数，其用法如下 ：
if (tid == 0)
{
atomicAdd(d_y, s_y[0]);
}原子函数 atomicAdd(address, val) 的第一个参数是待累加变量的地址 address，第二个 参数是累加的值 val。该函数的作用是将地址 address 中的旧值 old 读出，计算 old + val， 然后将计算的值存入地址 address。这些操作在一次原子事务(atomic transaction)中完成， 不会被别的线程中的原子操作所干扰。原子函数不能保证各个线程的执行具有特定的次序， 但是能够保证每个线程的操作一气呵成，不被其他线程干扰，所以能够保证得到正确的结果。
使用原子函数后总时间变为 2.8 ms，相比于之前的方式，性能提升接近3倍，GPU中操作耗时如下：
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 66.95% 194.36ms 100 1.9436ms 1.8185ms 2.0081ms reduce(float const *, float*, int)
33.00% 95.797ms 101 948.48us 1.4400us 95.642ms [CUDA memcpy HtoD]
0.05% 133.09us 100 1.3300us 1.2160us 2.6560us [CUDA memcpy DtoH]GPU中计算耗时从1.4ms增加到1.8ms，这是由于d_y也在GPU中计算导致的，另外比较明显的是由于从GPU输出结果的数量变少，DevicetoHost 的时间占比大幅缩减。
2.4.2 使用线程束函数与协作组
线程束(warp) 是 SM 中基本的执行单元。一个线程束由32个连续线程组成，这些线程按照单指令多线程(SIMT)方式执行（即所有线程执行相同指令，每个线程在私有数据上操作）。这样如果在条件语句中，同一线程束中的线程执行不同的指令，就会发生线程束分化(warp divergence) ，导致性能出现明显下降。
在归约问题中，当所涉及的线程都在一个线程束内时，可以将线程块同步函 数 __syncthreads 换成一个更加廉价的线程束同步函数 __syncwarp。
for (int offset = blockDim.x >> 1; offset >= 32; offset >>= 1)
{
if (tid < offset)
{
s_y[tid] += s_y[tid + offset];
}
__syncthreads();
}
for (int offset = 16; offset > 0; offset >>= 1)
{
if (tid < offset)
{
s_y[tid] += s_y[tid + offset];
}
__syncwarp();
}当 offset >= 32 时，我们在每一次折半求和后使用线程块同步函 数 __syncthreads; 当 offset <= 16 时，我们在每一次折半求和后使用束内同步函数 __syncwarp。
另外还可以利用线程束洗牌函数进行归约计算，函数 __shfl_down_sync 的作用是将高线程号的数据平移到低线程号中去，这正是归约问题中需要的操作
for (int offset = 16; offset > 0; offset >>= 1)
{
y += __shfl_down_sync(FULL_MASK, y, offset);
}相比之前的版本，有两处不同。第一，在进行线程束内的循环之前，将共享内存中的数据复制到了寄存器。在线程束内使用洗牌函数进行规约时，不再需要明显地使用共享内存。因为寄存器一般来说比共享内存更高效，所以能用寄存器就当然用寄存器 了。第二，去掉了束同步函数，这是因为洗牌函数能够自动处理同步与读-写竞争问题。
协作组(cooperative groups)可以看作是线程块和线程束同步机制的推广，它提供了更为灵活的线程协作方式，包括线程块内部的同步与协作、线程块之间的(网格级的)同步与协作及设备之间的同步与协作。
使用协作组的功能时需要在相关源文件包含如下头文件 ，并导入命名空间：
#include <cooperative_groups.h>
using namespace cooperative_groups;可以用函数 tiled_partition 将一个线程块划分为若干片(tile)，每一片构成一个 新的线程组。目前仅仅可以将片的大小设置为 2 的正整数次方且不大于 32，也就是 2、 4、8、16 和 32 。例如，如下语句通过函 数 tiled_partition 将一个线程块分割为我们熟知的线程束:
thread_group g32 = tiled_partition(this_thread_block(), 32);同时线程块片类型中也有洗牌函数，可以利用线程块片来进行数组归约的计算。
real y = s_y[tid];
thread_block_tile<32> g = tiled_partition<32>(this_thread_block());
for (int i = g.size() >> 1; i > 0; i >>= 1)
{
y += g.shfl_down(y, i);
}这3种方式的性能如下所示，性能相比之前的方式都有所提高：
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 29.49% 190.45ms 100 1.9045ms 1.8181ms 2.0072ms reduce_syncwarp(float const *, float*, int)
27.84% 179.82ms 100 1.7982ms 1.7960ms 1.8183ms reduce_shfl(float const *, float*, int)
27.82% 179.65ms 100 1.7965ms 1.7957ms 1.7976ms reduce_cp(float const *, float*, int)
14.80% 95.571ms 301 317.51us 1.4390us 95.122ms [CUDA memcpy HtoD]
0.06% 384.69us 300 1.2820us 1.2150us 1.7920us [CUDA memcpy DtoH]2.4.3 进一步分析和优化
在前边的例子中， 我们都使用大小为 128 的线程块，所以当 offset 等于 64 时，只用了 1/2 的线程进行计算，其余线程闲置。当 offset 等于 32 时，只用了 1/4 的线程进行计算，其余线程闲置。最终，当 offset 等于 1 时，只用了 1/128 的线程进行计算，其余线程闲置。归约过程一共用了 log2 128 = 7 步， 故归约过程中线程的平均利用率只有 (1/2 + 1/4 + ...)/7 ≈ 1/7 。
为了提高效率，可以考虑在归约之前将多个全局内存数组的数据累加到一个共享内存数组的一个元素中。用一个寄存器变量 y，用来在循环 体中对读取的全局内存数据进行累加， 在规约之前，必须将寄存器中的数据复制到共享内存。
real y = 0.0;
const int stride = blockDim.x * gridDim.x;
for (int n = bid * blockDim.x + tid; n < N; n += stride)
{
y += d_x[n];
}
s_y[tid] = y;
__syncthreads();另外还需要一个调用该核函数的包装函数，以返回最终的计算结果。这里，将GRID_SIZE取为10240，将BLOCK_SIZE取为128。在第10行，调用核函数将长一些的数 组 d_x 归约到短一些的数组 d_y 时，我们使用执行配置 <<<GRID_SIZE, BLOCK_SIZE>>>。当 数据量为 N = 100000000 时，在归约前每个线程将先累加几十个数据。再次调用同一 个核函数将数组 d_y 归约到最终结果(我们就将它保存在 d_y[0])时，仅使用一个线程块，但将线程块大小设置为所允许的最大值，即 1024。
real reduce(const real *d_x)
{
const int ymem = sizeof(real) * GRID_SIZE;
const int smem = sizeof(real) * BLOCK_SIZE;
real h_y[1] = {0};
real *d_y;
CHECK(cudaMalloc(&d_y, ymem));
reduce_cp<<<GRID_SIZE, BLOCK_SIZE, smem>>>(d_x, d_y, N);
reduce_cp<<<1, 1024, sizeof(real) * 1024>>>(d_y, d_y, GRID_SIZE);
CHECK(cudaMemcpy(h_y, d_y, sizeof(real), cudaMemcpyDeviceToHost));
CHECK(cudaFree(d_y));
return h_y[0];
}完整计算时间只需要0.85ms, GPU上核函数的计算时间更是极大缩减。
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 62.32% 95.388ms 1 95.388ms 95.388ms 95.388ms [CUDA memcpy HtoD]
37.59% 57.529ms 200 287.64us 6.4000us 572.79us reduce_cp(float const *, float*, int)
0.09% 135.46us 100 1.3540us 1.3110us 2.5600us [CUDA memcpy DtoH]在上面的包装函数 reduce 中，需要为数组 d_y 分配与释放设备内存。实际上，设备内存的分配与释放是比较耗时的。一种优化方案是使用静态全局内存代替这里的动态全局内存，因为静态内存是编译期间就会分配好的，不会在运行程序时反复地分配，故比动态内存分配高效很多。
利用函数 cudaGetSymbolAddress 将该指针与静态全局变量 static_y 联系起来，更新 reduce 包装函数如下：
__device__ real static_y[GRID_SIZE];
real reduce(const real *d_x)
{
real *d_y;
CHECK(cudaGetSymbolAddress((void**)&d_y, static_y));
const int smem = sizeof(real) * BLOCK_SIZE;
reduce_cp<<<GRID_SIZE, BLOCK_SIZE, smem>>>(d_x, d_y, N);
reduce_cp<<<1, 1024, sizeof(real) * 1024>>>(d_y, d_y, GRID_SIZE);
real h_y[1] = {0};
CHECK(cudaMemcpy(h_y, d_y, sizeof(real), cudaMemcpyDeviceToHost));
// CHECK(cudaMemcpyFromSymbol(h_y, static_y, sizeof(real)); // also ok
return h_y[0];
}计算时间从 0.85ms 缩短到了 0.6 ms。
总结一下以上各种方法的结果及性能：
#轻量级模型设计/部署
这里在一些经典的轻量级网络论文阅读理解的基础上，对不同硬件平台部署轻量级模型相关的经验进行了总结。
轻量级网络的核心是在尽量保持精度的前提下，从体积和速度两方面对网络进行轻量化改造。关于如何手动设计轻量级网络的研究，目前还没有广泛通用的准则，只有一些指导思想，和针对不同芯片平台（不同芯片架构）的一些设计总结，建议大家从经典论文中吸取指导思想和建议，然后自己实际做各个硬件平台的部署和模型性能测试。
一些关键字定义及理解
计算量 FLOPs
FLOPs：floating point operations指的是浮点运算次数，理解为计算量，可以用来衡量算法/模型时间的复杂度。FLOPS：（全部大写），Floating-point Operations Per Second，每秒所执行的浮点运算次数，理解为计算速度, 是一个衡量硬件性能/模型速度的指标，即一个芯片的算力。MACCs：multiply-accumulate operations，乘-加操作次数，MACCs大约是FLOPs的一半。将 w[0]∗x[0]+... 视为一个乘法累加或1个MACC。
内存访问代价 MAC
MAC: Memory Access Cost 内存访问代价。指的是输入单个样本（一张图像），模型/卷积层完成一次前向传播所发生的内存交换总量，即模型的空间复杂度，单位是 Byte。
FLOPs 和 MAC 的计算方式，请参考我之前写的文章神经网络模型复杂度分析：https://github.com/HarleysZhang/cv_note/blob/79740428b6162630eb80ed3d39052cac52f60c32/9-model_deploy/B-%E7%A5%9E%E7%BB%8F%E7%BD%91%E7%BB%9C%E6%A8%A1%E5%9E%8B%E5%A4%8D%E6%9D%82%E5%BA%A6%E5%88%86%E6%9E%90.md。
GPU 内存带宽
GPU的内存带宽决定了它将数据从内存 (vRAM) 移动到计算核心的速度，是比GPU内存速度更具代表性的指标。GPU的内存带宽的值取决于内存和计算核心之间的数据传输速度，以及这两个部分之间总线中单独并行链路的数量。
NVIDIA RTX A4000 建立在 NVIDIA Ampere 架构之上，其芯片规格如下所示:
英伟达 RTX A4000 芯片规格
A4000 芯片配备 16 GB 的 GDDR6 显存、256 位显存接口（GPU 和 VRAM 之间总线上的独立链路数量），因为这些与显存相关的特性，所以A4000 内存带宽可以达到 448 GB/s。
Latency and Throughput
参考英伟达-Ashu RegeDirector of Developer Technology 的
ppt文档 An Introduction to Modern GPU Architecture：https://download.nvidia.com/developer/cuda/seminar/TDCI_Arch.pdf。
深度学习领域延迟 Latency 和吞吐量 Throughput的一般解释：
- 延迟 (
Latency): 人和机器做决策或采取行动时都需要反应时间。延迟是指提出请求与收到反应之间经过的时间。大部分人性化软件系统（不只是 AI 系统），延迟都是以毫秒来计量的。 - 吞吐量 (
Throughput): 在给定创建或部署的深度学习网络规模的情况下，可以传递多少推断结果。简单理解就是在一个时间单元（如：一秒）内网络能处理的最大输入样例数。
CPU 是低延迟低吞吐量处理器；GPU 是高延迟高吞吐量处理器。
Volatile GPU Util
一般，很多人通过 nvidia-smi 命令查看 Volatile GPU Util 数据来得出 GPU 利用率，但是！关于这个利用率(GPU Util)，容易产生两个误区：
- 误区一:
GPU的利用率 =GPU内计算单元干活的比例。利用率越高，算力就必然发挥得越充分。 - 误区二: 同条件下，利用率越高，耗时一定越短。
但实际上，GPU Util 的本质只是反应了，在采样时间段内，一个或多个内核（kernel）在 GPU 上执行的时间百分比，采样时间段取值 1/6s~1s。
原文为 Percent of time over the past sample period during which one or more kernels was executing on the GPU. The sample period may be between 1 second and 1/6 second depending on the product. 来源文档 nvidia-smi.txt（https://developer.download.nvidia.com/compute/DCGM/docs/nvidia-smi-367.38.pdf)
通俗来讲，就是，在一段时间范围内， GPU 内核运行的时间占总时间的比例。比如 GPU Util 是 69%，时间段是 1s，那么在过去的 1s 中内，GPU 内核运行的时间是 0.69s。如果 GPU Util 是 0%，则说明 GPU 没有被使用，处于空闲中。
也就是说它并没有告诉我们使用了多少个 SM 做计算，或者程序有多“忙”，或者内存使用方式是什么样的，简而言之即不能体现出算力的发挥情况。
GPU Util的本质参考知乎文章-教你如何继续压榨GPU的算力(https://zhuanlan.zhihu.com/p/346389176) 和 stackoverflow 问答(https://link.zhihu.com/?target=https%3A//stackoverflow.com/questions/40937894/nvidia-smi-volatile-gpu-utilization-explanation)。
英伟达 GPU 架构
GPU 设计了更多的晶体管（transistors）用于数据处理（data process）而不是数据缓冲（data caching）和流控（flow control），因此 GPU 很适合做高度并行计算（highly parallel computations）。同时，GPU 提供比 CPU 更高的指令吞吐量和内存带宽（instruction throughput and memory bandwidth）。
CPU 和 GPU 的直观对比图如下所示
图片来源 CUDA C++ Programming Guide(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html)
最后简单总结下英伟达 GPU 架构的一些特点:
SIMT(Single Instruction Multiple Threads) 模式，即多个Core同一时刻只能执行同样的指令。虽然看起来与现代CPU的SIMD（单指令多数据）有些相似，但实际上有着根本差别。- 更适合计算密集与数据并行的程序，原因是缺少
Cache和Control。
2008-2020 英伟达 GPU 架构进化史如下图所示:
2008-2020 英伟达 GPU 架构进化史
另外，英伟达 GPU 架构从 2010 年开始到 2020 年这十年间的架构演进历史概述，可以参考知乎的文章-英伟达GPU架构演进近十年，从费米到安培(https://zhuanlan.zhihu.com/p/413145211)。
GPU 架构的深入理解可以参考博客园的文章-深入GPU硬件架构及运行机制()。
CNN 架构的理解
在一定的程度上，网络越深越宽，性能越好。宽度，即通道(channel)的数量，网络深度，及 layer 的层数，如 resnet18 有 18 层网络。注意我们这里说的和宽度学习一类的模型没有关系，而是特指深度卷积神经网络的（通道）宽度。
- 网络深度的意义：
CNN的网络层能够对输入图像数据进行逐层抽象，比如第一层学习到了图像边缘特征，第二层学习到了简单形状特征，第三层学习到了目标形状的特征，网络深度增加也提高了模型的抽象能力。 - 网络宽度的意义：网络的宽度（通道数）代表了滤波器（
3维）的数量，滤波器越多，对目标特征的提取能力越强，即让每一层网络学习到更加丰富的特征，比如不同方向、不同频率的纹理特征等。
手动设计高效 CNN 架构建议
一些结论
- 分析模型的推理性能得结合具体的推理平台（常见如：英伟达
GPU、移动端ARMCPU、端侧NPU芯片等）；目前已知影响CNN模型推理性能的因素包括: 算子计算量FLOPs（参数量Params）、卷积block的内存访问代价（访存带宽）、网络并行度等。但相同硬件平台、相同网络架构条件下，FLOPs加速比与推理时间加速比成正比。 - 建议对于轻量级网络设计应该考虑直接
metric（例如速度speed），而不是间接metric（例如FLOPs）。 FLOPs低不等于latency低，尤其是在有加速功能的硬体 (GPU、DSP与TPU)上不成立，得结合具硬件架构具体分析。- 不同网络架构的
CNN模型，即使是FLOPs相同，但其MAC也可能差异巨大。 - 大部分时候，对于
GPU芯片，Depthwise卷积算子实际上是使用了大量的低FLOPs、高数据读写量的操作。这些具有高数据读写量的操作，加上多数时候GPU芯片算力的瓶颈在于访存带宽，使得模型把大量的时间浪费在了从显存中读写数据上，导致GPU的算力没有得到“充分利用”。结论来源知乎文章-FLOPs与模型推理速度：https://zhuanlan.zhihu.com/p/122943688
一些建议
- 在大多数的硬件上，
channel数为16的倍数比较有利高效计算。如海思351x系列芯片，当输入通道为4倍数和输出通道数为16倍数时，时间加速比会近似等于FLOPs加速比，有利于提供NNIE硬件计算利用率。 - 低
channel数的情况下 (如网路的前几层)，在有加速功能的硬件（NPU 芯片）上使用普通convolution通常会比separable convolution有效率。（来源 MobileDets 论文(https://medium.com/ai-blog-tw/mobiledets-flops%E4%B8%8D%E7%AD%89%E6%96%BClatency-%E8%80%83%E9%87%8F%E4%B8%8D%E5%90%8C%E7%A1%AC%E9%AB%94%E7%9A%84%E9%AB%98%E6%95%88%E6%9E%B6%E6%A7%8B-5bfc27d4c2c8） - shufflenetv2 论文(https://arxiv.org/pdf/1807.11164.pdf) 提出的四个高效网络设计的实用指导思想: G1同样大小的通道数可以最小化
MAC、G2-分组数太多的卷积会增加MAC、G3-网络碎片化会降低并行度、G4-逐元素的操作不可忽视。 GPU芯片上 3\times 33\times 3 卷积非常快，其计算密度（理论运算量除以所用时间）可达 1\times 11\times 1 和 5\times 55\times 5 卷积的四倍。（来源 RepVGG 论文(https://zhuanlan.zhihu.com/p/344324470)）- 从解决梯度信息冗余问题入手，提高模型推理效率。比如 CSPNet(https://arxiv.org/pdf/1911.11929.pdf) 网络。
- 从解决
DenseNet的密集连接带来的高内存访问成本和能耗问题入手，如 VoVNet(https://arxiv.org/pdf/1904.09730.pdf) 网络，其由OSA（One-Shot Aggregation，一次聚合）模块组成。
轻量级网络模型部署总结
在阅读和理解经典的轻量级网络 mobilenet 系列、MobileDets、shufflenet 系列、cspnet、vovnet、repvgg 等论文的基础上，做了以下总结：
- 低算力设备-手机移动端
cpu硬件，考虑mobilenetv1(深度可分离卷机架构-低FLOPs)、低FLOPs和 低MAC的shuffletnetv2（channel_shuffle算子在推理框架上可能不支持）。 - 专用
asic硬件设备-npu芯片（地平线x3/x4等、海思3519、安霸cv22等），目标检测问题考虑cspnet网络(减少重复梯度信息)、repvgg（直连架构-部署简单，网络并行度高有利于发挥GPU算力，量化后有掉点风险） 。 - 英伟达
gpu硬件-t4芯片，考虑repvgg网络（类vgg卷积架构-高并行度带来高速度、单路架构省显存/内存）。
MobileNet block (深度可分离卷积 block, depthwise separable convolution block)在有加速功能的硬件（专用硬件设计-NPU 芯片）上比较没有效率。
这个结论在 CSPNet(https://arxiv.org/pdf/1911.11929.pdf) 和 MobileDets(https://arxiv.org/pdf/2004.14525.pdf) 论文中都有提到。
除非芯片厂商做了定制优化来提高深度可分离卷积 block 的计算效率，比如地平线机器人 x3 芯片对深度可分离卷积 block 做了定制优化。
下表是 MobileNetv2 和 ResNet50 在一些常见 NPU 芯片平台上做的性能测试结果。
深度可分离卷积和常规卷积模型在不同NPU芯片平台上的性能测试结果
以上，均是看了轻量级网络论文总结出来的一些不同硬件平台部署轻量级模型的经验，实际结果还需要自己手动运行测试。
参考资料
- An Introduction to Modern GPU Architecture：
https://download.nvidia.com/developer/cuda/seminar/TDCI_Arch.pdf - 轻量级网络论文解析合集：
https://github.com/HarleysZhang/cv_note/tree/79740428b6162630eb80ed3d39052cac52f60c32/7-model_compression/%E8%BD%BB%E9%87%8F%E7%BA%A7%E7%BD%91%E7%BB%9C%E8%AE%BA%E6%96%87%E8%A7%A3%E6%9E%90
