ysh329 commented 3 years ago

MindSpore Lite OpenCL AutoTune策略分析

mindspore的LiteKernel父类，会被相应设备继承，如OpenCL后端就会有子类class OpenCLKernel继承LiteKernel。这点确保了不会对不同设备后端侵入公共父类，造成代码污染，也让MindSpore的代码看着非常干净，class OpenCLKernel作为所有具体OpenCL Kernel类的公共父类，里面有一个Tune方法，各个具体的Op如Conv2DOpenCLKernel在实现时候可以不实现该Tune方法也可以实现，在调优时即op->kernel->Tune()。

1. 调优Tune思路和流程

该OpenCLKernel父类的实现位于lite/src/runtime/kernel/opencl/opencl_kernel.cc，下面我们从Tune()方法开始，看看其流程和思路。

下面是Tune()方法的实现，可以看到首先进去会判断是否开启了Profiling，因为调优需要根据不同参数设置下的OpenCL的Kernel执行时间，获取该时间就需要在创建OpenCL Command Queue时对属性设置加上Profiling的Flag，且对要获取时间的Kernel在入队时设置event，才能获取GPU的计算时间。

紧接着是TuneMode的判断，根据查询MindSpore Lite的代码，发现其有3种TunningMode：

TuningMode::DEFAULT：不做调优，直接返回；
TuningMode::FAST：只对对FAST_MODE_OPS做调优，目前看到这是一个包含了3个OP的列表：DepthwiseConv2D、Conv2D、DeConv2D。在流程上，当进入该方法，会先判断是否是该模式，同时判断当前OP是否在此列表内的OP，若不在就直接返回。看来这也是FAST的精髓所在，同时后文也没发现模式上有什么区别；
TuningMode::EXTREME：对模型中所有OpenCL OP调优。

后面就是执行流程，具体看代码和我加入的注释：

int OpenCLKernel::Tune() {
  // 判断command queue是否开启Profiling
  if (!ocl_runtime_->isProfiling()) {
    MS_LOG(WARNING) << "Tuning mode require opencl runtime profiling.";
    return RET_OK;
  }

  // 获取当前的TunningMode
  // 判断是否是DEFAULT模式若是则直接返回
  // 判断是否是FAST模式若是则只对Conv/DwConv/DeConv调优其他OP则直接返回
  lite::opencl::TuningMode mode = ocl_runtime_->GetTuningMode();
  if (mode == lite::opencl::TuningMode::DEFAULT) {
    return RET_OK;
  }
  static const std::set<int> FAST_MODE_OPS = {schema::PrimitiveType_Conv2D, schema::PrimitiveType_DepthwiseConv2D,
                                              schema::PrimitiveType_DeConv2D};
  if (mode == lite::opencl::TuningMode::FAST && FAST_MODE_OPS.find(op_parameter_->type_) == FAST_MODE_OPS.end()) {
    return RET_OK;
  }

  // 生成候选调优参数即LocalWorkSize
  auto tuning_params = GenerateTuningParam();
  if (tuning_params.empty()) {
    MS_LOG(WARNING) << "Tuning param size is 0.";
    return RET_OK;
  }

  // 依次对生成的候选参数执行并获取时间
  // 目前看来每种情况只会执行1次即执行1次Run()方法
  // 记录当前执行时间并刷新最短时间的候选参数索引
  int index = -1;
  double min_time = MAX_PROFILING_TIME_MILLI_SECOND;
  for (int i = 0; i < tuning_params.size(); i++) {
    AssignTuningParam(tuning_params[i]);
    auto ret = Run();
    if (ret != RET_OK) {
      MS_LOG(ERROR) << "Tuning " << name() << " failed for tuning param " << tuning_params[i];
      return ret;
    }
    double current_time = GetProfilingTimeMs();
    MS_LOG(DEBUG) << "Tuning " << name() << " param (" << tuning_params[i] << ") exectime " << current_time << "ms";
    if (current_time < min_time) {
      min_time = current_time;
      index = i;
    }
  }

  // 保存最短时间的LocalSize候选参数作为默认参数
  if (index != -1) {
    MS_LOG(INFO) << "Tuning " << name() << " result: param (" << tuning_params[index] << ") exectime " << min_time
                 << "ms";
    AssignTuningParam(tuning_params[index]);
  } else {
    MS_LOG(WARNING) << "Cannot find suitable param.";
  }
  return RET_OK;
}

ysh329 commented 3 years ago

2. 生成候选参数的流程

