开一个做网站的工作室,北京营销网站建设公司,百度指数的搜索指数代表什么,wordpress更改上传前几天在知乎上看到有知友提问#xff0c;什么是 GPU 算力。当时简单回答了一下#xff0c;今天有空#xff0c;在这里详细谈谈算力。算力也是做高性能计算的核心概念和指标。设备算力分为两部分#xff0c;其一是设备#xff0c;其二是算力。设备主要是指 CPU、GPU、DSP、…前几天在知乎上看到有知友提问什么是 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 中包含有 FP16FP32 等核心用于计算。不同型号的 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。时钟频率有三种可定制分别为650MHz800MHz 和 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 pipelineFMA 的 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 pipelineFMA 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访存数量为 8Mx4Bytes32M计算强度为 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)还比较低。至此我们详细的介绍了硬件算力的概念以及硬件峰值算力和硬件实际算力之间的差距并简单介绍了该如何通过正确的理论来知道我们对性能进行优化。感谢阅读欢迎斧正。
