前几天在知乎上看到有知友提问，什么是 GPU 算力。当时简单回答了一下，今天有空，在这里详细谈谈算力。算力也是做高性能计算的核心概念和指标。设备算力分为两部分，其一是设备，其二是算力。
设备主要是指 CPU、GPU、DSP、NUP、FPGA 等等可以进行数学运算的硬件设备。算力的字面意思很好的解释了自己，就是指计算能力，展开来讲，算力就是指设备进行某种运算的能力。
常见的运算无外乎加减乘除，在硬件设计中加法减法是同一运算，且占用硬件资源很少，而乘法则需要相对较多的硬件资源，因此一般硬件中都提供专门的乘法计算单元。所以我们绝大多数情况下，所说的算力是指乘加运算的能力。算力的度量单位是 FLOPS/s，做一次加法的算量是 1 FLOP，做一次乘法的算量也是 1 FLOP，因此一次乘加运算的算量就是 2 FLOPS，如果某个硬件 1s 能够完成一次乘法一次加法，那么这个硬件的算力就是 2 FLOPS / 1s = 2 FLOPS/s。
为了简单起见，我们约定把算力的单位写作 FLOPs，标识 FLOPS/s，把算量的单位写作 FLOPS。在举个小例子，一个 M 行 K 列的二维矩阵 A 乘以一个 K 行 N 列的二维矩阵 B，得到 M 行 N 列的二维矩阵 C，这个计算过程的算量是多少呢？我们知道 A 矩阵的一行和 B 矩阵的一列 对应元素相乘然后累加得到 C 中的一个元素，A 的一行有 K 个元素，B 的一列也有 K 个元素，那么计算 C 中的一个元素需要 K 次乘法 K 次加法，因此计算一个元素的算量就是 2K，那么算出 C 整个矩阵的算量就是 MxNx2K，写作 2MNK FLOPS。
以上介绍了如何算力的概念以及如何计算算量来获取算力，可以看出，算量是无关数据类型的，它只考虑运算的次数。那么算力要不要考虑数据类型呢？答案当然是肯定的，因为算力是来衡量硬件能力的，不同长度的数据类型，计算开销肯定是不同的。
所以，在询问某种设备算力的时候，我们需要明确是什么数据类型的算力。比如是 FP32 的算力，还是 FP16 的。此时大家可能已经知道什么是算力了。
那么对于具体的一款硬件设备我们该如获取他的算力呢？
图形 硬件设备的理论算力
算力作为硬件的基础指标，一定是与具体硬件相关的。在高性能计算中，具体算法表现所能达到峰值算力的百分比是衡量硬件算力利用率的核心指标，我们在最后一部分会详细介绍。本节主要介绍如何评估一款具体硬件的理论算力也叫做峰值算力。以常见的 CPU 和 GPU 为例，为了提升设备的 AI 运算能力，较新型号的设备都会提供 FP32、FP16、INT8 等三种乘加硬件单元以提升设备 AI 能力，无论那种数据类型，理论算力的计算方法是一样的，所以本文主要以 FP32 或者 FP16 浮点运算单元为例进行说明。
接下来我们以 Imagination 公司的 PowerVR GPU 为例介绍一下硬件算力的计算方法。PowerVR 的 GPU 作为 Apple 早期的御用 GPU 足见其在 GPU 设计上的优秀能力，A7/A8/A9 处理器都是使用的 PowerVR 的 GPU。oppo、vivo 和小米等国内厂商大量使用过 PowerVR 的 GE8320/8322 和 GM9446 等中低端 GPU。除了在移动端设备上，在汽车上 PowerVR 也有不少的 GPU 出货量。稍微介绍一点 Imagenation 公司的渊源，该公司是一家全球的半导体及软件设计公司，主要业务是设计 PowerVR 系列移动端 GPU，以及 MIPS CPU 相关业务。2017 年该公司由中资背景基金凯桥收购。为了规避风险，仅收购 PowerVR GPU 相关业务，MIPS CPU 相关业务由美国公司收购。近年来随着国产化推进，以及大量的自研 GPU 初创公司，其中相当量的企业就是通过购买 PowerVR 的 GPU 快速推出新品的。
下图是 Imagination 公司的 PowerVR GPU Series7XT系列中的GT7600 的统一渲染簇 (USC, Unified Shading Cluster) 的示意图。Series7XT 属于 Rogue 架构，GT7600 属于 PowerVR 的一款高端 GPU，在 Qualcomm 的 Adreno 和 Arm 的 Mali GPU 在 2-300 GFLOPS 的时期，其算力已经达到接近 1 TFLOPS 算力。当时用在魅族以及美图手机的旗舰机型上，后来虽然因为种种原因这两款设备也没有做起来(毕竟手机能否做好也不靠 GPU 算力)。
ALU 是 USC 中的计算单元。Rogue架构中每个 USC 包含 16 个 ALU Pipelines，每个 ALU 中包含有 FP16，FP32 等核心用于计算。不同型号的 GPU 包含的 ALU core 类型不同，数量也不同。例如在Series7XT Plus GPU 中增加了 Integer Pipelines，能够支持 Int8和 Int16 等整数类型，可以大幅度提升 GPU 性能，在 Series7XT GPU中增加了 FP64。FP16 核心也不是 Rogue 架构的标配，早期型号也是没有的。可以看出 PowerVR GPU 的先进性，在 17 年的 GPU 上就有了 Int8 等整型的加速，而 Mali 和 Adreno 都是在 2020/2021 后才开始陆续增加的。
Series7XT 系列中的 GT7600 包含 6 个 USC，一共有 192 个 FP32 core or 384 个 FP16 core。时钟频率有三种可定制，分别为650MHz，800MHz 和 1GHz。以 1 GHZ FP32 为例，6 x 192 x 1 x 2 = 2304 GFLOP/s, 最低配置的 650MHz 也可以达到 1497 GFLOP/s。
这样就可以得到硬件的理论峰值，其实就是根据计算单元的数量和计算单元的算力再乘以频率就可以得到峰值算力。
对于 CPU 其实也一样，只是 CPU 由于有大量的面积作为缓存和控制电路，因此其算力 pipeline 相对于 GPU 较少(涉及 GPU 和 CPU 分工和设计理念上的差异)。以 Arm 的 A76 为例，其拥有 2 条 128 bit FMA pipeline，FMA 的 throughput 是 2，其 FP32 的峰值算力为 2 x 4 x 2 FLOPS。所以 A76 的算力是 16 x 频率。
以上是关于理论算力的计算方法，那么在实际使用中是否能够达到这一水平呢？这其中涉及程序员的代码编写方式、编译器行为以及硬件资源的发挥水平，接下来我们会介绍一下设备实际算力的情况。
硬件设备的实际算力
关于优化技术的细节介绍，在之前的文章和以后的文章中，笔者都会继续详细介绍，本文仅仅以简单的方式展示不同实现方式所带来的性能差距。
首先给出矩阵乘法的朴素实现，C += AxB，代码部分十分简单这里不再赘述，直接给出代码。
如下:
void AddDot(int K, float *x, int incy, float *y, float *gamma)
{int p;for (p = 0; p < K; p++){*gamma += x[p] * Y(p);}
}void Gemm(int M, int N, int K, float *a, int lda, float *b, int ldb, float *c, int ldc)
{// C += A x Bfor (int m0 = 0; m0 < M; ++m0){for (int n0 = 0; n0 < N; ++n0){AddDot(K, &MA(m0, 0), ldb, &MB(0, n0), &MC(m0, n0));}}
}
之所以将实现分为两部分，也是为了之后的分块实现更方便。大家都知道，CPU 中不同层级存储的带宽差异是巨大的，例如 L1 Cache 的带宽就会远远优于 DDR 的带宽，因此在计算优化中，为了增加 L1 Cache 的命中率，会将矩阵分块，使计算拥有更好的局部性。这也是由矩阵计算的特性决定的(矩阵计算本身有很好的分块特性)。
下面向大家展示一段简单实现的 SIMD 实现的矩阵乘法。该实现只是简单的将矩阵计算在 M 和 N 维度上分为 4x4 大小的块，同时使用 SIMD 向量指令来加速。而这只是真正优化的起点，因此该实现的算力利用率也并不高，关于进一步对矩阵计算的优化将在后续文章中继续介绍。
void AddDot4x4(int k, float *a, int lda, float *b, int ldb, float *c, int ldc)
{float *a_0p_pntr, *a_1p_pntr, *a_2p_pntr, *a_3p_pntr;a_0p_pntr = &MA(0, 0);a_1p_pntr = &MA(1, 0);a_2p_pntr = &MA(2, 0);a_3p_pntr = &MA(3, 0);float32x4_t c_p0_sum = {0};float32x4_t c_p1_sum = {0};float32x4_t c_p2_sum = {0};float32x4_t c_p3_sum = {0};register float a_0p_reg, a_1p_reg, a_2p_reg, a_3p_reg;for (int p = 0; p < k; ++p){float32x4_t b_reg = vld1q_f32(&MB(p, 0));a_0p_reg = *a_0p_pntr++;a_1p_reg = *a_1p_pntr++;a_2p_reg = *a_2p_pntr++;a_3p_reg = *a_3p_pntr++;c_p0_sum = vmlaq_n_f32(c_p0_sum, b_reg, a_0p_reg);c_p1_sum = vmlaq_n_f32(c_p1_sum, b_reg, a_1p_reg);c_p2_sum = vmlaq_n_f32(c_p2_sum, b_reg, a_2p_reg);c_p3_sum = vmlaq_n_f32(c_p3_sum, b_reg, a_3p_reg);}float *c_pntr = 0;c_pntr = &MC(0, 0);float32x4_t c_reg = vld1q_f32(c_pntr);c_reg = vaddq_f32(c_reg, c_p0_sum);vst1q_f32(c_pntr, c_reg);c_pntr = &MC(1, 0);c_reg = vld1q_f32(c_pntr);c_reg = vaddq_f32(c_reg, c_p1_sum);vst1q_f32(c_pntr, c_reg);c_pntr = &MC(2, 0);c_reg = vld1q_f32(c_pntr);c_reg = vaddq_f32(c_reg, c_p2_sum);vst1q_f32(c_pntr, c_reg);c_pntr = &MC(3, 0);c_reg = vld1q_f32(c_pntr);c_reg = vaddq_f32(c_reg, c_p3_sum);vst1q_f32(c_pntr, c_reg);
}void GemmB4x4(int M, int N, int K, float *a, int lda, float *b, int ldb, float *c, int ldc)
{// C = A x Bfor (int m0 = 0; m0 < M; m0 += 4){for (int n0 = 0; n0 < N; n0 += 4){AddDot4x4(K, &MA(m0, 0), lda, &MB(0, n0), ldb, &MC(m0, n0), ldc);}}
}笔者在联发科的天玑1000+ 设备上对两段实现进行了测试，天玑1000+ 是 4 个 A55(2000 MHz), 4 个 A77(2600 MHz) 的架构。程序均运行在小核 A55 上，频率达到 1791 MHz。
性能对比如下
由上图可以看到，虽然优化后的利用率也不高，但是相对于 Native 实现有 2-3 倍的性能提升。优化效果还是很明显的。A55 拥有 2 条 64 bit 的 FMA pipeline，FMA throughput 为 1，按照我们达到的 1791 MHz 的频率算，其峰值计算能力为 1 x 4 x 2 x 1791 / 1024 = 13.99 GFLOPS。可以看到当前的优化，算力利用率还非常的低。因此还需要做大量的工作才能进一步提升其算力利用率。
算力指标和性能优化的关系
程序要在硬件上运行，因此程序运行的各项指标是受硬件条件限制和制约的。单从性能优化角度来讲，程序的性能受硬件的带宽和算力制约，根据程序规模的不同性能会受到从硬盘到 L1 cache 等各级存储的带宽影响。算力的影响则是主要是受 ALU 或者其他超越函数硬件单元算力上限的影响。关于程序性能与算力和带宽的关系，在 HPC 领域很早已经给出了相应的理论 Roofline Model (https://en.wikipedia.org/wiki/Roofline_model)。
Roofline 模型给出了程序性能与硬件带宽和硬件算力的关系，并且能够帮助我们区分性能是受限于内存带宽，还是算力。该理论可以很好的评估我们程序的性能水平以及优化上限。下面我们以前一节介绍的两种 GEMM 实现为例，结合 Roofline 模型分析器性能优化情况及接下来的优化方向。
前一节给出了，两种方案的算力对比，下图给出两种方案的带宽利用率。
可以看到 SIMD 方案的带宽也是有不少提升的，但是还是很低。笔者测试天玑 1000+ 的读写带宽可以达到 12 GB/s.下图为Roofline 模型的示意图。
其中绿色的线表示算力的峰值，单位为 FLOPs or GFLOPs，实验平台为 13.99 GFLOPs. 斜率为带宽，单位为 B/s or GB/s， 实验平台为 12 GB/s。横轴标识操作强度(Operational Intensity, OI)，为 FLOPS / B。落在红色区域表示当前实现为 Memory Bound，需要优化访存，落在绿色区域表示当前实现为 Compute Bound，达到了硬件的算力峰值，无法进一步提升算力了。下一步我们使用该模型分析一下目前两种优化方案(假设矩阵是方阵 M/N/K == M，不影响结论)，首先看 Native 版本，计算一个输出点需要读 2M 个数据进行 M 次乘加，那么计算强度是多少呢？算量为 Mx2FLOPS，访存为 2Mx4Bytes = 8M，所以计算强度 OI = (2M)/(8M) = 0.25, 此时要达到峰值算力 13.99 GFLOPs，对应需要的带宽是 13.99 / 0.25 = 55.96 GB/s,而上图 Native 版本的带宽仅有 0.01 GB/s 左右，说明我们带宽远远不足，所以接下来需要进一步优化访存，提升数据访问的速度。那么 SIMD 版本的情况怎么样呢？SIMD 版本每计算 16 个点需要读取 8M 个数，算量为 16Mx2FLOPs，访存数量为 8Mx4Bytes=32M，计算强度为 1，那么为了达到峰值算力他所需要的带宽为 13.99 / 1 = 13.99 GB/s,而当前 SIMD 版本的带宽为 0.02 GB/s，因此依然远远无法达到峰值性能。但是相比较于 Native 版本，实际带宽提升到0.02，所需带宽降低为 13.99，其实已经有很大优化了，接下来需要做的是进一步优化访存同时增加计算强度，增加计算强度可以有效降低所需带宽的数值，优化访存可以提升实际带宽。同时按照当前的带宽和峰值算力情况，达到峰值所需要的计算强度是 13.99 / 12 = 1.165, 所以目前我们的计算强度已经基本接近理论值，需要进一步优化带宽，以获取更好的性能。由于这里没有考虑 cache 的影响，所以我们实际可用带宽可以更高，在真正达到峰值算力的时候，计算强度可以略有波动。后续的文章中我们将继续探讨如何进一步优化来提升当前 SIMD 版本的性能。当前计算平台的 Roofline 模型折线图如下图所示。
目前的 OI 都是落在红色区域属于Memory Bound，因此我们需要进一步优化访存，以提升程序性能，cache 命中率(包括 L1/L2)还比较低。至此我们详细的介绍了硬件算力的概念以及硬件峰值算力和硬件实际算力之间的差距，并简单介绍了该如何通过正确的理论来知道我们对性能进行优化。感谢阅读，欢迎斧正。
![[C++]MySQL数据库操作实例](http://www.cnblogs.com/Emoticons/msn/wink_smile.gif)