ysh329
ysh329
paddledetection【TODO】 背景：发现profiler里tf和lite卷积个数不一致，反馈给青青方面，提供新的ssd mobilenetv3模型，比较lite与tflite在armv8 cpu单线程上的性能； 状态：ssd-mobilenetv3青青方面@于广华 提供新模型，还在查模型paddle与tf的差异
### 【MEG-地图-RONIN模型】（已上线） - 背景：在客户端加载过程中遇到了crash的情况，arm-cpu部署； - 需求：解决部署问题； - 解决bug1：缺少arm cpu cast的一种情况，dev已merge，release2.6已merge； - 解决bug2：业务方反馈调用时，试用包armv8，在RuntimeProgram挂掉，但demo无法复现应该是版本与RD测试不一致； - 问题讨论1：加载非法文件卡主，如非模型文件，目前在armv8发生未解密模型文件，predictor加载卡死问题：1.加解密和mml高度耦合，2.本身校验有一定性能损失对已有框架优化所有业务带来加载性能增加，3.开发周期长，建议交给mml来做预先加密模型文件的判断； - 收益：gps信号弱场景下 地图定位精度上升5% - owner：panyu
## CLTune 作者在CLBlast的文章里并没有谈及较为细致的tune说明，而在这篇CLTune，作者在实验部分以矩阵乘法和二维卷积为例，讲了自己CLTune的工作，在2D卷积和GEMM的实验结果上，都达到甚至超过现今最好性能，实验在NVIDIA/AMD/Intel GPU上进行，且统一为FP32精度。  **图：需要auto-tuner的理由** 作者在2016年NVIDIA举办的GPU技术大会上也做了主题为**《Better Than All the Rest: Finding Maximum-Performance GPU Kernels Using Auto-Tuning》**的演讲。下面的内容将结合作者的文章摘要、演讲Slide以及我浅薄的理解。 CLTune与现有很多Tuner工具不同的地方，在其对Kernel的普适性tune支持、易用、支持多种搜索策略（随机/粒子群PSO/模拟退火）且开源。多种搜索策略，尤其是启发式算法，也是由于搜索空间太大而不得不选择的，例如GEMM的搜索空间就达到20万种组合。说到这里，不得不提一下AutoTune的使用场景，也是作者在设计之初考虑的： 1. 搜索空间巨大。如向量宽度（vector width），workgroup，线程的工作粒度(work per thread)，是否使用local memory（cuda里称为shared memory），就光是目前说到的这几个参数，组合起来的情况是巨大的； 2. 各种设备上都要保证高效运行。如GPU厂商对自家GPU优化确保性能符合预期需考虑不同架构代数，手机厂商需要支持发布的多款手机需考虑不同架构及驱动，APP应用开发者需考虑兼容性和不同设备上的性能； 3. 最优参数的设置随输入数据改变的情况。kernel要能以不变应万变，不变的是性能一直保持领先如硬件峰值百分之XX以上，而输入的数据是会变化的。当这种输入数据改变时，也需要调整最优的参数设置； 4....
## 1. 模板化实现Kernel 为有一个直观的描述，下面从一个简单例子`copy`这个io密集型的kernel开始。 ### 举例1：copy 在kernel实现之初，就以类似**模板kernel**的形式来实现，其**模板参数**`WPT`（work per thread）单位线程的工作量，表示了每个线程做多少个元素的拷贝，该参数可以是`{1,2,4}`等，由于对任务数量做了切分，因而主机端设置全局线程数量时，就需要对原始的`GlobalSize`除以`WPT`单位线程的工作量，得到实际需要的全局线程总数。 ```cpp __kernel void copy(__global float* in, __global float* out) { const int tid = get_global_id(0); for (int w = 0; w...
### 举例2：matvec_tiled 第二个例子是矩阵mat_a（M行N列）与向量vec_x（N列1行）的乘操作，且结果为向量vec_y（M行1列），属于BLAS level2 routine。其kernel模板化的实现中有一个可调优的参数TS，即tile size，该参数用于对vec_x的分片缓存，即预先放到local memory中，再在与矩阵mat_a计算的时warp内的work item就能共享使用。vec_x的维度是tile_size的整数倍，且vec_x能被分成N/TS个tile size，cl kernel执行一次（即一个线程执），其内部内的for循环会将所有N/TS个tile size的第一个元素保存到对应local mem的tile_x中，供一个warp内的work item共享。  **图：matvec_tiled kernel的host端和kernel代码** 可以看到这个mat_tiled的外部global work size用到了1维，gid(0)遍历mat_a的M行即[0,M-1)，用到local work size的1维，lid(0)遍历分片（tiled size）的尺寸即[0,TS-1)，遍历的目的是对vec_x做部分的缓存（local memory），缓存大小即分片大小TS行1列，由于是local memory，因而这些是一个warp内所有work item所共享的，且需要在填充完设置barrier(CLK_LOCAL_MEM_FENCE)，用于后续计算vec_y元素的部分结果，即mat_a的1行TS列，与tile_x的TS行1列的尺寸相对应。  **图：matvec_tiled kernel计算示意图** 将上面的代码画成了示意图： 左侧代码，蓝色的表示对vec_x做local...
## 2. 搜索过程中的主要耗时 但当用户量达到一定规模时，这种方式也可以，但数量太过庞大，可能需要on-line在线方式调优，这就要考虑GPU可用性和兼容性，也要考虑到在线调优的时长。 像上面以宏参数的形式传入调优的各种值，是比较好的，但是每次需要编译，在手机上编译一次入mobilenetv1模型，如骁龙8系列的BuildProgram就要100ms这个数量级，模型更大的情况下OpenCL的Program Build如Yolov3模型则500ms到1秒之间，这还是复用了编译过的Program的情况。 所以在移动端上做on-line tune，可能就需考虑避免二次Build Program的调优，可以尝试将原本的宏参数改为`setKernelArg`，以参数的形式来做如在较小粒度上调优是否使用某种inline的方法的哪一种实现（当然这种方式在使用上不如加`build_options`来的方便，毕竟kernel代码里一堆`if-else`的也影响性能），或是调优不需要二次编译Program的local work size（即work group size），还有可以在更大粒度上调优选择要执行同一个Op的kernel的多种实现如卷积的不同实现方式等。 但搜索空间特别庞大时，即使是离线，考虑调优的时间包括：  **图：Profiling flags** cl_profiling_info类型（剖析时间信息） | 解释 -- | -- CL_PROFILING_COMMAND_QUEUED | 当主机（host）将由事件（Event）标识的命令（command）排入命令队列（command queue）时。该值为64bit值，当前设备以纳秒为单位的时间计数器，这些信息下同，略 CL_PROFILING_COMMAND_SUBMIT | 主机将带有时间标识的命令从host提交到device相关联的命令队列中。因为考虑到host上可能有多个device，每个device可能都有各自的命令队列。命令队列由设备、上下文创建，上下文由设备创建，API是这么定义的，一般每个设备只有一个命令队列、一个上下文...
## 3. 搜索空间的特点 搜索过程不是基于一堆已有的性能数据和选项做预测最佳设定，即没有性能数据库，而是基于候选的选项如WPT各种候选值、VW各种候选值等在这些设定下，跑出最好的性能。即使如此，也有一些人为的设定限制，但即使在有这些限制下，搜索空间还是很大，如下图是5个参数下，排列组合且去除不合理设定下仍有3424种组合。  **图：直接卷积的实现下的搜索空间** 这其中也能发现一些空间上的规律： 1. 每种参数实际上候选值是有限的：比方指令宽度（VW）往往是1，2，4，而每线程的工作量往往在2到8，work group的大小也是在2的5次幂，3个维度且再算上默认的(0,0,0)就是16种，是否做for循环的展开，是否使用local mem等等； 2. 卷积的搜索空间只有5维，但若是写的更复杂些如达到10个参数即10维度以上，是轻而易举的； 3. 参数离散且非线性：如WPT可以是1，2，4，8，而且对性能来说，从4到8很可能由于寄存器压力从4到8导致性能急剧下降； 4. 参数间的强相关。  **表：在矩阵乘法中7项参数在不同硬件上的最佳选择** 由于非线性（且值非常接近）和布尔变量参数值的存在，基于导数、自动微分、无导数来寻找最优值的三种方法也不适用。因而选择启发式、以及随机搜索的方式。其实随机搜索是最简单的策略，其采样并测试随机的组合情况。其执行效率完全取决于搜索空间的形状，如果高性能排列组合的参数在搜索空间里挨得近，那么搜索（到高性能的参数）自然效率就低。
## 4. 矩阵乘法 作者介绍了两个例子：2D卷积和矩阵乘法。矩阵乘法介绍的更详细一些，这里我展开一下。 矩阵乘法也是计算密集型算子，且作为2D卷积的实现方式之一，在深度学习和机器学习领域被广泛使用。也是大多数BLAS调优库的重点优化对象。矩阵乘法可以表示为`C = α * A^T B + β C`，其中`α`、`β`为常数，`A^T`为转置后的矩阵A，假设矩阵维度是2次幂，且维度是tile size（后续会说道）的整数倍。 在调优参数上，为了尽可能粒度能细一些，实现了[一个高度可调优的版本](https://github.com/CNugteren/CLTune/blob/master/samples/gemm/gemm.opencl)，其中调优参数有14个： ```c++ // Parameters determined by the tuner // 1. MWG : M维度上的Tile-size，如64, 128 // 2....
## 5. 卷积    