前文，分析了整体流程，下面我们看下具体对LocalWorkSize的生成，这部分在GenerateTunningParam方法里，该方法在OpenCLKernel基类里就有实现的，作为子类如Conv2D等对其进行了重写。在通用性上，基类里的方法更具有普适性。

下面我们先来看看通用的生成策略，再看针对Conv2D的生成策略。

2.1 通用的OpenCLKernel::GenerateTuningParam

这部分的实现可以在：mindspore/lite/src/runtime/kernel/opencl/opencl_kernel.cc#L231找到，在生成GenerateTuningParam过程中，生成过程完全依赖人工设定的规则和GlobalWorkSize：

人工规则结合了暴力搜索和些许的经验，这部分在后文GenerateLocalByGlobal有写到；
GlobalWorkSize的设定，则是根据任务量在GPU上的分配，这个和OpenCL kernel的实现有直接关系，这里不做展开。

本文更关注生成调优参数的策略和流程，下面是通用的OpenCLKernel父类在GenerateTuningParam方法的实现代码：

std::vector<BaseTuningParameter> OpenCLKernel::GenerateTuningParam() {
  size_t ndim = global_size_.size();
  std::vector<BaseTuningParameter> tuning_params = {};
  if (ndim == 0) {
    MS_LOG(ERROR) << "Generate tuning param failed, global_size_ is null.";
    return tuning_params;
  }
  BaseTuningParameter default_tuning_param = BaseTuningParameter();
  default_tuning_param.local_size = local_size_;
  tuning_params.push_back(default_tuning_param);
  std::vector<size_t> max_work_items = ocl_runtime_->GetWorkItemSize();
  size_t max_workgroup_size = ocl_runtime_->GetMaxWorkGroupSize(kernel_);
  const size_t MIN_WORKGROUP_SIZE = 8;
  std::set<size_t> candidate_x = GenerateLocalByGlobal(global_size_[0]);
  std::set<size_t> candidate_y = {1};
  std::set<size_t> candidate_z = {1};
  if (ndim > 1) {
    candidate_y = GenerateLocalByGlobal(global_size_[1]);
  }
  if (ndim > 2) {
    candidate_z = GenerateLocalByGlobal(global_size_[2]);
  }
  for (auto x : candidate_x) {
    if (x <= max_work_items[0]) {
      for (auto y : candidate_y) {
        if (y <= max_work_items[1]) {
          for (auto z : candidate_z) {
            auto group_size = x * y * z;
            if (z <= max_work_items[2] && group_size <= max_workgroup_size && group_size >= MIN_WORKGROUP_SIZE) {
              BaseTuningParameter tuning_param = BaseTuningParameter();
              tuning_param.local_size = {x, y, z};
              tuning_params.push_back(tuning_param);
            }
          }
        }
      }
    }
  }
  return tuning_params;
}

在生成候选LocalSize过程中，不做调优的默认local_size_作为第一个候选，之后便会根据人工设定的经验自动生成。在生成前，会获取当前设备的硬件参数，作为生成时候的规则限制：

硬件参数1：在3个方向上的最大支持work-item个数，作为生成的单独每个LocalSize（local_x、local_y、local_z）的上界。上界的三个值一般是相同的，如Adreno540上均为1024，目前我所见到的GPU的该值都是1024；
硬件参数2：此kernel的MAX_WORK_GROUP_SIZE，作为生成的LocalSize（local_x*local_y*local_z）连乘积的上界。

此外，也有人工经验上的设定，设置了MIN_WORKGROUP_SIZE=8，作为LocalSize（local_x*local_y*local_z）连乘积的下界。

生成成对的<local_x、local_y、local_z>前，会先生成单独的local_x、local_y、local_z候选值，然后对其排列组合为成对的最终候选结果对<local_x、local_y、local_z>。下面对上述的代码整理，梳理为如下流程：

生成单独的local_x、local_y、local_z候选值：通过调用GenerateLocalByGlobal分别得到三个方向的候选值（后面有详细说明该生成过程）；
3个for循环依次遍历步骤1中的候选值，生成成对的结果，该过程会根据前面所提到的硬件2个限制（3方向的work-item上界、KERNEL_MAX_WORKGROUP_SIZE下界），人工设定的KERNEL_MIN_WORKGROUP_SIZE下界，对LocalSize单独的三个方向以及连乘积加以限制，满足条件则生成候选对<local_x、local_y、local_z>。

2.2 针对Conv2D的GenerateTunningParam

下面讲一下Conv2DOpenCLKernel::GenerateTuningParam，该方法位于mindspore/lite/src/runtime/kernel/opencl/kernel/conv2d.cc，是对父类的同名方法进行了重写，差别仅仅在不需要对local_c调优，该位置的local_c复用默认的local_size_[2]。这里就不贴代码了。

2.3 GenerateLocalByGlobal 分析

GenerateLocalByGlobal是生成LocalSize单独一个方向候选值的重要步骤，会根据Global的设置生成local的候选值，代码如下：

std::set<size_t> OpenCLKernel::GenerateLocalByGlobal(size_t global_i) {
  std::set<size_t> local_ = {};
  int index = 1;
  while (index <= global_i) {
    local_.insert(index);
    index *= 2;
  }
  for (size_t i = 1; i <= 16; i++) {
    if (global_i % i == 0) {
      local_.insert(i);
    }
  }
  return local_;
}

为了更具体的说明生成规则，举例如输入的global_i=32，则GenerateLocalByGlobal在生成候选std::set<size_t> local_时，会有两种策略：

从1开始每次左移1位（乘2），作为候选local值，直到大于global_i时，即第一个策略生成的std::set<size_t> local_={1, 2, 4, 8, 16}；
从1开始到16每次自增1，用global_i%i == 0则作为候选local值，即第二个策略生成的std::set<size_t> local_={1,2,4,8,16}；

因为是集合容器即set，所以两次有交叠，最后global_i=32得到的候选local_={1,2,4,8,16}，因为这个例子中global_i是偶数，所以无论是策略2在对i取余结果为0的结果候选较多。

但若是奇数如global_i=33，则策略1生成的候选是不变的，仍旧是local_={1,2,,4,8,16}，但策略2由于是奇数生成的就很有限local_= {1,3,11}。

这两种从Global Size生成候选Local Size的方式，特点可以简单总结如下：

策略1：偏向默认的（必选）候选，即在[1, global_i]的范围内，下一个候选为上一个值的2倍数，选取范围大。确保了候选值必定存在，且确保local work的值与2的幂次对齐以保证效率；
策略2：偏向较小的因子（取余的被除数范围是[1,16]，选取范围小），当能被整除，将被除数作为候选，虽然选取范围小，但是在Global Size为奇数时往往能获取与策略1不同的奇数候选值，以囊括特殊的边角情况。虽然在策略2上也有必选，如1必然在策略2的候选集合中出现，但是策略1已经包括。

ysh329 commented 3 years ago

3. 性能变化

下面分别在麒麟820和骁龙855上以armv7平台上，选取了2个模型mobilenetv1和v2在两个框架上情形，即总共4个模型。并做了2个TuningMode的测试：default和extreme，测试过程中二者CPU都绑定大核，且设置单线程，确保性能稳定。

kirin820/armv7/bigcore/st

根据kirin820上的两个调优模式的增益来看（第三行：performance%），arm mali gpu对local work size的调优，可能并不敏感：较大的提升仅6.5%，且测试过程中发现存在调优性能和不调优性能，都存在一定波动，其它的提升均较小。

TuningMode\Model	caffe_mobilenetv1	caffe_mobilenetv2	tf_mobilenetv1	tf_mobilenetv2
DEFAULT	15.3 ms	18.4 ms	12.8 ms	13.6 ms
EXTREME	14.3 ms	18.2 ms	12.4 ms	13.5 ms
performance %	6.5%	1.0%	3.1%	0.7%

sd855/armv7/bigcore/st

高通Adreno GPU这边调优后的性能，平均在10%左右

TuningMode\Model	caffe_mobilenetv1	caffe_mobilenetv2	tf_mobilenetv1	tf_mobilenetv2
DEFAULT	12.5 ms	14.0 ms	10.6 ms	9.7 ms
EXTREME	11.3 ms	12.5 ms	9.6 ms	8.6 ms
performance %	9.6%	10.7%	9.4%	11.3%

根据上述测试结果，总结：

arm mali gpu在对local work调优有效，但提升有限，大多在5%内。本身arm mali的gpu性能波动就较大，无论是否调优的情况；
高端adreno gpu在对local work调优有效且提升明显，10%左右。性能更稳定；
低端adreno gpu在对local work调优非常明显，骁龙625 gpu，caffe_mobilenetv1调优前250ms，调优后170ms。

ysh329 / OpenCL-101

【竞品调研】MindSpore Lite的OpenCL AutoTune 策略 #48

MindSpore Lite OpenCL AutoTune策略分析

1. 调优Tune思路和流程

2. 生成候选参数的流程

2.1 通用的OpenCLKernel::GenerateTuningParam

2.2 针对Conv2D的GenerateTunningParam

2.3 GenerateLocalByGlobal 分析

3. 性能变化

kirin820/armv7/bigcore/st

sd855/armv7/bigcore/st