提交 1eb9779f 编写于 作者: T tienfeek

Merge branch 'fpga_pr' into new_dev

# 如何增加Layout
Paddle-Lite中Place包含了Target、Layout、Precision信息,用来注册和选择模型中的具体Kernel。下面以增加Place中的layout:`ImageDefault``ImageFolder``ImageNW`为例,讲解如何增加新Layout。
根据在`lite/core/``lite/api`目录下以`NHWC`为关键词检索代码,发现需要分别在以下的文件中加入Layout内容:
1. lite/api/paddle_place.h
2. lite/api/paddle_place.cc
3. lite/api/python/pybind/pybind.cc
4. lite/core/op_registry.h
5. lite/core/op_registry.cc
## 1. lite/api/paddle_place.h
`enum class DataLayoutType`中加入对应的Layout,注意已有的Layout不能改变值,增加新Layout递增即可:
```cpp
enum class DataLayoutType : int {
kUnk = 0,
kNCHW = 1,
kNHWC = 3,
kImageDefault = 4, // for opencl image2d
kImageFolder = 5, // for opencl image2d
kImageNW = 6, // for opencl image2d
kAny = 2, // any data layout
NUM = 7, // number of fields.
};
```
## 2. lite/api/paddle_place.cc
本文件有3处修改,注意在` DataLayoutToStr`函数中加入对应Layout的字符串名,顺序为`lite/api/paddle_place.h`中枚举值的顺序:
```cpp
// 该文件第1处
const std::string& DataLayoutToStr(DataLayoutType layout) {
static const std::string datalayout2string[] = {
"unk", "NCHW", "any", "NHWC", "ImageDefault", "ImageFolder", "ImageNW"};
auto x = static_cast<int>(layout);
CHECK_LT(x, static_cast<int>(DATALAYOUT(NUM)));
return datalayout2string[x];
}
// 该文件第2处
const std::string& DataLayoutRepr(DataLayoutType layout) {
static const std::string datalayout2string[] = {"kUnk",
"kNCHW",
"kAny",
"kNHWC",
"kImageDefault",
"kImageFolder",
"kImageNW"};
auto x = static_cast<int>(layout);
CHECK_LT(x, static_cast<int>(DATALAYOUT(NUM)));
return datalayout2string[x];
}
// 该文件第3处
std::set<DataLayoutType> ExpandValidLayouts(DataLayoutType layout) {
static const std::set<DataLayoutType> valid_set({DATALAYOUT(kNCHW),
DATALAYOUT(kAny),
DATALAYOUT(kNHWC),
DATALAYOUT(kImageDefault),
DATALAYOUT(kImageFolder),
DATALAYOUT(kImageNW)});
if (layout == DATALAYOUT(kAny)) {
return valid_set;
}
return std::set<DataLayoutType>({layout});
}
```
## 3. lite/api/python/pybind/pybind.cc
```cpp
// DataLayoutType
py::enum_<DataLayoutType>(*m, "DataLayoutType")
.value("NCHW", DataLayoutType::kNCHW)
.value("NHWC", DataLayoutType::kNHWC)
.value("ImageDefault", DataLayoutType::kImageDefault)
.value("ImageFolder", DataLayoutType::kImageFolder)
.value("ImageNW", DataLayoutType::kImageNW)
.value("Any", DataLayoutType::kAny);
```
## 4. lite/core/op_registry.h
找到KernelRegister final中的`using any_kernel_registor_t =`,加入下面修改信息:
```cpp
// 找到KernelRegister final中的`using any_kernel_registor_t =`
// 加入如下内容:
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFP16),
DATALAYOUT(kNCHW)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFP16),
DATALAYOUT(kNHWC)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFP16),
DATALAYOUT(kImageDefault)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFP16),
DATALAYOUT(kImageFolder)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFP16),
DATALAYOUT(kImageNW)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageFolder)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageNW)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kAny),
DATALAYOUT(kImageDefault)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kAny),
DATALAYOUT(kImageFolder)> *, //
KernelRegistryForTarget<TARGET(kOpenCL),
PRECISION(kAny),
DATALAYOUT(kImageNW)> *, //
```
## 5. lite/core/op_registry.cc
该文件有2处修改:
```cpp
// 该文件第1处
#define CREATE_KERNEL1(target__, precision__) \
switch (layout) { \
case DATALAYOUT(kNCHW): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kNCHW)>(op_type); \
case DATALAYOUT(kAny): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kAny)>(op_type); \
case DATALAYOUT(kNHWC): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kNHWC)>(op_type); \
case DATALAYOUT(kImageDefault): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kImageDefault)>(op_type); \
case DATALAYOUT(kImageFolder): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kImageFolder)>(op_type); \
case DATALAYOUT(kImageNW): \
return Create<TARGET(target__), \
PRECISION(precision__), \
DATALAYOUT(kImageNW)>(op_type); \
default: \
LOG(FATAL) << "unsupported kernel layout " << DataLayoutToStr(layout); \
}
// 该文件第2处
// 找到文件中的下面的函数
KernelRegistry::KernelRegistry()
: registries_(static_cast<int>(TARGET(NUM)) *
static_cast<int>(PRECISION(NUM)) *
static_cast<int>(DATALAYOUT(NUM)))
// 在该函数中加入新增Layout的下面内容
INIT_FOR(kOpenCL, kFP16, kNCHW);
INIT_FOR(kOpenCL, kFP16, kNHWC);
INIT_FOR(kOpenCL, kFP16, kImageDefault);
INIT_FOR(kOpenCL, kFP16, kImageFolder);
INIT_FOR(kOpenCL, kFP16, kImageNW);
INIT_FOR(kOpenCL, kFloat, kImageDefault);
INIT_FOR(kOpenCL, kFloat, kImageFolder);
INIT_FOR(kOpenCL, kFloat, kImageNW);
INIT_FOR(kOpenCL, kAny, kImageDefault);
INIT_FOR(kOpenCL, kAny, kImageFolder);
INIT_FOR(kOpenCL, kAny, kImageNW);
```
# 新增Pass方法
本文从三个方面介绍了`Lite`中的`Pass`结构:**Pass是什么****Pass的实现与接口****Pass的一般注册流程**。最后以`Fc_fuse_pass`为例介绍了`fusion_pass`的作用与注册方法。
## 前述:Pass是什么?
**CxxPredictor加载模型后,在执行预测前会先优化模型。模型优化过程是通过Pass实现的。**
具体调用关系如下:
![图片](https://user-images.githubusercontent.com/45189361/69638690-20d21880-1096-11ea-8169-1d2c7e1a1609.png)
- `CreatePredictor(CxxConfig)`函数调用了Predictor->Build(CxxConfig)
- CxxPredictor的构建过程(Build)分为两步:
- Predictor->LoadModel() 加载模型文件到program中
- Predicotr->optimizer_.Run() 对Program中的原始图形结构进行优化
- 对图结构的优化是通过调用 `Pass->Apply(const std::unique_ptr<SSAGraph>& graph)`方法实现的。
**每一类Pass定义了一种优化过程**,包括:原模型中的kernel选取、OP融合、冗余OP去除、子图创建、内存优化、类型推导、类型转换等。
## Pass的实现与接口 :Pass基类、PassManager和Pass注册
### 1、Pass基类:`paddle::lite::mir::Pass`
```c++
class Pass {
public:
// Pass的类型,Pass按照作用的不同可以分为三种
enum class Kind { //种类的作用不太清楚
// 1. 修改模型中的图拓扑结构的Pass
kProgramWise = 0,
// 2. 不修改图结构,修改状态的Pass
kStmtWise,
// 3. 不修改 IR,用于搜集信息和可视化信息的Pass.
kDebug,
};
// 主要实现函数:Apply 函数定义了 Pass 运行时执行的操作
virtual void Apply(const std::unique_ptr<SSAGraph>& graph) = 0;
bool is_program_pass() const { return kind_ == Kind::kProgramWise; }
bool is_stmt_pass() const { return kind_ == Kind::kStmtWise; }
virtual ~Pass() = default;
private:
const Kind kind_; // pass 的种类
std::string name_; // pass 的名称
std::set<TargetType> bound_targets_; // 指定了Pass运行的硬件平台,模型优化过程会根据当前硬件平台是否匹配筛选Pass。
std::unordered_map<std::string, std::set<lite_api::Place>> bound_kernels_; // 绑定的kernel
};
// Different kinds.
class ProgramPass : public Pass {
public:
ProgramPass() : Pass(Kind::kProgramWise) {}
};
class StmtPass : public Pass {
public:
StmtPass() : Pass(Kind::kStmtWise) {}
};
class DebugPass : public Pass {
public:
DebugPass() : Pass(Kind::kDebug) {}
};
```
**代码位置**`lite/core/mir/pass.h`
**主要类成员**
`const Kind kind_` : Pass类型。pass 有三种基本基本类型 :修改图结构的`ProgramPass`、修改状态量的`StmtPass`和Debug过程采集信息与控制可视化的`DebugPass`
`std::string name_` :pass 的名称
`std::set<TargetType> bound_targets_` : Pass运行的硬件平台,optimizer.Run()优化过程会根据硬件平台选择匹配的Pass。------根据硬件平台自动选择需要的pass
`std::unordered_map<std::string, std::set<lite_api::Place>> bound_kernels_` : Pass 绑定的kernel (what's this used for)
**主要接口**
`Pass::Apply(const std::unique_ptr& graph)` : Pass优化过程的具体操作,是新注册Pass需要实现的接口。输入为`SSAGraph`型指针,是对模型结构的拓扑表示。
### 2、Pass管理 `paddle::lite::mir::PassManager`
```c++
class PassManager {
public:
// 内部静态变量PassManager,用来存储使用的Pass和图优化操作
static PassManager& Global() {
static PassManager x;
return x;
}
// 执行所有的 Pass
void Run(const std::unique_ptr<SSAGraph>& graph) {
for (auto& pass : passes_) {
LOG(INFO) << "Running MIR pass " << pass->name();
pass->Apply(graph);
}
private:
std::list<std::unique_ptr> passes_; //存储所有的 Pass
std::map<std::string, mir::Pass*> pass_map_; //使用map变量存储 PassName::Pass
}
```
**代码位置**`lite/core/mir/pass_manager.h`
**主要类成员**
`std::list:unique_ptr> passes_;` : List类型,存储了所有已注册Pass。
`std::map<std::string, mir::Pass*> pass_map_; ` : Map类型,存储了所有"Pass名称-Pass类"键对,用于根据名称查找Pass。
**主要接口**
`static PassManager& Global()` 返回PassManager全局静态变量,该变量存储了所有已注册的Pass
` bool AddNewPass(const std::string& name, Pass* pass)` 添加新的Pass到PassManager中
### 3、 Pass 注册 `paddle::lite::mir::PassRegistry`
**代码位置**`lite/core/mir/pass_registry.h`
**主要接口**
`REGISTER_MIR_PASS(name__, class__)` :宏定义函数,用于注册Pass。注册Pass过程实现的是 `PassManager::Global().AddNewPass(name__, class__)`,将新注册Pass添加到全局变量`PassManager`中。
## Pass的一般注册流程与使用方法
### 1. Pass 注册流程
`lite/core/mir`或其子目录下继承`Pass基类`,实现`Pass::Apply`接口,并使用宏`REGISTER_MIR_PASS(name__, class__)`将Pass注册到`PassManager`即完成了新Pass注册。
**以新建 **`new_demo_pass`**为例**,具体流程如下:
(1)在`lite/core/mir`路径下新建`example_pass.cc``new_demo_pass.h` 文件
(2)在`example_pass.h` 文件中继承Pass基类(ProgramPass、StmtPass或DebugPass)定义自己的Pass类。
```c++
#include "lite/core/mir/pass.h"
namespace paddle {
namespace lite {
namespace mir {
class ExamplePass : public ProgramPass {
void Apply(const std::unique_ptr<SSAGraph> &graph) override {}
...
};
} // namespace mir
} // namespace lite
} // namespace paddle
```
(3)在`example_pass.cc` 文件中实现`ExamplePass::Apply()`接口,并注册`ExamplePass`
```c++
#include "lite/core/mir/pass_registry.h"
#include "lite/core/mir/example_pass.h"
namespace paddle {
namespace lite {
namespace mir {
void ExamplePass::Apply(const std::unique_ptr<SSAGraph>& graph) {
...
}
} // namespace mir
} // namespace lite
} // namespace paddle
REGISTER_MIR_PASS(example_pass, paddle::lite::mir::ExamplePass)
.BindTargets({TARGET(kARM)}); // Pass执行的目标硬件平台
// .BindKernel("conv2d"); //Pass绑定的 kernel
```
(4)修改`lite/core/mir/CMakeLists.txt`文件,将`example_pass.cc` 编译到`mir_passes`库中
```cmake
lite_cc_library(mir_passes
SRCS
demo_pass.cc // 新建的Pass文件
...
memory_optimize_pass.cc
DEPS mir_pass types context ${mir_fusers} ${subgraph_passes})
```
### 2. Pass使用流程
将Pass注册到PassManager后不会自动生效。需要在`optimizer->run()` 函数中添加该Pass才会在模型优化过程中调用。
(1)在`paddle_use_passes.h`文件中调用该Pass
```cmake
#include "paddle_lite_factory_helper.h" // NOLINT
...
USE_MIR_PASS(new_demo_pass); //调用 new_demo_pass
```
(2)要想在优化模型时调用该Pass,需要在`optimizer->run()`函数中手动添加调用。
修改`lite/core/optimizer.h`文件,添加`new_demo_pass``Optimizer::Run()`函数;
```c++
class Optimizer {
public:
void Run(...) {
...
if (passes.empty()) {
RunPasses(std::vector<std::string>{
{"new_demo_pass" //将新注册的Pass添加在这里
...
}
...
}
```
(3)只有CxxPredictor才会在模型加载后根据Pass优化模型。
```c++
...
#include "paddle_use_passes.h" // 引用Pass优化模型
void RunModel() {
// 1. 创建 CxxConfig
CxxConfig config;
config.set_model_dir(FLAGS_model_dir);
config.set_valid_places(Place{TARGET(kARM), PRECISION(kFloat)});
// 2. 创建CxxPredictor,该过程包括加载模型和用Pass优化模型
std::shared_ptr> predictor =
Creat<CxxConfig>(config);
}
```
## Fusion Pass的定义与注册
`Fusion Pass`是一种常见图结构优化Pass,可将多个连续OP融合成单个等效OP,减少数据交换并简化图结构。Pass运行时调用`Fuser`自动查找并替换指定图结构,所以注册`FuserPass`时还需要实现对应的Fuser类。
下面以`fc_fuse_pass`为例,详细说明`FusionPass`的效果和注册方法。
### `fc_fuse_pass`的作用
将相邻的`mul`算子和 `element_wise add `算子 融合成一个 `FC` 算子
```c++
mul(X) = X * W
elementwise_add( mul(x) ) = X * W + Bias
//----------> after fusion
FC(X) = X * W +Bias
```
Pass 运行效果如下:
![图片](https://user-images.githubusercontent.com/45189361/69639193-12383100-1097-11ea-9063-21f030414080.png)
mul和elementwise_add的原有参数映射到FC的参数上:
![图片](https://user-images.githubusercontent.com/45189361/69638836-74446680-1096-11ea-9cdc-a961fa995dfe.png)
### `fc_fuse_pass`的注册方法
#### 1、创建FcFuser
(1)在`lite/core/mir/fusion`路径下新建`fc_fuser.cc``fc_fuser.h` 文件
(2)在`fc_fuser.h` 文件中继承`FuseBase`定义自己的Fuser类。
```c++
#include "lite/core/mir/pattern_matcher_high_api.h"
namespace paddle {
namespace lite {
namespace mir {
namespace fusion {
class FcFuser : public FuseBase {
public:
void BuildPattern() override;
void InsertNewNode(SSAGraph* graph, const key2nodes_t& matched) override;
private:
cpp::OpDesc GenOpDesc(const key2nodes_t& matched) override;
};
} // namespace fusion
} // namespace mir
} // namespace lite
} // namespace paddle
```
**主要接口**
`FuseBase::BuildPattern` : 描述需要替换位置的图结构(pattern),Fuser运行时会自动查找并替换该pattern。
`FuseBase::GenOpDesc` : 创建融合后的等效Fused_op。
`FuseBase::InsertNewNode` :用Fused_op替换原始图结构(pattern)。
对于 `FcFuser`:BuildPattern描述的Pattern是`mul+elementwise add`,GenOpDesc创建的FC_op,InsertNewNode函数的效果是用新建的`FC_op`替换模型中的`mul+elementwise add` pattern。
(3) 在`fc_fuser.cc`文件中实现 `BuildPattern()``GenOpDesc()``InsertNewNode() `接口
下面以FcFuser为例介绍三种接口的实现:
```c++
// 1. BuildPattern函数,描述需要替换的图结构
// FcFuser::BuildPattern() 描述了 mul + element_wise add 图结构
void FcFuser::BuildPattern() {
// (1) 用OpNode描述和VarNode
// mul OP
auto* mul = OpNode("mul", "mul");
// mul OP 的输入和输出
auto* x = VarNode("x")->assert_is_op_input("mul", "X");
auto* W = VarNode("W")->assert_is_op_input("mul", "Y");
auto* mul_out = VarNode("mul_out");
// elementwise_add OP
auto* add = OpNode("add", "elementwise_add");
//elementwise_add 的输入
auto* b = VarNode("b")->assert_is_persistable_var();
// elementwise_add OP的输出(最终输出)
auto* Out = VarNode("Out");
//(2) 描述拓扑连接 (Fuse之前mul 和elementwise_add的连接)
std::vector<PMNode*> mul_inputs{W, x};
std::vector<PMNode*> add_inputs{mul_out, b};
mul_inputs >> *mul >> *mul_out;
add_inputs >> *add >> *Out;
//(3) 声明新的拓扑结构中将会被移除的节点,包括被fuse的OP和OP之间的中间变量
mul_out->AsIntermediate();
mul->AsIntermediate();
add->AsIntermediate();
}
// 2. GenOpDesc函数新建等效 Fused_op
// FcFuser::GenOpDesc() 新建了Fc_op
cpp::OpDesc FcFuser::GenOpDesc(const key2nodes_t& matched) {
// (1) 得到第一个OP节点的 OpDesc ,并清空输入输出信息
cpp::OpDesc op_desc = *matched.at("mul")->stmt()->op_info();
op_desc.mutable_inputs()->clear();
op_desc.mutable_outputs()->clear();
// (2) 修改OpDesc , 将OpType设置为 "fc" (FC OP 的OP_type),
op_desc.SetType("fc");
// (3) 设置OpDesc中的Input、Output、Attrbute。分别连接到BuildPattern()函数中创建的VarNode
op_desc.SetInput("Input", {matched.at("x")->arg()->name});
op_desc.SetInput("W", {matched.at("W")->arg()->name});
op_desc.SetInput("Bias", {matched.at("b")->arg()->name});
op_desc.SetOutput("Out", {matched.at("Out")->arg()->name});
op_desc.SetAttr(
"in_num_col_dims",
matched.at("mul")->stmt()->op_info()->GetAttr<int>("x_num_col_dims"));
return op_desc;
}
// 3. InsertNewNode函数用Fused OP 替换模型图中的原始 Pattern
// FcFuser::InsertNewNode() 用Fc_OP替换原始模型图中的 " mul + element_wise add "
void FcFuser::InsertNewNode(SSAGraph* graph, const key2nodes_t& matched) {
// (1) 创建FC OP的参数(OpDesc)
auto op_desc = GenOpDesc(matched);
// 创建一个 FC OP
auto fc_op = LiteOpRegistry::Global().Create("fc");
// 找到原拓扑结构中的scope (作用域)和 valid_places (可支持设备类型)
auto mul = matched.at("mul")->stmt()->op();
auto* scope = mul->scope();
auto& valid_places = mul->valid_places();
// (2) 将 FC OP的 scope和 valid_places设置与fuse前相同,并在图中创建该节点(node)
fc_op->Attach(op_desc, scope);
auto* new_op_node = graph->GraphCreateInstructNode(fc_op, valid_places);
// (3) 将FC节点连接到输入输出(var_node)
IR_NODE_LINK_TO(matched.at("W"), new_op_node);
IR_NODE_LINK_TO(matched.at("x"), new_op_node);
IR_NODE_LINK_TO(matched.at("b"), new_op_node);
IR_NODE_LINK_TO(new_op_node, matched.at("Out"));
}
```
#### 2、注册fc_fuse_pass
(1)在`lite/core/mir/fusion`路径下新建`fc_fuse_pass.cc``fc_fuse_pass.h` 文件
(2)在`fc_fuse_pass.h` 文件中,继承`ProgramPass`定义`FcFusePass`
```c++
#include "lite/core/mir/pass.h"
namespace paddle {
namespace lite {
namespace mir {
class FcFusePass : public ProgramPass {
public:
void Apply(const std::unique_ptr<SSAGraph>& graph) override; namespace mir namespace lite namespace paddle
```
(3)在`fc_fuse_pass.cc` 文件中实现`FcFusePass::Apply()`接口,并注册`FcFusePass`
```c++
#include "lite/core/mir/pass_registry.h"
#include "lite/core/mir/example_pass.h"
namespace paddle {
namespace lite {
namespace mir {
void FcFusePass::Apply(const std::unique_ptr<SSAGraph>& graph) {
fusion::FcFuser fuser;
fuser(graph.get());namespace mir
} // namespace lite
} // namespace paddle
REGISTER_MIR_PASS(lite_fc_fuse_pass, paddle::lite::mir::FcFusePass)
.BindTargets({TARGET(kAny)}) // FcFusePass 可以在任何硬件平台执行
.BindKernel("fc"); // FcFusePass 绑定 fc_kernel
```
(4)修改`lite/core/mir/fusion/CMakeLists.txt`文件,将`fc_fuser.cc` 编译到`mir_fusers`
```cmake
lite_cc_library(fuse_fc
SRCS fc_fuser.cc
DEPS pattern_matcher_high_api)
set(mir_fusers
fuse_fc
...
CACHE INTERNAL "fusers")
```
(5)修改`lite/core/mir/CMakeLists.txt`文件,将`fc_fuse_pass.cc` 编译到`mir_pass`
```cmake
lite_cc_library(mir_passes
SRCS
fusion/fc_fuse_pass.cc
...
DEPS mir_pass types context ${mir_fusers} ${subgraph_passes})
```
#### 3、使用 fc_fuse_pass
(1) `lite/api/paddle_use_passes.h`使用`USE_LITE_PASS`宏来引入新加入的pass
```c++
USE_MIR_PASS(lite_fc_fuse_pass);
```
(2) 在`lite/core/optimizer.h`文件的`Optimizer::Run()`函数中添加新注册的pass
```C++
class Optimizer {
public:
void Run(Program&& program,
const std::vector<Place>& valid_places,
core::KernelPickFactor kernel_pick_factor,
const std::vector<std::string>& passes = {}) {
...
if (passes.empty()) {
RunPasses(std::vector<std::string>{
{"lite_fc_fuse_pass", // the newly registered pass
...
"argument_type_display_pass"}});
} else {
RunPasses(passes);
}
exec_scope_ = program.exec_scope();
}
```
(3) 以上修改完成后,在CreatePredictor(CxxConfig)创建CxxPredictor时,模型优化过程会调用`lite_fc_fuse_pass `,扫描`mul + element_wise add`结构并替换为等效的Fc_OP。
# CV 图像预处理API接口介绍
请把编译脚本`Paddle-Lite/lite/too/build.sh``BUILD_CV`变量设置为`ON`, 其他编译参数设置请参考[源码编译](../source_compile), 以确保 Lite 可以正确编译。这样`CV`图像的加速库就会编译进去,且会生成`paddle_image_preprocess.h`的API文件
- 硬件平台: `ARM`
- 操作系统:`MAC``LINUX`
## CV 图像预处理功能
Lite 支持不同颜色空间的图像相互转换 `Convert` 、缩放 `Resize` 、翻转 `Flip`、旋转 `Rotate` 和图像数据转换为 `Tensor` 存储`ImageToTensor` 功能,下文将详细介绍每个功能的API接口。
### CV 枚举变量和结构体变量
- 颜色空间
```cpp
enum ImageFormat {
RGBA = 0,
BGRA,
RGB,
BGR,
GRAY,
NV21 = 11,
NV12,
};
```
- 翻转参数
```cpp
enum FlipParam {
X = 0, // flip along the X axis
Y, // flip along the Y axis
XY // flip along the XY axis
};
```
- 转换参数
```cpp
typedef struct {
int ih; // input height
int iw; // input width
int oh; // outpu theight
int ow; // output width
FlipParam flip_param; // flip, support x, y, xy
float rotate_param; // rotate, support 90, 180, 270
} TransParam;
```
### ImagePreprocess 类的成员变量
`ImagePreprocess` 类含有以下三个私有成员变量,通过构造函数进行初始化。
```cpp
private:
ImageFormat srcFormat_; // input image color format
ImageFormat dstFormat_; // output image color format
TransParam transParam_; // image transform parameter
// init
ImagePreprocess::ImagePreprocess(ImageFormat srcFormat, ImageFormat dstFormat, TransParam param) {
this->srcFormat_ = srcFormat;
this->dstFormat_ = dstFormat;
this->transParam_ = param;
}
```
### 颜色空间转换 Convert
`Convert` 函数支持颜色空间:GRAY、NV12(NV21)、RGB(BGR)和RGBA(BGRA)
+ 目前支持以下颜色空间的相互转换:
- GRAY2BGR
- GRAY2RGB
- BGR2RGB
- BGRA2BGR
- BGRA2RGB
- RGBA2RGB
- RGBA2BGR
- BGRA2RGBA
+ 目前支持以下颜色空间的单向转换:
- NV12—BGR
- NV21—BGR
- NV12—RGB
- NV21—RGB
- NV12—BGRA
- NV21—BGRA
- NV12—RGBA
- NV21—RGBA
+ `Convert` 功能的API接口
```cpp
// 方法一
void ImagePreprocess::imageCovert(const uint8_t* src, uint8_t* dst);
// 方法二
void ImagePreprocess::imageCovert(const uint8_t* src,
uint8_t* dst, ImageFormat srcFormat, ImageFormat dstFormat);
```
+ 第一个 `imageCovert` 接口,缺省参数来源于 `ImagePreprocess` 类的成员变量。故在初始化 `ImagePreprocess` 类的对象时,必须要给以下成员变量赋值:
- param srcFormat:`ImagePreprocess` 类的成员变量`srcFormat_`
- param dstFormat:`ImagePreprocess` 类的成员变量`dstFormat_`
- 第二个`imageCovert` 接口,可以直接使用
### 缩放 Resize
`Resize` 功能支持颜色空间:GRAY、NV12(NV21)、RGB(BGR)和RGBA(BGRA)
`Resize` 功能目前支持的方法:`bilinear`
+ `Resize` 功能的API接口
```cpp
// 方法一
void ImagePreprocess::imageResize(const uint8_t* src, uint8_t* dst);
// 方法二
void ImagePreprocess::imageResize(const uint8_t* src, uint8_t* dst, ImageFormat srcFormat, ImageFormat srcFormat, int srcw, int srch, int dstw, int dsth);
```
+ 第一个`imageResize` 接口,缺省参数来源于`ImagePreprocess` 类的成员变量。故在初始化`ImagePreprocess` 类的对象时,必须要给以下成员变量赋值:
- param srcFormat:`ImagePreprocess` 类的成员变量`dstFormat_`
- param srcw:`ImagePreprocess` 类的成员变量`transParam_.iw`
- param srch:`ImagePreprocess` 类的成员变量`transParam_.ih`
- param dstw:`ImagePreprocess` 类的成员变量`transParam_.ow`
- param dsth:`ImagePreprocess` 类的成员变量`transParam_.ow`
- 第二个`imageResize` 接口,可以直接使用
### 旋转 Rotate
`Rotate` 功能支持颜色空间:GRAY、RGB(BGR)和RGBA(BGRA)
`Rotate` 功能目前支持的角度:90、180 和 270
+ `Rotate` 功能的API接口
```cpp
// 方法一
void ImagePreprocess::imageRotate(const uint8_t* src, uint8_t* dst);
// 方法二
void ImagePreprocess::imageRotate(const uint8_t* src, uint8_t* dst, ImageFormat srcFormat, ImageFormat srcFormat, int srcw, int srch, float degree);
```
+ 第一个`imageRotate` 接口,缺省参数来源于`ImagePreprocess` 类的成员变量。故在初始化`ImagePreprocess` 类的对象时,必须要给以下成员变量赋值:
- param srcFormat:`ImagePreprocess` 类的成员变量`dstFormat_`
- param srcw:`ImagePreprocess` 类的成员变量`transParam_.ow`
- param srch:`ImagePreprocess` 类的成员变量`transParam_.oh`
- param degree:`ImagePreprocess` 类的成员变量`transParam_.rotate_param`
- 第二个`imageRotate` 接口,可以直接使用
### 翻转 Flip
`Flip` 功能支持颜色空间:GRAY、RGB(BGR)和RGBA(BGRA)
`Flip` 功能目前支持的功能:沿X轴翻转、沿Y轴翻转和沿XY轴翻转
+ `Flip` 功能的API接口
```cpp
// 方法一
void ImagePreprocess::imageFlip(const uint8_t* src, uint8_t* dst);
// 方法二
void ImagePreprocess::imageFlip(const uint8_t* src, uint8_t* dst, ImageFormat srcFormat, ImageFormat srcFormat, int srcw, int srch, FlipParam flip_param);
```
+ 第一个`imageFlip` 接口,缺省参数来源于`ImagePreprocess` 类的成员变量。故在初始化`ImagePreprocess` 类的对象时,必须要给以下成员变量赋值:
- param srcFormat:`ImagePreprocess` 类的成员变量`dstFormat_`
- param srcw:`ImagePreprocess` 类的成员变量`transParam_.ow`
- param srch:`ImagePreprocess` 类的成员变量`transParam_.oh`
- param flip_param:`ImagePreprocess` 类的成员变量`transParam_.flip_param`
- 第二个`imageFlip` 接口,可以直接使用
### Image2Tensor
`Image2Tensor` 功能支持颜色空间:RGB(BGR)和RGBA(BGRA)
`Image2Tensor` 功能目前支持的Layout:`NCHW``NHWC`
`Image2Tensor` 不仅完成图像转换为`Tensor`数据处理,而且还完成了图像数据的归一化处理
+ `Image2Tensor` 功能的API接口
```cpp
// 方法一
void ImagePreprocess::image2Tensor(const uint8_t* src, Tensor* dstTensor, LayoutType layout, float* means, float* scales);
// 方法二
void ImagePreprocess::image2Tensor(const uint8_t* src, Tensor* dstTensor, ImageFormat srcFormat, srcw, int srch, LayoutType layout, float* means, float* scales;
```
+ 第一个`image2Tensor` 接口,缺省参数来源于`ImagePreprocess` 类的成员变量。故在初始化`ImagePreprocess` 类的对象时,必须要给以下成员变量赋值:
- param srcFormat:`ImagePreprocess` 类的成员变量`dstFormat_`
- param srcw:`ImagePreprocess` 类的成员变量`transParam_.ow`
- param srch:`ImagePreprocess` 类的成员变量`transParam_.oh`
- 第二个`image2Tensor` 接口,可以直接使用
## CV 图像预处理 Demo 示例
例子:输入 `1920x1080` 大小的 `NV12` 图像src,输出 `960x540` 大小 `RGB` 格式的图像dst;然后,完成 `90` 度旋转和沿 `X` 轴翻转功能;最后,用 `NHWC` 格式存储在Tensor里。
定义 `ImagePreprocess` 类的对象,初始化成员变量
```cpp
// init
srcFormat = ImageFormat::NV12;
dstFormat = ImageFormat::RGB;
srch = 1920;
srcw = 1080;
dsth = 960;
dstw = 540;
flip_param = FlipParam::X;
degree = 90;
layout = LayoutType::NHWC
// 方法一:
TransParam tparam;
tparam.ih = srch;
tparam.iw = srcw;
tparam.oh = dsth;
tparam.ow = dstw;
tparam.flip_param = flip_param;
tparam.rotate_param = degree;
ImagePreprocess image_preprocess(srcFormat, dstFormat, tparam);
// 方法二:
ImagePreprocess image_preprocess();
```
### imageConvert Demo
```cpp
// 方法一:
image_preprocess.imageCovert(src, lite_dst);
// 方法二:
image_preprocess.imageCovert(src, lite_dst, (ImageFormat)srcFormat, (ImageFormat)dstFormat);
```
### imageResize Demo
```cpp
// 方法一:
image_preprocess.imageResize(lite_dst, resize_tmp);
// 方法二:
image_preprocess.imageResize(lite_dst,resize_tmp, (ImageFormat)dstFormat, srcw,
srch, dstw, dsth);
```
### imageRotate Demo
```cpp
// 方法一:
image_preprocess.imageRotate(resize_tmp, tv_out_ratote);
// 方法二:
image_preprocess.imageRotate(resize_tmp,tv_out_ratote, (ImageFormat)dstFormat, dstw, dsth, degree);
```
### imageFlip Demo
```cpp
// 方法一:
image_preprocess.imageFlip(tv_out_ratote, tv_out_flip);
// 方法二:
image_preprocess.imageFlip(tv_out_ratote, tv_out_flip, (ImageFormat)dstFormat dstw, dsth, flip_param);
```
### image2Tensor Demo
```cpp
// 方法一:
image_preprocess.image2Tensor(tv_out_flip, &dst_tensor, layout, means, scales);
// 方法二:
image_preprocess.image2Tensor(tv_out_flip, &dst_tensor,(ImageFormat)dstFormat, dstw, dsth, layout, means, scales);
```
# 使用华为NPU
Paddle Lite是首款支持华为自研达芬奇架构NPU(Kirin 810/990 SoC搭载的NPU)的预测框架。
原理是在线分析Paddle模型,将Paddle算子转成HiAI IR后,调用HiAI IR/Builder/Runtime APIs生成并执行HiAI模型。
## 已支持的设备
- 华为nova5、nova5i pro、mate30、mate30 pro、mate30 5G、荣耀v30,以及即将推出的mate40、p40。据华为透露,今后上市的大部分手机都会搭载其自研达芬奇架构NPU。
## 已支持的模型
- MobileNetV1
- MobileNetV2
- ResNet-18/50
- ShuffleNetV2
- CycleGAN (暂时需要华为内部rom的支持)
- 百度内部业务模型(由于涉密,不方便透露具体细节)
## 已支持(或部分支持)的Paddle算子
- sigmoid
- relu
- tanh
- relu_clipped
- leaky_relu
- softsign
- hard_sigmoid
- batch_norm
- concat
- conv2d
- depthwise_conv2d
- conv2d_transpose
- dropout
- elementwise_add
- elementwise_sub
- elementwise_mul
- elementwise_div
- fusion_elementwise_add_activation
- fusion_elementwise_sub_activation
- fusion_elementwise_mul_activation
- fusion_elementwise_div_activation
- fc
- bilinear_interp
- nearest_interp
- matmul
- mul
- pad2d
- pool2d
- reduce_mean
- reshape
- reshape2
- scale
- shuffle_channel
- softmax
- split
- sqrt
- square
- transpose
- transpose2
- unsqueeze
- unsqueeze2
- instance_norm (暂时需要华为内部rom的支持)
- layer_norm (暂时需要华为内部rom的支持)
## 编译支持NPU的Paddle Lite库
- 从https://developer.huawei.com/consumer/cn/hiai/下载华为HiAI DDK后解压到任意路径(注意:华为提供了多个版本的DDK,我们需要下载针对麒麟810/990芯片HiAI Foundation开发套件,例如最新的[DDK V310版本](https://obs.cn-north-2.myhwclouds.com/hms-ds-wf/sdk/hwhiai-ddk-100.310.011.010.zip))。
- 将HiAI DDK中的ai_ddk_lib目录拷贝至Paddle Lite源码根目录后,使用[NPU编译脚本](https://github.com/PaddlePaddle/Paddle-Lite/blob/develop/lite/tools/build_npu.sh)编译full_publish和tiny_publish。
注意:以下是HiAI DDK V310版解压后的目录结构,需要将ai_ddk_lib目录拷贝至Paddle Lite源码根目录。
```shell
- app_sample
- ddk
- ai_ddk_lib
- include
- lib # for armv7
- lib64 # for armv8
- document
- tools
```
- full_publish and tiny_publish for armv8,由于HiAI DDK的armv7和armv8的so库均基于c++_shared构建,因此,建议使用c++_shared编译Paddle Lite。
```shell
$ ./lite/tools/build_npu.sh --arm_os=android --arm_abi=armv8 --arm_lang=gcc --android_stl=c++_shared full_publish
$ ./lite/tools/build_npu.sh --arm_os=android --arm_abi=armv8 --arm_lang=gcc --android_stl=c++_shared tiny_publish
```
- full_publish and tiny_publish for armv7
```shell
$ ./lite/tools/build_npu.sh --arm_os=android --arm_abi=armv7 --arm_lang=gcc --android_stl=c++_shared full_publish
$ ./lite/tools/build_npu.sh --arm_os=android --arm_abi=armv7 --arm_lang=gcc --android_stl=c++_shared tiny_publish
```
注意:为了保证编译环境一致,建议参考[源码编译](../installation/source_compile)中的Docker开发环境进行配置,然后再执行上述命令。
## 优化生成NPU模型
- model_optimize_tool工具已经支持生成NPU模型,仅需要将valid_targets设置为npu,arm即可,具体参考[模型转化方法](../user_guides/model_optimize_tool)
```shell
./model_optimize_tool --model_dir=<model_param_dir> \
--model_file=<model_path> \
--param_file=<param_path> \
--optimize_out_type=(protobuf|naive_buffer) \
--optimize_out=<output_optimize_model_dir> \
--valid_targets=npu,arm \
--prefer_int8_kernel=(true|false) \
--record_tailoring_info =(true|false)
```
- model_optimize_tool生成的模型只是标记了NPU支持的Paddle算子,并没有真正生成NPU HiAI模型,只有在执行时才会将标记的Paddle算子转成HiAI IR,最终生成并执行HiAI模型,具体实现参考PR[2576](https://github.com/PaddlePaddle/Paddle-Lite/pull/2576)
- 不同模型,不同型号(ROM版本)的华为手机,在执行阶段,由于某些Paddle算子无法完全转成HiAI IR,或目标手机的HiAI版本过低等原因,可能导致HiAI模型无法成功生成,在这种情况下,Paddle Lite会调用CPU版算子进行运算完成整个预测任务。
## 通过JAVA接口加载并执行NPU模型
- 使用方法和[Java实例](../user_guides/java_demo)一致,无需额外设置任何参数,只需将模型换成NPU模型即可。[Paddle-Lite-Demo](https://github.com/PaddlePaddle/Paddle-Lite-Demo)中的Image Classification Demo for Android是同时支持CPU和NPU两种模型的图像分类Demo。
注意:在拷贝libpaddle_lite_jni.so的时候,由于依赖HiAI DDK so和libc++_shared.so库,需要将HiAI DDK中ai_ddk_lib/lib或ai_ddk_lib/lib64目录下的所有so和libc++_shared.so,拷到libpaddle_lite_jni.so同级目录下。
## 通过C++接口加载并执行NPU模型
- 使用方法和[C++实例](../user_guides/cpp_demo)一致,同样无需额外设置任何参数,只需将模型换成NPU模型即可。
注意:1)不能使用安卓模拟器,需要使用真实设备,且必须是支持NPU的华为手机。2)在使用adb push命令向手机推送目标程序时,需要将HiAI DDK中ai_ddk_lib/lib或ai_ddk_lib/lib64目录下的所有so和libc++_shared.so,推送到目标程序同级目录下。
## 其它说明
- 华为达芬奇架构的NPU内部大量采用float16进行运算,因此,预测结果会存在偏差,但大部分情况下精度不会有较大损失,可参考[Paddle-Lite-Demo](https://github.com/PaddlePaddle/Paddle-Lite-Demo)中Image Classification Demo for Android对同一张图片CPU与NPU的预测结果。
- 华为Kirin 810/990 Soc搭载的自研达芬奇架构的NPU,与Kirin 970/980 Soc搭载的寒武纪NPU不一样,同样的,与Hi3559A、Hi3519A使用的NNIE也不一样,Paddle Lite只支持华为自研达芬奇架构NPU。
- 我们正在持续增加能够适配HiAI IR的Paddle算子bridge/converter,以便适配更多Paddle模型,同时华为研发同学也在持续对HiAI IR性能进行优化。
......@@ -39,6 +39,7 @@
- expand
- fake_channel_wise_dequantize_max_abs
- fake_dequantize_max_abs
- fake_quantize_abs_max
- fake_quantize_dequantize_moving_average_abs_max
- fake_quantize_moving_average_abs_max
- fake_quantize_range_abs_max
......@@ -56,10 +57,11 @@
- fusion_elementwise_mul_activation
- fusion_elementwise_sub_activation
- gather
- gelu
- generate_proposals
- graph_op
- greater_equal
- greater_than
- grid_sampler
- gru
- gru_unit
- hard_sigmoid
......@@ -93,7 +95,7 @@
- nearest_interp
- negative
- norm
- notequal
- not_equal
- pad2d
- pool2d
- power
......@@ -127,6 +129,7 @@
- sequence_expand
- sequence_expand_as
- sequence_pool
- sequence_pool_concat
- sequence_reshape
- sequence_reverse
- sequence_softmax
......@@ -144,6 +147,7 @@
- squeeze
- squeeze2
- stack
- subgraph
- swish
- tanh
- top_k
......@@ -216,6 +220,7 @@
- generate_proposals
- greater_equal
- greater_than
- grid_sampler
- gru
- gru_unit
- hard_sigmoid
......@@ -301,6 +306,9 @@
- gelu
- gru
- layer_norm
- leaky_relu
- lookup_table
- lookup_table_v2
- match_matrix_tensor
- matmul
- mul
......@@ -378,9 +386,11 @@
- yolo_box
### OpenCL kernels
- concat
- conv2d
- depthwise_conv2d
- elementwise_add
- elementwise_mul
- fc
- fusion_elementwise_add_activation
- layout
......@@ -388,5 +398,10 @@
- io_copy
- io_copy_once
- mul
- nearest_interp
- pool2d
- relu
- reshape
- reshape2
- scale
- sigmoid
# 使用X86预测库
Paddle-Lite 支持在Docker或Linux环境编译x86预测库。环境搭建参考[环境准备](../installation/source_compile)
(注意:非docker Linux环境需要是Ubuntu16.04)
## 编译
1、 下载代码
```bash
git clone https://github.com/PaddlePaddle/Paddle-Lite.git
#需要切换到 release/v2.0.0之后版本
git checkout <release_tag>
```
2、 源码编译
```bash
cd Paddle-Lite
./lite/tools/build.sh x86
```
## 编译结果说明
x86编译结果位于 `build.lite.x86/inference_lite_lib`
**具体内容**说明:
1、 `bin`文件夹:可执行工具文件 `test_model_bin`
2、 `cxx`文件夹:包含c++的库文件与相应的头文件
- `include` : 头文件
- `lib` : 库文件
- 打包的静态库文件:
- `libpaddle_api_full_bundled.a` :包含 full_api 和 light_api 功能的静态库
- `libpaddle_api_light_bundled.a` :只包含 light_api 功能的静态库
- 打包的动态态库文件:
- `libpaddle_full_api_shared.so` :包含 full_api 和 light_api 功能的动态库
- `libpaddle_light_api_shared.so`:只包含 light_api 功能的动态库
3、 `third_party` 文件夹:第三方库文件
## x86预测API使用示例
```c++
#include <gflags/gflags.h>
#include <iostream>
#include <vector>
#include "paddle_api.h" // NOLINT
#include "paddle_use_kernels.h" // NOLINT
#include "paddle_use_ops.h" // NOLINT
#include "paddle_use_passes.h" // NOLINT
using namespace paddle::lite_api; // NOLINT
DEFINE_string(model_dir, "", "Model dir path.");
DEFINE_string(optimized_model_dir, "", "Optimized model dir.");
DEFINE_bool(prefer_int8_kernel, false, "Prefer to run model with int8 kernels");
int64_t ShapeProduction(const shape_t& shape) {
int64_t res = 1;
for (auto i : shape) res *= i;
return res;
}
void RunModel() {
// 1. Set CxxConfig
CxxConfig config;
config.set_model_file(FLAGS_model_dir + "model");
config.set_param_file(FLAGS_model_dir + "params");
config.set_valid_places({
lite_api::Place{TARGET(kX86), PRECISION(kFloat)}
});
// 2. Create PaddlePredictor by CxxConfig
std::shared_ptr<PaddlePredictor> predictor =
CreatePaddlePredictor<CxxConfig>(config);
// 3. Prepare input data
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
input_tensor->Resize(shape_t({1, 3, 224, 224}));
auto* data = input_tensor->mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// 4. Run predictor
predictor->Run();
// 5. Get output
std::unique_ptr<const Tensor> output_tensor(
std::move(predictor->GetOutput(0)));
std::cout << "Output dim: " << output_tensor->shape()[1] << std::endl;
for (int i = 0; i < ShapeProduction(output_tensor->shape()); i += 100) {
std::cout << "Output[" << i << "]:" << output_tensor->data<float>()[i] << std::endl;
}
}
int main(int argc, char** argv) {
google::ParseCommandLineFlags(&argc, &argv, true);
RunModel();
return 0;
}
```
# C++ API文档
## CreatePaddlePredictor
```c++
template <typename ConfigT>
std::shared_ptr<PaddlePredictor> CreatePaddlePredictor(const ConfigT&);
```
`CreatePaddlePredictor`用来根据`MobileConfig`构建预测器。
示例:
```c++
// 设置MobileConfig
MobileConfig config;
config.set_model_dir(FLAGS_model_dir);
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
```
参数:
- `config(MobileConfig)` - 用于构建Predictor的配置信息。
返回:`PaddlePredictor`指针
返回类型:`std::shared_ptr<PaddlePredictor>`
## CxxConfig
```c++
class CxxConfig;
```
`CxxConfig`用来配置构建CxxPredictor的配置信息,如protobuf格式的模型地址、能耗模式、工作线程数、place信息等等。
示例:
```c++
config = CxxConfig()
# 设置模型目录,加载非combined模型时使用
config.set_model_dir(<your_model_dir_path>)
# 设置工作线程数
config.set_threads(4);
# 设置能耗模式
config.set_power_mode(PowerMode.LITE_POWER_NO_BIND)
# 设置valid places
places = [Place(TargetType.ARM, PrecisionType.FP32)]
config.set_valid_places(places)
# 根据CxxConfig创建CxxPredictor
predictor = create_paddle_predictor(config)
```
### `set_model_dir(model_dir)`
设置模型文件夹路径,当需要从磁盘加载非combined模型时使用。
参数:
- `model_dir(str)` - 模型文件夹路径
返回:`None`
返回类型:`None`
### `model_dir()`
返回设置的模型文件夹路径。
参数:
- `None`
返回:模型文件夹路径
返回类型:`str`
### `set_model_file(model_file)`
设置模型文件路径,加载combined形式模型时使用。
参数:
- `model_file(str)` - 模型文件路径
返回类型:`None`
### `model_file()`
获取设置模型文件路径,加载combined形式模型时使用。
参数:
- `None`
返回:模型文件路径
返回类型:`str`
### `set_param_file(param_file)`
设置模型参数文件路径,加载combined形式模型时使用。
参数:
- `param_file(str)` - 模型文件路径
返回类型:`None`
### `param_file()`
获取设置模型参数文件路径,加载combined形式模型时使用。
参数:
- `None`
返回:模型参数文件路径
返回类型:`str`
### `set_valid_places(valid_places)`
设置可用的places列表。
参数:
- `valid_places(list)` - 可用place列表。
返回类型:`None`
示例:
```c++
config = CxxConfig()
# 设置模型目录,加载非combined模型时使用
config.set_model_dir(<your_model_dir_path>)
# 设置valid places
# 注意,valid_places列表中Place的排序表明了用户对Place的偏好程度,如用户想优先使用ARM上Int8精度的
# kernel,则应把Place(TargetType.ARM, PrecisionType.INT8)置于valid_places列表的首位。
places = [Place(TargetType.ARM, PrecisionType.INT8),
Place(TargetType.ARM, PrecisionType.FP32)]
config.set_valid_places(places)
# 根据CxxConfig创建CxxPredictor
predictor = create_paddle_predictor(config)
```
### `set_power_mode(mode)`
设置CPU能耗模式。若不设置,则默认使用`PowerMode.LITE_POWER_HIGH`
*注意:只在开启`OpenMP`时生效,否则系统自动调度。此函数只在使用`LITE_WITH_ARM`编译选项下生效。*
参数:
- `mode(PowerMode)` - CPU能耗模式
返回:`None`
返回类型:`None`
### `power_mode()`
获取设置的CPU能耗模式。
*注意:此函数只在使用`LITE_WITH_ARM`编译选项下生效。*
参数:
- `None`
返回:设置的CPU能耗模式
返回类型:`PowerMode`
### `set_threads(threads)`
设置工作线程数。若不设置,则默认使用单线程。
*注意:只在开启`OpenMP`的模式下生效,否则只使用单线程。此函数只在使用`LITE_WITH_ARM`编译选项下生效。*
参数:
- `threads(int)` - 工作线程数
返回:`None`
返回类型:`None`
### `threads()`
获取设置的工作线程数。
*注意:此函数只在使用`LITE_WITH_ARM`编译选项下生效。*
参数:
- `None`
返回:工作线程数
返回类型:`int`
### `set_x86_math_library_num_threads(threads)`
设置CPU Math库线程数,CPU核心数支持情况下可加速预测。默认为1,并且仅在x86下有效。
参数:
- `threads(int)` - CPU Math库线程数。
返回:`None`
返回类型:`None`
### `x86_math_library_num_threads()`
返回CPU Math库线程数,CPU核心数支持情况下可加速预测。仅在x86下有效。
参数:
- `None`
返回:CPU Math库线程数。
返回类型:`int`
## MobileConfig
```c++
class MobileConfig;
```
`MobileConfig`用来配置构建轻量级PaddlePredictor的配置信息,如NaiveBuffer格式的模型地址、模型的内存地址(从内存加载模型时使用)、能耗模式、工作线程数等等。
*注意:输入的模型需要使用[Model Optimize Tool](../model_optimize_tool)转化为NaiveBuffer格式的优化模型。*
示例:
```c++
MobileConfig config;
// 设置NaiveBuffer格式模型目录,从文件加载模型时使用
config.set_model_dir(FLAGS_model_dir);
// 设置工作线程数
config.set_threads(4);
// 设置能耗模式
config.set_power_mode(LITE_POWER_HIGH);
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
```
### `set_model_from_file(model_dir)`
设置模型文件,当需要从磁盘加载模型时使用。
参数:
- `model_dir(std::string)` - 模型文件路径
返回:`None`
返回类型:`void`
### `set_model_dir(model_dir)`
**注意**:Lite模型格式在release/v2.3.0之后修改,本接口为加载老格式模型的接口,将在release/v3.0.0废弃。建议替换为`set_model_from_file`接口。
设置模型文件夹路径,当需要从磁盘加载模型时使用。
参数:
- `model_dir(std::string)` - 模型文件夹路径
返回:`None`
返回类型:`void`
### `model_dir()`
返回设置的模型文件夹路径。
参数:
- `None`
返回:模型文件夹路径
返回类型:`std::string`
### `set_model_from_buffer(model_buffer)`
设置模型的内存数据,当需要从内存加载模型时使用。
参数:
- `model_buffer(std::string)` - 内存中的模型数据
返回:`None`
返回类型:`void`
### `set_model_buffer(model_buffer, model_buffer_size, param_buffer, param_buffer_size)`
**注意**:Lite模型格式在release/v2.3.0之后修改,本接口为加载老格式模型的接口,将在release/v3.0.0废弃。建议替换为`set_model_from_buffer`接口。
设置模型、参数的内存地址,当需要从内存加载模型时使用。
示例:
```c++
// 读取模型文件到内存
std::string model_buffer = ReadFile(FLAGS_model_path);
std::string params_buffer = lite::ReadFile(FLAGS_params_path);
// 设置MobileConfig
lite_api::MobileConfig config;
config.set_model_buffer(model_buffer.c_str(), model_buffer.size(),
params_buffer.c_str(), params_buffer.size());
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
```
参数:
- `model_buffer(const char*)` - 内存中模型结构数据。
- `model_buffer_size(size_t)` - 内存中模型结构数据的大小。
- `param_buffer(const char*)` - 内存中模型参数数据。
- `param_buffer_size(size_t)` - 内存中模型参数数据的大小。
返回:`None`
返回类型:`Void`
### `model_from_memory()`
是否从内存中加载模型,当使用`set_model_buffer`接口时返回`true`
参数:
- `None`
返回:是否从内存加载模型
返回类型:`bool`
### `model_buffer()`
获取内存中模型结构数据。
参数:
- `None`
返回:内存中模型结构数据
返回类型:`const std::string&`
### `param_buffer()`
获取内存中模型参数数据。
参数:
- `None`
返回:内存中模型结构数据
返回类型:`const std::string&`
### `set_power_mode(mode)`
设置CPU能耗模式。若不设置,则默认使用`LITE_POWER_HIGH`
*注意:只在开启`OpenMP`时生效,否则系统自动调度。*
参数:
- `mode(PowerMode)` - CPU能耗模式
返回:`None`
返回类型:`void`
### `power_mode()`
获取设置的CPU能耗模式。
参数:
- `None`
返回:设置的CPU能耗模式
返回类型:`PowerMode`
### `set_threads(threads)`
设置工作线程数。若不设置,则默认使用单线程。
*注意:只在开启`OpenMP`的模式下生效,否则只使用单线程。*
参数:
- `threads(int)` - 工作线程数
返回:`None`
返回类型:`void`
### `threads()`
获取设置的工作线程数。
参数:
- `None`
返回:工作线程数
返回类型:`int`
## PaddlePredictor
```c++
class PaddlePredictor
```
`PaddlePredictor`是Paddle-Lite的预测器,由`CreatePaddlePredictor`根据`MobileConfig`进行创建。用户可以根据PaddlePredictor提供的接口设置输入数据、执行模型预测、获取输出以及获得当前使用lib的版本信息等。
示例:
```c++
int64_t ShapeProduction(const shape_t& shape) {
int64_t res = 1;
for (auto i : shape) res *= i;
return res;
}
// 设置MobileConfig
MobileConfig config;
config.set_model_dir(FLAGS_model_dir);
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
// 获得模型的输入和输出名称
std::vector<std::string> input_names = predictor->GetInputNames();
for (int i = 0; i < input_names.size(); i ++) {
printf("Input name[%d]: %s\n", i, input_names[i].c_str());
}
std::vector<std::string> output_names = predictor->GetOutputNames();
for (int i = 0; i < output_names.size(); i ++) {
printf("Output name[%d]: %s\n", i, output_names[i].c_str());
}
// 准备输入数据
// (1)根据index获取输入Tensor
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
// (2)根据名称获取输入Tensor
// std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInputByName(input_names[0])));
input_tensor->Resize({1, 3, 224, 224});
auto* data = input_tensor->mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// 执行预测
predictor->Run();
// 获取输出
// (1)根据index获取输出Tensor
std::unique_ptr<const Tensor> output_tensor(std::move(predictor->GetOutput(0)));
// (2)根据名称获取输出Tensor
// std::unique_ptr<const Tensor> output_tensor(std::move(predictor->GetOutput(output_names[0])));
printf("Output dim: %d\n", output_tensor->shape()[1]);
for (int i = 0; i < ShapeProduction(output_tensor->shape()); i += 100) {
printf("Output[%d]: %f\n", i, output_tensor->data<float>()[i]);
}
```
### `GetInput(index)`
获取输入Tensor指针,用来设置模型的输入数据。
参数:
- `index(int)` - 输入Tensor的索引
返回:第`index`个输入`Tensor`的指针
返回类型:`std::unique_ptr<Tensor>`
### `GetOutput(index)`
获取输出Tensor的指针,用来获取模型的输出结果。
参数:
- `index(int)` - 输出Tensor的索引
返回:第`index`个输出Tensor`的指针
返回类型:`std::unique_ptr<Tensor>`
### `GetInputNames()`
获取所有输入Tensor的名称。
参数:
- `None`
返回:所有输入Tensor的名称
返回类型:`std::vector<std::string>`
### `GetOutputNames()`
获取所有输出Tensor的名称。
参数:
- `None`
返回:所有输出Tensor的名称
返回类型:`std::vector<std::string>`
### `GetInputByName(name)`
根据名称获取输出Tensor的指针,用来获取模型的输出结果。
参数:
- `name(const std::string)` - 输入Tensor的名称
返回:输入Tensor`的指针
返回类型:`std::unique_ptr<Tensor>`
### `GetTensor(name)`
根据名称获取输出Tensor的指针。
**注意**`GetTensor`接口是为开发者设计的调试接口,可以输出[转化](../model_optimize_tool)后模型中的任一节点。如果出现`GetTensor(InputName)`返回值为空`Tensor`,可能原因是以该`InputName`命名的Tensor在模型转化的**子图融合**过程被融合替换了。
参数:
- `name(const std::string)` - Tensor的名称
返回:指向`const Tensor`的指针
返回类型:`std::unique_ptr<const Tensor>`
### `Run()`
执行模型预测,需要在***设置输入数据后***调用。
参数:
- `None`
返回:`None`
返回类型:`void`
### `GetVersion()`
用于获取当前lib使用的代码版本。若代码有相应tag则返回tag信息,如`v2.0-beta`;否则返回代码的`branch(commitid)`,如`develop(7e44619)`
参数:
- `None`
返回:当前lib使用的代码版本信息
返回类型:`std::string`
## TargetType
```c++
class TargetType;
```
`TargetType`为目标设备硬件类型,用户可以根据应用场景选择硬件平台类型。
枚举型变量`TargetType`的所有可能取值包括:
`{X86, CUDA, ARM, OpenCL, FPGA, NPU}`
## PrecisionType
```c++
class PrecisionType {FP32};
```
`PrecisionType`为模型中Tensor的数据精度,默认值为FP32(float32)。
枚举型变量`PrecisionType`的所有可能取值包括:
`{FP32, INT8, INT32, INT64}`
## DataLayoutType
```c++
class DataLayoutType {NCHW};
```
`DataLayoutType`为Tensor的数据格式,默认值为NCHW(number, channel, height, weigth)。
枚举型变量`DataLayoutType`的所有可能取值包括:
` {NCHW, NHWC}`
## Place
```c++
class Place{
TargetType target;
PrecisionType precision{FP32};
DataLayoutType layout{NCHW}
}
```
`Place``TargetType``PrecisionType``DataLayoutType`的集合,说明运行时的设备类型、数据精度和数据格式。
示例:
```C++
Place{TargetType(ARM), PrecisionType(FP32), DataLayoutType(NCHW)}
```
## PowerMode
```c++
enum PowerMode;
```
`PowerMode`为ARM CPU能耗模式,用户可以根据应用场景设置能耗模式获得最优的能效比。
示例:
```c++
MobileConfig config;
// 设置NaiveBuffer格式模型目录
config.set_model_dir(FLAGS_model_dir);
// 设置能耗模式
config.set_power_mode(LITE_POWER_HIGH);
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
```
PowerMode详细说明如下:
| 选项 | 说明 |
| :------------------: | ------------------------------------------------------------ |
| LITE_POWER_HIGH | 绑定大核运行模式。如果ARM CPU支持big.LITTLE,则优先使用并绑定Big cluster。如果设置的线程数大于大核数量,则会将线程数自动缩放到大核数量。如果系统不存在大核或者在一些手机的低电量情况下会出现绑核失败,如果失败则进入不绑核模式。 |
| LITE_POWER_LOW | 绑定小核运行模式。如果ARM CPU支持big.LITTLE,则优先使用并绑定Little cluster。如果设置的线程数大于小核数量,则会将线程数自动缩放到小核数量。如果找不到小核,则自动进入不绑核模式。 |
| LITE_POWER_FULL | 大小核混用模式。线程数可以大于大核数量。当线程数大于核心数量时,则会自动将线程数缩放到核心数量。 |
| LITE_POWER_NO_BIND | 不绑核运行模式(推荐)。系统根据负载自动调度任务到空闲的CPU核心上。 |
| LITE_POWER_RAND_HIGH | 轮流绑定大核模式。如果Big cluster有多个核心,则每预测10次后切换绑定到下一个核心。 |
| LITE_POWER_RAND_LOW | 轮流绑定小核模式。如果Little cluster有多个核心,则每预测10次后切换绑定到下一个核心。 |
## Tensor
```c++
class Tensor
```
Tensor是Paddle-Lite的数据组织形式,用于对底层数据进行封装并提供接口对数据进行操作,包括设置Shape、数据、LoD信息等。
*注意:用户应使用`PaddlePredictor`的`GetInput`和`GetOuput`接口获取输入/输出的`Tensor`。*
示例:
```c++
int64_t ShapeProduction(const shape_t& shape) {
int64_t res = 1;
for (auto i : shape) res *= i;
return res;
}
// 设置MobileConfig
MobileConfig config;
config.set_model_dir(FLAGS_model_dir);
// 根据MobileConfig创建PaddlePredictor
std::shared_ptr<PaddlePredictor> predictor = CreatePaddlePredictor<MobileConfig>(config);
// 准备输入数据, 获取输入Tensor
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
// 设置输入Tensor维度信息
input_tensor->Resize({1, 3, 224, 224});
// 设置输入数据
auto* data = input_tensor->mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// 执行预测
predictor->Run();
// 获取输出Tensor
std::unique_ptr<const Tensor> output_tensor(std::move(predictor->GetOutput(0)));
// 获取输出Tensor维度
printf("Output dim: %d\n", output_tensor->shape()[1]);
// 获取输出Tensor数据
for (int i = 0; i < ShapeProduction(output_tensor->shape()); i += 100) {
printf("Output[%d]: %f\n", i, output_tensor->data<float>()[i]);
}
```
### `Resize(shape)`
设置Tensor的维度信息。
参数:
- `shape(std::vector<int64_t>)` - 维度信息
返回:`None`
返回类型:`void`
### `shape()`
获取Tensor的维度信息。
参数:
- `None`
返回:Tensor的维度信息
返回类型:`std::vector<int64_t>`
### `data<T>()`
```c++
template <typename T>
const T* data() const;
```
获取Tensor的底层数据的常量指针,根据传入的不同模型类型获取相应数据。用于读取Tensor数据。
示例:
```c++
std::unique_ptr<const Tensor> output_tensor(std::move(predictor->GetOutput(0)));
// 如果模型中输出为float类型
output_tensor->data<float>()
```
参数:
- `None`
返回:`Tensor`底层数据常量指针
返回类型:`const T*`
### `mutable_data<T>()`
```c++
template <typename T>
T* mutable_data() const;
```
获取Tensor的底层数据的指针,根据传入的不同模型类型获取相应数据。用于设置Tensor数据。
示例:
```c++
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
// 如果模型中输出为float类型
auto* data = input_tensor->mutable_data<float>();
// 设置Tensor数据
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
```
参数:
- `None`
返回:`Tensor`底层数据指针
返回类型:`T*`
### `SetLoD(lod)`
设置Tensor的LoD信息。
参数:
- `lod(std::vector<std::vector<uint64_t>>)` - Tensor的LoD信息
返回:`None`
返回类型:`void`
### `lod()`
获取Tensor的LoD信息
参数:
- `None`
返回:`Tensor`的LoD信息
返回类型:`std::vector<std::vector<uint64_t>>`
......@@ -13,6 +13,7 @@ Welcome to Paddle-Lite's documentation!
introduction/tech_highlights
introduction/architecture
introduction/support_hardware
.. toctree::
:maxdepth: 1
......@@ -34,7 +35,13 @@ Welcome to Paddle-Lite's documentation!
:caption: 使用指南
:name: sec-user-guides
user_guides/model_optimize_tool
user_guides/library_tailoring
user_guides/cuda
user_guides/fpga
user_guides/opencl
user_guides/cpp_demo
user_guides/java_demo
.. toctree::
:maxdepth: 1
......@@ -42,7 +49,12 @@ Welcome to Paddle-Lite's documentation!
advanced_user_guides/support_operation_list
advanced_user_guides/add_operation
advanced_user_guides/add_layout
advanced_user_guides/model_quantization
advanced_user_guides/add_new_pass
advanced_user_guides/npu
advanced_user_guides/x86
advanced_user_guides/cv
.. toctree::
:maxdepth: 1
......@@ -52,6 +64,8 @@ Welcome to Paddle-Lite's documentation!
:maxdepth: 1
:caption: API文档
api_reference/cxx_api_doc
.. toctree::
:maxdepth: 1
:caption: FAQ
......
# 预测库说明
Paddle-Lite的编译结果为预测库文件(包括静态库和动态库),具体编译过程参考[源码编译](./source_compile)
Lite预测库分为**基础预测库****全量预测库**:基础预测库只打包了基础模型需要的基础算子,预测库体积较小;全量预测库打包了所有的Lite算子,可以支持更多的模型,但是预测库的体积也更大。 编译时由编译选项 `build_extra`(默认为OFF)控制,`--build_extra=OFF`时编译基础预测库,`--build_extra=ON`时编译全量的预测库。
## 基础预测库
### 编译方法
编译时设置`--build_extra=OFF` (默认值) 或不指定即可编译出基础预测库。例如:
```
./lite/tools/build.sh --arm_os=android --arm_abi=armv8 --arm_lang=gcc --android_stl=c++_static tiny_publish
```
### 基础预测库支持的功能
(1)支持基础CV模型
(2)支持基础的in8量化模型
(3)支持[benchmark测试](../benchmark/benchmark)
### 基础预测库支持的基础模型:
1. fluid基础模型(paddle model 提供的基础模型9个)
```
mobileNetV1 mnasnet yolov3 ssd_mobilenetv1 shufflenet_v2
mobileNetV2 resnet50 unet squeezenet_v11
```
2. int8量化模型模型
```
mobilenet_v1 mobilenet_v2 resnet50
```
### 特点
轻量级预测库,体积更小,支持常用的基础模型。
## 全量预测库
### 编译方法
编译时设置`--build_extra=ON` 即可编译出全量预测库。例如:
```
./lite/tools/build.sh --arm_os=android --arm_abi=armv8 --arm_lang=gcc --android_stl=c++_static --build_extra=ON tiny_publish
```
### 全量预测库功能
(1) 基础预测库所有功能
(2)支持所有Paddle-Lite中注册的所有算子
### 特点
支持更多的硬件平台和算子,可以支持更多模型但体量更大。
# 源码编译
Paddle-Lite 提供了移动端的一键源码编译脚本 `lite/tools/build.sh`,编译流程如下:
1. 环境准备(选择其一):Docker交叉编译环境、Linux交叉编译环境
2. 编译:调用`build.sh`脚本一键编译
## 一、环境准备
目前支持三种编译的环境:
1. Docker 容器环境,
2. Linux(推荐 Ubuntu 16.04)环境,
3. Mac OS 环境。
### 1、 Docker开发环境
[Docker](https://www.docker.com/) 是一个开源的应用容器引擎, 使用沙箱机制创建独立容器,方便运行不同程序。Docker初学者可以参考[Docker使用方法](https://thenewstack.io/docker-station-part-one-essential-docker-concepts-tools-terminology/)正确安装Docker。
#### 准备Docker镜像
有两种方式准备Docker镜像,推荐从Dockerhub直接拉取Docker镜像
```shell
# 方式一:从Dockerhub直接拉取Docker镜像
docker pull paddlepaddle/paddle-lite:2.0.0_beta
# 方式二:本地源码编译Docker镜像
git clone https://github.com/PaddlePaddle/Paddle-Lite.git
cd Paddle-Lite/lite/tools
mkdir mobile_image
cp Dockerfile.mobile mobile_image/Dockerfile
cd mobile_image
docker build -t paddlepaddle/paddle-lite .
# 镜像编译成功后,可用`docker images`命令,看到`paddlepaddle/paddle-lite`镜像。
```
#### 进入Docker容器
在拉取Paddle-Lite仓库代码的上层目录,执行如下代码,进入Docker容器:
```shell
docker run -it \
--name paddlelite_docker \
-v $PWD/Paddle-Lite:/Paddle-Lite \
--net=host \
paddlepaddle/paddle-lite /bin/bash
```
该命令的含义:将容器命名为`paddlelite_docker``<container-name>`,将当前目录下的`Paddle-Lite`文件夹挂载到容器中的`/Paddle-Lite`这个根目录下,并进入容器中。至此,完成Docker环境的准备。
#### Docker常用命令
```shell
# 退出容器但不停止/关闭容器:键盘同时按住三个键:CTRL + q + p
# 启动停止的容器
docker start <container-name>
# 从shell进入已启动的容器
docker attach <container-name>
# 停止正在运行的Docker容器
docker stop <container-name>
# 重新启动正在运行的Docker容器
docker restart <container-name>
# 删除Docker容器
docker rm <container-name>
```
### 2、Linux 开发环境
#### Android
##### 交叉编译环境要求
- gcc、g++、git、make、wget、python、adb
- Java environment
- cmake(建议使用3.10或以上版本)
- Android NDK (建议ndk-r17c)
##### 具体步骤
安装软件部分以 Ubuntu 为例,其他 Linux 发行版类似。
```shell
# 1. Install basic software
apt update
apt-get install -y --no-install-recommends \
gcc g++ git make wget python unzip adb curl
# 2. Prepare Java env.
apt-get install -y default-jdk
# 3. Install cmake 3.10 or above
wget -c https://mms-res.cdn.bcebos.com/cmake-3.10.3-Linux-x86_64.tar.gz && \
tar xzf cmake-3.10.3-Linux-x86_64.tar.gz && \
mv cmake-3.10.3-Linux-x86_64 /opt/cmake-3.10 && \
ln -s /opt/cmake-3.10/bin/cmake /usr/bin/cmake && \
ln -s /opt/cmake-3.10/bin/ccmake /usr/bin/ccmake
# 4. Download Android NDK for linux-x86_64
# Note: Skip this step if NDK installed
# recommand android-ndk-r17c-darwin-x86_64
# ref: https://developer.android.com/ndk/downloads
cd /tmp && curl -O https://dl.google.com/android/repository/android-ndk-r17c-linux-x86_64.zip
cd /opt && unzip /tmp/android-ndk-r17c-linux-x86_64.zip
# 5. Add environment ${NDK_ROOT} to `~/.bashrc`
echo "export NDK_ROOT=/opt/android-ndk-r17c" >> ~/.bashrc
source ~/.bashrc
```
#### ARM Linux
适用于基于 ARMv8 和 ARMv7 架构 CPU 的各种开发板,例如 RK3399,树莓派等,目前支持交叉编译和本地编译两种方式,对于交叉编译方式,在完成目标程序编译后,可通过 scp 方式将程序拷贝到开发板运行。
##### 交叉编译
###### 编译环境要求
- gcc、g++、git、make、wget、python、scp
- cmake(建议使用3.10或以上版本)
###### 具体步骤
安装软件部分以 Ubuntu 为例,其他 Linux 发行版类似。
```shell
# 1. Install basic software
apt update
apt-get install -y --no-install-recommends \
gcc g++ git make wget python unzip
# 2. Install arm gcc toolchains
apt-get install -y --no-install-recommends \
g++-arm-linux-gnueabi gcc-arm-linux-gnueabi \
g++-arm-linux-gnueabihf gcc-arm-linux-gnueabihf \
gcc-aarch64-linux-gnu g++-aarch64-linux-gnu
# 3. Install cmake 3.10 or above
wget -c https://mms-res.cdn.bcebos.com/cmake-3.10.3-Linux-x86_64.tar.gz && \
tar xzf cmake-3.10.3-Linux-x86_64.tar.gz && \
mv cmake-3.10.3-Linux-x86_64 /opt/cmake-3.10 && \
ln -s /opt/cmake-3.10/bin/cmake /usr/bin/cmake && \
ln -s /opt/cmake-3.10/bin/ccmake /usr/bin/ccmake
```
##### 本地编译(直接在RK3399或树莓派上编译)
###### 编译环境要求
- gcc、g++、git、make、wget、python
- cmake(建议使用3.10或以上版本)
###### 具体步骤
安装软件部分以 Ubuntu 为例,其他 Linux 发行版本类似。
```shell
# 1. Install basic software
apt update
apt-get install -y --no-install-recomends \
gcc g++ make wget python unzip
# 2. install cmake 3.10 or above
wget https://www.cmake.org/files/v3.10/cmake-3.10.3.tar.gz
tar -zxvf cmake-3.10.3.tar.gz
cd cmake-3.10.3
./configure
make
sudo make install
```
之后可通过cmake --version查看cmake是否安装成功。
至此,完成 Linux 交叉编译环境的准备。
### 3、Mac OS 开发环境
#### 交叉编译环境要求
- gcc、git、make、curl、unzip、java
- cmake(Android编译请使用3.10版本,IOS编译请使用3.15版本)
- 编译Android: Android NDK (建议ndk-r17c)
- 编译IOS: XCode(Version 10.1)
#### 具体步骤
```bash
# 1. Install basic software
brew install curl gcc git make unzip wget
# 2. Install cmake: mac上实现IOS编译和Android编译要求的cmake版本不一致,可以根据需求选择安装。
# (1)在mac环境编译 Paddle-Lite 的Android版本,需要安装cmake 3.10
# mkdir /usr/local/Cellar/cmake/ && cd /usr/local/Cellar/cmake/
# wget https://cmake.org/files/v3.10/cmake-3.10.2-Darwin-x86_64.tar.gz
# tar zxf ./cmake-3.10.2-Darwin-x86_64.tar.gz
# mv cmake-3.10.2-Darwin-x86_64/CMake.app/Contents/ ./3.10.2
# ln -s /usr/local/Cellar/cmake/3.10.2/bin/cmake /usr/local/bin/cmake
# (2)在mac环境编译 Paddle-Lite 的IOS版本,需要安装cmake 3.15
# mkdir /usr/local/Cellar/cmake/ && cd /usr/local/Cellar/cmake/
# cd /usr/local/Cellar/cmake/
# wget https://cmake.org/files/v3.15/cmake-3.15.2-Darwin-x86_64.tar.gz
# tar zxf ./cmake-3.15.2-Darwin-x86_64.tar.gz
# mv cmake-3.15.2-Darwin-x86_64/CMake.app/Contents/ ./3.15.2
# ln -s /usr/local/Cellar/cmake/3.15.2/bin/cmake /usr/local/bin/cmake
# 3. Download Android NDK for Mac
# recommand android-ndk-r17c-darwin-x86_64
# ref: https://developer.android.com/ndk/downloads
# Note: Skip this step if NDK installed
cd ~/Documents && curl -O https://dl.google.com/android/repository/android-ndk-r17c-darwin-x86_64.zip
cd ~/Library && unzip ~/Documents/android-ndk-r17c-darwin-x86_64.zip
# 4. Add environment ${NDK_ROOT} to `~/.bash_profile`
echo "export NDK_ROOT=~/Library/android-ndk-r17c" >> ~/.bash_profile
source ~/.bash_profile
# 5. Install Java Environment
brew cask install java
# 6. 编译IOS需要安装XCode(Version 10.1),可以在App Store里安装。安装后需要启动一次并执行下面语句。
# sudo xcode-select -s /Applications/Xcode.app/Contents/Developer
```
至此,完成 Mac 交叉编译环境的准备。
**注意**: Mac上编译Paddle-Lite的full_publish版本时,Paddle-Lite所在路径中不可以含有中文字符
## 二、编译PaddleLite
### 下载代码
```shell
git clone https://github.com/PaddlePaddle/Paddle-Lite.git
cd Paddle-Lite
git checkout <release-version-tag>
```
### 编译模式与参数
编译脚本`./lite/tools/build.sh`,支持三种编译模式:
| 编译模式 | 介绍 | 适用对象 |
|:-------:|-----|:-------:|
| tiny_publish | 编译移动端部署库,无第三方库依赖 | 用户 |
| full_publish | 编译移动端部署库,有第三方依赖如protobuf、glags等,含有可将模型转换为无需protobuf依赖的naive buffer格式的工具,供tiny_publish库使用 | 用户 |
| test | 编译指定`arm_os``arm_abi`下的移动端单元测试 | 框架开发者 |
编译脚本`./lite/tools/build.sh`,追加参数说明:
| 参数 | 介绍 | 值 |
|-----------|-------------|-------------|
| --arm_os |必选,选择安装平台 | `android``ios``ios64``armlinux` |
| --arm_abi |必选,选择编译的arm版本,其中`armv7hf`为ARMLinux编译时选用| `armv8``armv7``armv7hf`(仅`armlinux`支持) |
| --arm_lang |arm_os=android时必选,选择编译器 | `gcc``clang`(`clang`当前暂不支持) |
| --android_stl |arm_os=android时必选,选择静态链接STL或动态链接STL | `c++_static``c++_shared`|
| --build_java | 可选,是否编译java预测库(默认为OFF) | `ON``OFF` |
| --build_extra | 可选,是否编译全量预测库(默认为OFF)。详情可参考[预测库说明](./library.html)。 | `ON``OFF` |
| target |必选,选择编译模式,`tiny_publish`为编译移动端部署库、`full_publish`为带依赖的移动端部署库、`test`为移动端单元测试、`ios`为编译ios端`tiny_publish` | `tiny_publish``full_publish``test``ios` |
### 编译代码
**<font color="orange" >注意</font>**<font color="orange" >:非开发者建议在编译前使用</font>[**“加速第三方依赖库的下载”**](#id22)<font color="orange" >的方法,加速工程中第三方依赖库的下载与编译。 </font>
#### 编译`tiny publish`动态库
##### Android
```shell
./lite/tools/build.sh \
--arm_os=android \
--arm_abi=armv8 \
--build_extra=OFF \
--arm_lang=gcc \
--android_stl=c++_static \
--build_extra=OFF \
tiny_publish
```
##### IOS
```shell
./lite/tools/build.sh \
--arm_os=ios64 \
--arm_abi=armv8 \
--build_extra=OFF \
ios
```
**注意:mac环境编译IOS 时,cmake版本需要高于cmake 3.15;mac环境上编译Android时,cmake版本需要设置为cmake 3.10。**
ios tiny publish支持的编译选项:
* `--arm_os`: 可选ios或者ios64
* `--arm_abi`: 可选armv7和armv8(**注意**:当`arm_os=ios`时只能选择`arm_abi=armv7`,当`arm_os=ios64`时只能选择`arm_abi=armv8`
* 如果mac编译过程中报错:"Invalid CMAKE_DEVELOPER_ROOT: does not exist", 运行:
```shell
sudo xcode-select -s /Applications/Xcode.app/Contents/Developer
```
##### ARMLinux
```shell
./lite/tools/build.sh \
--build_extra=OFF \
--arm_os=armlinux \
--arm_abi=armv7hf \
--arm_lang=gcc \
--build_extra=OFF \
tiny_publish
```
- `--arm_abi`: 树莓派3b使用armv7hf,RK3399使用armv8
#### 编译`full publish`动态库
##### Android
```shell
./lite/tools/build.sh \
--arm_os=android \
--arm_abi=armv8 \
--build_extra=OFF \
--arm_lang=gcc \
--android_stl=c++_static \
--build_extra=OFF \
full_publish
```
##### ARMLinux
```shell
./lite/tools/build.sh \
--arm_os=armlinux \
--arm_abi=armv7hf \
--arm_lang=gcc \
--build_extra=OFF \
full_publish
```
- `--arm_abi`: 树莓派3b使用armv7hf,RK3399使用armv8
### 编译结果说明
**编译最终产物位置**`build.lite.xxx.xxx.xxx` 下的 `inference_lite_lib.xxx.xxx` ,如 Android 下 ARMv8 的产物位于`inference_lite_lib.android.armv8`
![](https://user-images.githubusercontent.com/45189361/65375706-204e8780-dccb-11e9-9816-ab4563ce0963.png)
**目录内容**(可能)如下:
**Full_publish编译结果:**
![](https://user-images.githubusercontent.com/45189361/65375704-19c01000-dccb-11e9-9650-6856c7a5bf82.png)
**Tiny_publish结果:**
![](https://user-images.githubusercontent.com/45189361/65375726-3bb99280-dccb-11e9-9903-8ce255371905.png)
**IOS编译结果:**
![](https://user-images.githubusercontent.com/45189361/65375726-3bb99280-dccb-11e9-9903-8ce255371905.png)
**具体内容**说明:
1、 `bin`文件夹:可执行工具文件 `paddle_code_generator``test_model_bin`
2、 `cxx`文件夹:包含c++的库文件与相应的头文件
- `include` : 头文件
- `lib` : 库文件
- 打包的静态库文件:
- `libpaddle_api_full_bundled.a` :包含 full_api 和 light_api 功能的静态库
- `libpaddle_api_light_bundled.a` :只包含 light_api 功能的静态库
- 打包的动态态库文件:
- `libpaddle_full_api_shared.so` :包含 full_api 和 light_api 功能的动态库
- `libpaddle_light_api_shared.so`:只包含 light_api 功能的动态库
3、 `demo`文件夹:示例 demo ,包含 C++ demo 和 Java demo。
- `cxx` : C++示例 demo
- `mobile_full` : full_api 的使用示例
- `mobile_light` : light_api的使用示例
- `java` :Java 示例 demo
- `android` : Java的 Android 示例
4、 `java` 文件夹:包含 Jni 的动态库文件与相应的 Jar 包
- `jar` : `PaddlePredictor.jar`
- `so` : Jni动态链接库 `libpaddle_lite_jni.so`
5、 `third_party` 文件夹:第三方库文件`gflags`
**注意:**
1、 只有当`--arm_os=android` 时才会编译出:
- Java库文件与示例:`Java``demo/java`
- 动态库文件:`libpaddle_full_api_shared.so`,`libpaddle_light_api_shared.so`
2、 `tiny_publish`编译结果不包括 C++ demo和 C++ 静态库,但提供 C++ 的 light_api 动态库、 Jni 动态库和Java demo
### 加速第三方依赖库的下载
移动端相关编译所需的第三方库均位于 `<PaddleLite>/third-party` 目录下,默认编译过程中,会利用`git submodule update --init --recursive`链上相关的第三方依赖的仓库。
为加速`full_publish``test`编译模式中对`protobuf`等第三方依赖的下载,`build.sh``ci_build.sh`支持了从国内 CDN 下载第三方依赖的压缩包。
使用方法:`git clone``Paddle-Lite`仓库代码后,手动删除本地仓库根目录下的`third-party`目录:
```shell
git clone https://github.com/PaddlePaddle/Paddle-Lite.git
git checkout <release-version-tag>
cd Paddle-Lite
rm -rf third-party
```
之后再根据本文档,进行后续编译时,便会忽略第三方依赖对应的`submodule`,改为下载第三方压缩包。
# 支持硬件列表
## ARM CPU
Paddle Lite支持[ARM Cortex-A系列处理器](https://en.wikipedia.org/wiki/ARM_Cortex-A),支持列表如下:
### 32bit(ARMv7a)
- Cortex-A5
- Cortex-A7
- Cortex-A8
- Cortex-A9
- Cortex-A12
- Cortex-A15
- Cortex-A17(RK3288)
- Cortex-A32
### 64bit(ARMv7a, ARMv8a)
- Cortex-A35
- Cortex-A53(树莓派3)
- Cortex-A55
- Cortex-A57(Nvidia tx1,Nvidia tx2, 高通810等)
- Cortex-A72(麒麟95X,高通820, RK3399,树莓派4等)
- Cortex-A73(麒麟960,麒麟970,高通835, 联发科X30等)
- Cortex-A75(高通845等)
- Cortex-A76(麒麟980,麒麟990,高通855,高通730,联发科G90等)
- Cortex-A77
- ARMv8-A compatible(Apple A系列处理器, Nvidia tegra, Qualcomm Kryo, Falkor, Samsung Mongoose)
## 移动端GPU
Paddle Lite支持移动端GPU和Nvidia端上GPU设备,支持列表如下:
- ARM Mali G 系列
- Qualcomm Adreno 系列
- Nvida tegra系列: tx1, tx2, nano, xavier
# C++ Demo
## 编译
首先按照[PaddleLite 源码编译](https://github.com/PaddlePaddle/Paddle-Lite/wiki/source_compile)准备交叉编译环境,之后拉取最新[PaddleLite release发布版代码](https://github.com/PaddlePaddle/Paddle-Lite)。下面以Android-ARMv8架构为例,介绍编译过程,并最终在手机上跑通MobilNetv1模型。
进入 Paddle-Lite 目录,运行以下命令编译代码(**需加编译选项`--build_extra=ON`确保完整编译**):
```
./lite/tools/build.sh \
--arm_os=android \
--arm_abi=armv8 \
--arm_lang=gcc \
--android_stl=c++_static \
--build_extra=ON \
full_publish
```
编译完成后 `./build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/` 文件夹下包含:
- cxx
- include (头文件文件夹)
- lib (库文件文件夹)
- libpaddle_api_full_bundled.a
- libpaddle_api_light_bundled.a
- libpaddle_light_api_shared.so
- libpaddle_full_api_shared.so
- demo
- cxx (C++ demo)
- mobile_light (light api demo)
- mobile_full (full api demo)
- mobile_detection (detection model api demo)
- mobile_classify (classify model api demo)
- Makefile.def
- include
- third_party (第三方库文件夹)
- gflags
## 准备执行环境
执行环境有两种:使用安卓手机;若没安卓手机,也可在安卓模拟器中执行。
### 环境一:使用安卓手机
将手机连上电脑,在手机上打开选项 -> 开启-开发者模式 -> 开启-USB调试模式。确保 `adb devices` 能够看到相应的设备。
### 环境二:使用安卓模拟器
运行下面命令,分别创建安卓armv8、armv7架构的模拟器。若需在真机测试,将模拟器换成相应架构的真机环境即可。
```
*android-armv8*
adb kill-server
adb devices | grep emulator | cut -f1 | while read line; do adb -s $line emu kill; done
echo n | avdmanager create avd -f -n paddle-armv8 -k "system-images;android-24;google_apis;arm64-v8a"
echo -ne '\n' | ${ANDROID_HOME}/emulator/emulator -avd paddle-armv8 -noaudio -no-window -gpu off -port 5554 &
sleep 1m
```
```
*android-armv7*
adb kill-server
adb devices | grep emulator | cut -f1 | while read line; do adb -s $line emu kill; done
echo n | avdmanager create avd -f -n paddle-armv7 -k "system-images;android-24;google_apis;armeabi-v7a"
echo -ne '\n' | ${ANDROID_HOME}/emulator/emulator -avd paddle-armv7 -noaudio -no-window -gpu off -port 5554 &
sleep 1m
```
## 下载模型并运行示例
```
cd inference_lite_lib.android.armv8/demo/cxx/mobile_full
wget http://paddle-inference-dist.bj.bcebos.com/mobilenet_v1.tar.gz
tar zxvf mobilenet_v1.tar.gz
make
adb push mobilenet_v1 /data/local/tmp/
adb push mobilenetv1_full_api /data/local/tmp/
adb shell chmod +x /data/local/tmp/mobilenetv1_full_api
adb shell "/data/local/tmp/mobilenetv1_full_api --model_dir=/data/local/tmp/mobilenet_v1 --optimized_model_dir=/data/local/tmp/mobilenet_v1.opt"
```
注:我们也提供了轻量级 API 的 demo、图像分类demo和目标检测demo,支持图像输入;
### Light API Demo
```
cd ../mobile_light
make
adb push mobilenetv1_light_api /data/local/tmp/
adb shell chmod +x /data/local/tmp/mobilenetv1_light_api
adb shell "/data/local/tmp/mobilenetv1_light_api --model_dir=/data/local/tmp/mobilenet_v1.opt "
```
### 图像分类 Demo
```
cd ../mobile_classify
wget http://paddle-inference-dist.bj.bcebos.com/mobilenet_v1.tar.gz
tar zxvf mobilenet_v1.tar.gz
make
adb push mobile_classify /data/local/tmp/
adb push test.jpg /data/local/tmp/
adb push labels.txt /data/local/tmp/
adb push ../../../cxx/lib/libpaddle_light_api_shared.so /data/local/tmp/
adb shell chmod +x /data/local/tmp/mobile_classify
adb shell "export LD_LIBRARY_PATH=/data/local/tmp/:$LD_LIBRARY_PATH && /data/local/tmp/mobile_classify /data/local/tmp/mobilenet_v1.opt /data/local/tmp/test.jpg /data/local/tmp/labels.txt"
```
### 目标检测 Demo
```
cd ../mobile_detection
wget https://paddle-inference-dist.bj.bcebos.com/mobilenetv1-ssd.tar.gz
tar zxvf mobilenetv1-ssd.tar.gz
make
adb push mobile_detection /data/local/tmp/
adb push test.jpg /data/local/tmp/
adb push ../../../cxx/lib/libpaddle_light_api_shared.so /data/local/tmp/
adb shell chmod +x /data/local/tmp/mobile_detection
adb shell "export LD_LIBRARY_PATH=/data/local/tmp/:$LD_LIBRARY_PATH && /data/local/tmp/mobile_detection /data/local/tmp/mobilenetv1-ssd /data/local/tmp/test.jpg"
adb pull /data/local/tmp/test_detection_result.jpg ./
```
## Demo 程序运行结果
### light API Demo 运行结果
运行成功后 ,将在控制台输出预测结果的前10个类别的预测概率:
```
Output dim: 1000
Output[0]: 0.000191
Output[100]: 0.000160
Output[200]: 0.000264
Output[300]: 0.000211
Output[400]: 0.001032
Output[500]: 0.000110
Output[600]: 0.004829
Output[700]: 0.001845
Output[800]: 0.000202
Output[900]: 0.000586
```
### 图像分类 Demo 运行结果
运行成功后 ,将在控制台输出预测结果的前5个类别的类型索引、名字和预测概率:
```
parameter: model_dir, image_path and label_file are necessary
parameter: topk, input_width, input_height, are optional
i: 0, index: 285, name: Egyptian cat, score: 0.482870
i: 1, index: 281, name: tabby, tabby cat, score: 0.471593
i: 2, index: 282, name: tiger cat, score: 0.039779
i: 3, index: 287, name: lynx, catamount, score: 0.002430
i: 4, index: 722, name: ping-pong ball, score: 0.000508
```
### 目标检测 Demo 运行结果
运行成功后 ,将在控制台输出检测目标的类型、预测概率和坐标:
```
running result:
detection image size: 935, 1241, detect object: person, score: 0.996098, location: x=187, y=43, width=540, height=592
detection image size: 935, 1241, detect object: person, score: 0.935293, location: x=123, y=639, width=579, height=597
```
## 如何在代码中使用 API
在C++中使用PaddleLite API非常简单,不需要添加太多额外代码,具体步骤如下:
- 加入头文件引用
```
#include <iostream>
#include <vector>
#include "paddle_api.h"
#include "paddle_use_kernels.h"
#include "paddle_use_ops.h"
#include "paddle_use_passes.h"
```
- 通过MobileConfig设置:模型文件位置(model_dir)、线程数(thread)和能耗模式( power mode )。输入数据(input),从 MobileConfig 创建 PaddlePredictor 并执行预测。 (注:Lite还支持从memory直接加载模型,可以通过MobileConfig::set_model_buffer方法实现)
代码示例:
```
// 1. Create MobileConfig
MobileConfig config;
// 2. Load model
config.set_model_dir("path to your model directory"); // model dir
/*load model: Lite supports loading model from file or from memory (naive buffer from optimized model)
//Method One: Load model from memory:
void set_model_buffer(const char* model_buffer,
size_t model_buffer_size,
const char* param_buffer,
size_t param_buffer_size)
//Method Two: Load model from file:
void set_model_dir(const std::string& model_dir) */
// 3. Set MobileConfig (or you can skip this step to use default value):
config.set_power_mode(LITE_POWER_HIGH); // power mode
/*power modes: Lite supports the following power modes
LITE_POWER_HIGH
LITE_POWER_LOW
LITE_POWER_FULL
LITE_POWER_NO_BIND
LITE_POWER_RAND_HIGH
LITE_POWER_RAND_LOW */
config.set_threads("num of threads"); // threads
// 4. Create PaddlePredictor by MobileConfig
std::shared_ptr<PaddlePredictor> predictor =
CreatePaddlePredictor<MobileConfig>(config);
// 5. Prepare input data
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
input_tensor->Resize({1, 3, 224, 224});
auto *data = input_tensor -> mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// 6. Run predictor
predictor->Run();
// 7. Get output
std::unique_ptr<const Tensor> output_tensor(std::move(predictor->GetOutput(0)));
```
## CxxConfig案例: OCR_model的运行
1. OCR 模型文件:
- 我们提供Pb格式的[ocr_attention_mode](https://paddle-inference-dist.cdn.bcebos.com/ocr_attention.tar.gz)l下载
- 也可以从[Paddle/model项目](https://github.com/PaddlePaddle/models/tree/develop/PaddleCV/ocr_recognition)中训练出模型
2. 示例代码:
```
#include "paddle_api.h" // NOLINT
#include "paddle_use_passes.h" // NOLINT
#include <gflags/gflags.h>
#include <stdio.h>
#include <vector>
using namespace paddle::lite_api; // NOLINT
DEFINE_string(model_dir, "", "Model dir path.");
DEFINE_bool(prefer_int8_kernel, false, "Prefer to run model with int8 kernels");
int64_t ShapeProduction(const shape_t &shape) {
int64_t res = 1;
for (auto i : shape)
res *= i;
return res;
}
void RunModel() {
// 1. Set CxxConfig
CxxConfig config;
config.set_model_dir(FLAGS_model_dir);
std::vector<Place> valid_places({Place{TARGET(kARM), PRECISION(kFloat)}});
if (FLAGS_prefer_int8_kernel) {
valid_places.insert(valid_places.begin(),
Place{TARGET(kARM), PRECISION(kInt8)});
}
config.set_valid_places(valid_places);
// 2. Create PaddlePredictor by CxxConfig
std::shared_ptr<PaddlePredictor> predictor =
CreatePaddlePredictor<CxxConfig>(config);
// 3. Prepare input data
// input 0
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
input_tensor->Resize(shape_t({1, 1, 48, 512}));
auto *data = input_tensor->mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// input1
std::unique_ptr<Tensor> init_ids(std::move(predictor->GetInput(1)));
init_ids->Resize(shape_t({1, 1}));
auto *data_ids = init_ids->mutable_data<float>();
for (int i = 0; i < ShapeProduction(init_ids->shape()); ++i) {
data_ids[i] = 0;
}
lod_t lod_i;
lod_i.push_back({0, 1});
lod_i.push_back({0, 1});
init_ids->SetLoD(lod_i);
// input2
std::unique_ptr<Tensor> init_scores(std::move(predictor->GetInput(2)));
init_scores->Resize(shape_t({1, 1}));
auto *data_scores = init_scores->mutable_data<float>();
for (int i = 0; i < ShapeProduction(init_scores->shape()); ++i) {
data_scores[i] = 0;
}
lod_t lod_s;
lod_s.push_back({0, 1});
lod_s.push_back({0, 1});
init_scores->SetLoD(lod_s);
// 4. Run predictor
predictor->Run();
// 5. Get output
std::unique_ptr<const Tensor> output_tensor(
std::move(predictor->GetOutput(0)));
for (int i = 0; i < ShapeProduction(output_tensor->shape()); i++) {
printf("Output[%d]: %f\n", i, output_tensor->data<float>()[i]);
}
}
int main(int argc, char **argv) {
google::ParseCommandLineFlags(&argc, &argv, true);
RunModel();
return 0;
}
```
3. 运行方法:
参考以上代码编译出可执行文件`OCR_DEMO`,模型文件夹为`ocr_attention`。手机以USB调试、文件传输模式连接电脑。
```
简单编译出`OCR_DEMO`的方法:用以上示例代码替换编译结果中`build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/demo/cxx/mobile_full/mobilenetv1_full_api.cc`文件的内容,终端进入该路径(`demo/cxx/mobile_full/`),终端中执行`make && mv mobilenetv1_full_api OCR_DEMO`即编译出了OCR模型的可执行文件`OCR_DEMO`
```
在终端中输入以下命令执行OCR model测试:
```
#OCR_DEMO为编译出的可执行文件名称;ocr_attention为ocr_attention模型的文件夹名称;libpaddle_full_api_shared.so是编译出的动态库文件,位于`build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/cxx/lib`
adb push OCR_DEMO /data/local/tmp
adb push ocr_attention /data/local/tmp
adb push libpaddle_full_api_shared.so /data/local/tmp/
adb shell 'export LD_LIBRARY_PATH=/data/local/tmp/:$LD_LIBRARY_PATH && cd /data/local/tmp && ./OCR_DEMO --model_dir=./OCR_DEMO'
```
4. 运行结果
<img src='https://user-images.githubusercontent.com/45189361/64398400-46531580-d097-11e9-9f1c-5aba1dfbc24f.png' align='left' width="150" height="200"/>
# Lite基于FPGA的模型预测
Paddle Lite支持基于arm的FPGA zu3/zu5/zu9的模型预测,提供armv8的交叉编译
Lite基于FPGA运行模型需要相应的FPGA驱动,目前只支持百度[Edgeboard开发板](https://ai.baidu.com/tech/hardware/deepkit)
## Lite实现FPGA简介
Lite支持FPGA作为后端硬件进行模型推理,其主要特性如下:
- Lite中FPGA的kernel(feed、fetch除外)均以FP16、NHWC的格式作为输入输出格式,所有的weights和bias仍为FP32、NCHW的格式,feed的输入和fetch的输出均为FP32、NCHW格式的数据,在提升计算速度的同时能做到用户对数据格式无感知
- 对于FPGA暂不支持的kernel,均会切回arm端运行,实现arm+FPGA混合布署运行
- 目前FPGA成本功耗都较低,Lite基于FPGA的模型性能远远好于arm端,可作为边缘设备首选硬件
## 编译
需要提前准备带有FPGAdrv.ko的FPGA开发板(如edgeboard开发板)和Lite代码
CMAKE编译选项:
- 设置`LITE_WITH_FPGA=ON``LITE_WITH_ARM=ON`
其他编译选项与ARM编译相同,可以参考[“Paddle Lite在Docker下的ARM编译”](../source_compile)
示例如下:
```shell
cmake .. \
-DWITH_GPU=OFF \
-DWITH_MKL=OFF \
-DWITH_LITE=ON \
-DLITE_WITH_CUDA=OFF \
-DLITE_WITH_X86=OFF \
-DLITE_WITH_ARM=ON \
-DLITE_WITH_OPENMP=ON \
-DLITE_WITH_LIGHT_WEIGHT_FRAMEWORK=ON \
-DWITH_TESTING=OFF \
-DLITE_WITH_FPGA=ON \
-DARM_TARGET_OS=armlinux
make publish_inference -j2
```
Lite提供FPGA编译脚本,位于lite/tools/build_FPGA.sh,在Lite根目录执行该脚本即可编译
## 运行示例
- **运行文件准备**
下面以Resnet50模型为例,介绍如何使用edgeboard开发板实现模型运行
```bash
#连接开发板,并利用screen命令启动 [本机执行]
screen /dev/cu.SLAB_USBtoUART 115200
#查看开发板ip并ssh登录到开发板,假设开发板ip为192.0.1.1 [本机执行]
ssh root@192.0.1.1
#在开发板上建立目录workspace,拷贝FPGA驱动FPGAdrv.ko到workspace目录 [开发板执行]
mkdir workspace && scp $DRIVER_PATH/FPGAdrv.ko workspace
#将Lite中编译好的测试程序拷贝到开发板workspace目录 [本机执行]
scp $LITE_ROOT/build_FPGA/lite/api/test_resnet50_FPGA root@$EDGEBOARD_IP:workspace/
#把Resnet50的模型和参数scp到开发板workspace目录 [本机执行]
scp -r $LITE_ROOT/build_FPGA/lite/third_party/install/resnet50/ root@$EDGEBOARD_IP:workspace/
#在运行模型前需要加载FPGA驱动 [开发板执行]
insmod FPGAdrv.ko
#给测试程序添加可运行权限 [开发板执行]
chmod +x test_resnet50_FPGA
```
- **使用FPGA进行模型预测**
```bash
#以下命令均在开发板上运行
#直接运行单测程序
./test_resnet50_FPGA --model_dir=resnet50
#如果需要测试性能,可以用repeats参数设置模型运行次数(如1000),同时可以设置预热次数(如10)来让硬件事先运行到稳定水平
./test_resnet50_FPGA --model_dir=resnet50 --repeats=1000 --warmup=10
```
## 如何在Code中使用
在Lite中使用FPGA与ARM相似,具体的区别如下:
- 由于fpga运行模式为fp16精度、nhwc布局,所以需要修改相应的`valid_place`
- fpga不需要device的初始化和运行模式设置
代码示例:
```cpp
lite::Predictor predictor;
std::vector<Place> valid_places(
{Place{TARGET(kFPGA), PRECISION(kFP16), DATALAYOUT(kNHWC)},Place{TARGET(kARM)});
predictor.Build(model_dir, "", "", valid_places);
auto* input_tensor = predictor.GetInput(0);
input_tensor->Resize(DDim(std::vector<DDim::value_type>({1, 3, 224, 224})));
auto* data = input_tensor->mutable_data<float>();
auto item_size = input_tensor->dims().production();
//假设设置输入数据全为1
for (int i = 0; i < item_size; i++) {
data[i] = 1;
}
predictor.Run();
auto* out = predictor.GetOutput(0);
```
# Java Demo
本节中,Java demo 完整代码位于 [demo/java](https://github.com/PaddlePaddle/Paddle-Lite/tree/develop/lite/demo/java)
要编译和跑起Android demo 程序 PaddlePredictor,你需要准备:
1. 一台能运行安卓程序的安卓手机
2. 一台带有AndroidStudio的开发机
## 编译
首先在PaddleLite的开发 [Docker镜像](../source_compile) 中,拉取最新PaddleLite代码,编译对应你手机架构的预测库,
下面我们以arm8 架构举例。进入paddlelite 目录,运行以下命令:
```shell
./lite/tools/build.sh \
--arm_os=android \
--arm_abi=armv8 \
--arm_lang=gcc \
--android_stl=c++_static \
tiny_publish
```
命令完成后查看要存在
```
./build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/java/so/libpaddle_lite_jni.so
./build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/java/jar/PaddlePredictor.jar
./build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/demo/java/android
```
libpaddle_lite_jni.so为 PaddleLite c++ 动态链接库,PaddlePredictor.jar为 Java jar 包,两者包含 PaddleLite Java API,接下来 Android Java 代码会使用这些api。android文件夹中则是Android demo。
## 准备 demo 需要的其他文件
Demo 除了代码,还需要准备在Android工程目录下配置好JNI .so 库(上节提到的`libpaddle_lite_jni.so`),Java .jar 包(上文提到的`PaddlePredictor.jar` ),和模型文件。我们提供了自动化的脚本和手动拷贝两种方法,用户可以根据自己需要选择:
### 脚本方法
进入 `build.lite.android.armv8.gcc/inference_lite_lib.android.armv8/demo/java/android`,我们准备了一个脚本`prepare_demo.bash`,脚本输入一个参数,为你要拷贝的.so 对应的架构文件夹名。
例如运行
```
bash prepare_demo.bash arm8
```
该脚本自动下载并解压缩模型文件,拷贝了 .jar 包进demo,还有生成的.so包进`PaddlePredictor/app/src/main/jinLibs/架构文件夹下`
在我们这个例子里,armv8 就是架构文件夹。备注:这种方式构建的 demo 在 armv8 手机运行正常。如果要demo 程序在别的手机架构(如 armv7)上也运行正常,需要添加别的架构。
### 手动拷贝方法
接下来我们介绍手动拷贝,如果使用了脚本,那么可以跳过以下手动方法的介绍。
### 把 .so 动态库和 .jar 拷贝进安卓demo程序:
1. 将PaddlePredictor 载入到AndroidStudio。
2.`libpaddle_lite_jni.so`拷贝进 `PaddlePredictor/app/src/main/jinLibs/架构文件夹下` ,比如文件夹arm8里要包含该 .so文件。
3.`PaddlePredictor.jar` 拷贝进 `PaddlePredictor/app/libs`
### 把demo使用到的模型文件拷贝进安卓程序:
下载我们的5个模型文件,并解压缩到 `PaddlePredictor/app/src/main/assets` 这个文件夹中
需要拷贝的模型文件和下载地址:
```
inception_v4_simple_opt.nb http://paddle-inference-dist.bj.bcebos.com/inception_v4_simple_opt.nb.tar.gz
lite_naive_model_opt.nb http://paddle-inference-dist.bj.bcebos.com/lite_naive_model_opt.nb.tar.gz
mobilenet_v1_opt.nb http://paddle-inference-dist.bj.bcebos.com/mobilenet_v1_opt.nb.tar.gz
mobilenet_v2_relu_opt.nb http://paddle-inference-dist.bj.bcebos.com/mobilenet_v2_relu_opt.nb.tar.gz
resnet50_opt.nb http://paddle-inference-dist.bj.bcebos.com/resnet50_opt.nb.tar.gz
```
下载完后,assets文件夹里要包含解压后的上面五个模型文件夹,但demo里不需要保存原压缩.tar.gz 文件。
注意:输入的模型要求为naive buffer存储格式,您可以通过 [**Model Optimize Tool**](../model_optimize_tool) 将fluid模型转为naive buffer存储格式。
## 运行 Android 程序结果
以上准备工作完成,就可以开始Build 、安装、和运行安卓demo程序。当你运行PaddlePredictor 程序时,大概会等10秒,然后看到类似以下字样:
```
lite_naive_model output: 50.213173, -28.872887
expected: 50.2132, -28.8729
inception_v4_simple test:true
time: xxx ms
resnet50 test:true
time: xxx ms
mobilenet_v1 test:true
time: xxx ms
mobilenet_v2 test:true
time: xxx ms
```
该 demo 程序跑我们的 5 个模型,第一个模型结果将真正的头两个数字输出,并在第二行附上期望的正确值。你应该要看到他们的误差小于0.001。后面四个模型如果你看到 `test:true` 字样,说明模型输出通过了我们在 demo 程序里对其输出的测试。time 代表该测试花费的时间。
# 裁剪预测库方法
Paddle-Lite支持**根据模型裁剪预测库**功能。Paddle-Lite的一般编译会将所有已注册的operator打包到预测库中,造成库文件体积膨胀;**裁剪预测库**能针对具体的模型,只打包优化后该模型需要的operator,有效降低预测库文件大小。
## 效果展示(Tiny_publish Android动态预测库体积)
| 测试模型 | 裁剪开关 | **libpaddle_lite_jni.so** |转化后模型中的OP|
| ------------------ | ---------------------------- | -------- |------------------|
| mobilenetv1(armv8) | 裁剪前--build_tailor=OFF | 1.5M | feed,etch,conv2d,depthwise_conv2d,fc,fpool2d,softmax |
| mobilenetv1(armv8) | 裁剪后--build_tailor=ON | 788K |feed,etch,conv2d,depthwise_conv2d,fc,fpool2d,softmax|
| mobilenetv2(armv8) | 裁剪前--build_tailor=OFF | 1.5M | feed,fetch,conv2d,depthwise_conv2d,elementwise_add,fc,pool2d,relu6,softmax |
| mobilenetv2(armv8) | 裁剪后--build_tailor=ON | 912K |feed,fetch,conv2d,depthwise_conv2d,elementwise_add,fc,pool2d,relu6,softmax|
| mobilenetv1(armv7) | 裁剪前--build_tailor=OFF | 938K |feed,fetch,concat,conv2d,dropout,fc,pool2d,softmax|
| mobilenetv1(armv7) | 裁剪后--build_tailor=ON | 607K |feed,fetch,concat,conv2d,dropout,fc,pool2d,softmax|
| mobilenetv2(armv7) | 裁剪前--build_tailor=OFF | 938K | feed,fetch,conv2d,depthwise_conv2d,elementwise_add,fc,pool2d,relu6,softmax |
| mobilenetv2(armv7) | 裁剪后--build_tailor=ON |687K |feed,fetch,conv2d,depthwise_conv2d,elementwise_add,fc,pool2d,relu6,softmax|
## 实现过程:
### 1、转化模型时记录优化后模型信息
说明:使用model_optimize_tool转化模型时,选择 `--record_tailoring_info =true` 会将优化后模型的OP和kernel信息保存到输出文件夹,这些信息将用于编译裁剪后的动态库。
注意:需要使用Paddle-Lite 最新版本(release/v2.0.0之后)代码编译出的model_optimize_tool
例如:
```bash
./model_optimize_tool --model_dir=./mobilenet_v1 --optimize_out_type=naive_buffer --optimize_out=mobilenet_v1NB --record_tailoring_info =true --valid_targets=arm
```
效果:优化后模型使用的OP和kernel信息被保存在 `mobilenet_v1NB`文件夹中的隐藏文件里了
### 2、根据模型信息编译裁剪后的预测库
说明:编译Paddle-Lite时选择`--build_tailor=ON` ,并且用 `–-opt_model_dir=` 指定优化后的模型的地址
例如:
```bash
./lite/tools/build.sh --arm_os=android --arm_abi=armv7 --arm_lang=gcc --android_stl=c++_static --build_extra=ON --build_tailor=ON --opt_model_dir=../mobilenet_v1NB full_publish
```
**注意**:上面命令中的`../mobilenet_v1NB`是第1步得到的转化模型的输出路径
**效果**:编译出来的动态库文件变小,且可以运行优化后的模型。
编译出的C++预测库文件位于 :
`build.lite.android.armv7.gcc/inference_lite_lib.android.armv7/cxx/lib/`
编译出的Java预测库文件位于:
`build.lite.android.armv7.gcc/inference_lite_lib.android.armv7/java/so/`
### 3、运行裁剪后的预测库文件
注意:基于某一模型裁剪出的预测库只能支持优化工具转化后的该模型,例如根据mobilenetV1裁剪出的 full_api预测库只能运行以protobuf格式转化出的模型mobilenetV1_opt_nb, 裁剪出的light_api预测库只能运行以naive_buffer格式转化出的模型mobilenetV1_opt_nb, 运行其他模型可能会出现`segementation fault:undifined op or kernel`。 模型转化方法参考:[使用opt转化模型](./model_optimize_tool))。
**示例1**:使用裁剪后的light_api预测库运行mobilenetv1
1、执行第二步编译后,light_api的C++ 示例位于
`/Paddle-Lite/build.lite.android.armv7.gcc/inference_lite_lib.android.armv7/demo/cxx/mobile_light`
输入`make`命令执行编译可编译出可执行文件mobilenetv1_light_api
2、使用adb将mobilenetV1_NB模型和mobilenetv1_light_api传到手机后执行demo:
`./mobilenetv1_light_api --model_dir=./mobilenetV1_NB`
注意:`mobilenetV1_NB`是用`mobilenetV1`模型转化的naive_buffer格式模型(不需要设置` --record_tailoring_info =true`,转化流程参考:[使用opt转化模型](./model_optimize_tool))。
**示例2**:使用裁剪后的full_api预测库运行mobilenetv1
1、执行第二步编译后,full_api的C++ 示例位于
`/Paddle-Lite/build.lite.android.armv7.gcc/inference_lite_lib.android.armv7/demo/cxx/mobile_light`
替换mobilenetv1_full_api.cc代码内容:
```C++
#include <gflags/gflags.h>
#include <stdio.h>
#include <vector>
#include "paddle_api.h" // NOLINT
#include "paddle_use_kernels.h" // NOLINT
#include "paddle_use_ops.h" // NOLINT
#include "paddle_use_passes.h" // NOLINT
using namespace paddle::lite_api; // NOLINT
DEFINE_string(model_dir, "", "Model dir path.");
int64_t ShapeProduction(const shape_t& shape) {
int64_t res = 1;
for (auto i : shape) res *= i;
return res;
}
void RunModel() {
// 1. Set CxxConfig
CxxConfig config;
config.set_model_file(FLAGS_model_dir + "model");
config.set_param_file(FLAGS_model_dir + "params");
std::vector<Place> valid_places{Place{TARGET(kARM), PRECISION(kFloat)}};
config.set_valid_places(valid_places);
// 2. Create PaddlePredictor by CxxConfig
std::shared_ptr<PaddlePredictor> predictor =
CreatePaddlePredictor<CxxConfig>(config);
// 3. Prepare input data
std::unique_ptr<Tensor> input_tensor(std::move(predictor->GetInput(0)));
input_tensor->Resize(shape_t({1, 3, 224, 224}));
auto* data = input_tensor->mutable_data<float>();
for (int i = 0; i < ShapeProduction(input_tensor->shape()); ++i) {
data[i] = 1;
}
// 4. Run predictor
predictor->Run();
// 5. Get output
std::unique_ptr<const Tensor> output_tensor(
std::move(predictor->GetOutput(0)));
printf("Output dim: %d\n", output_tensor->shape()[1]);
for (int i = 0; i < ShapeProduction(output_tensor->shape()); i += 100) {
printf("Output[%d]: %f\n", i, output_tensor->data<float>()[i]);
}
}
int main(int argc, char** argv) {
google::ParseCommandLineFlags(&argc, &argv, true);
RunModel();
return 0;
}
```
2、使用adb将mobilenetV1_PB模型和mobilenetv1_full_api传到手机后执行demo:
`./mobilenetv1_full_api --model_dir=./mobilenetV1_PB`
注意:`mobilenetV1_PB`是用`mobilenetV1`模型转化的protobuf格式模型(不需要设置` --record_tailoring_info =true`,转化流程参考:[使用opt转化模型](./model_optimize_tool))。
## 按模型集合裁剪预测库
为了方便用户使用,我们同时提供了按模型集合进行预测库裁剪的功能。用户可以提供一个模型集合,Model Optimize Tool会根据用户所指定的模型集合分析其**优化后的**模型所需要的算子信息对预测库进行裁剪。使用此功能用户根据自己的需要使用模型集合来对预测库中的算子进行任意裁剪。
使用方法如下所示:
```shell
# 非combined模型集合
./model_optimize_tool \
--model_set_dir=<your_model_set_dir> \
--optimize_out_type=naive_buffer \
--optimize_out=<output_model_set_dir> \
--record_tailoring_info=true \
--valid_targets=arm
# combined模型集合
./model_optimize_tool \
--model_set_dir=<your_model_set_dir> \
--optimize_out_type=naive_buffer \
--model_filename=<model_topo_filename> \
--param_filename=<model_param_filename> \
--optimize_out=<output_model_set_dir> \
--record_tailoring_info=true \
--valid_targets=arm
```
经过以上步骤后会在`<output_model_set_dir>`中生成模型集合中各模型对应的NaiveBuffer格式的优化模型。此步会对模型集合中所需算子信息进行搜集并存储到`<output_model_set_dir>`中。下一步编译预测库的流程与使用单模型进行预测库裁剪步骤相同。
**注意:**
1. 模型集合**必须**均为combined参数模型或均为非combined参数模型。
2. 使用非combined参数模型时,模型拓扑文件名应为`__model__`,使用非combined参数模型时,集合中各模型的拓扑与参数名应相同,分别由`--model_filename``--param_filename`指定。
3. 模型集合**必须**均为INT8量化模型或均为非INT8量化模型。
4. 需要使用Paddle-Lite 最新版本(release/v2.1.0之后)代码编译出的model_optimize_tool。
# 模型转化方法
Lite架构在预测过程中表现出来的高性能得益于其丰富的优化组件,其中包括量化、子图融合、混合调度、Kernel优选等等策略。为了使优化过程更加方便易用,我们提供了**opt**来自动完成优化步骤,输出一个轻量的、最优的可执行模型。具体使用方法介绍如下:
**注意**:release/v2.2.0之前的模型转化工具名称为`model_optimize_tool`,从release/v2.3开始模型转化工具名称修改为`opt`
## 准备opt
当前获得opt方法有三种:
1. 我们提供当前develop分支编译结果下载:[opt](https://paddlelite-data.bj.bcebos.com/model_optimize_tool/opt)[opt_mac](https://paddlelite-data.bj.bcebos.com/model_optimize_tool/opt_mac)
release/v2.2.0之前版本的model_optimize_tool: [model_optimize_tool](https://paddlelite-data.bj.bcebos.com/model_optimize_tool/model_optimize_tool)[model_optimize_tool_mac](https://paddlelite-data.bj.bcebos.com/model_optimize_tool/model_optimize_tool_mac)
2. 可以进入Paddle-Lite Github仓库的[release界面](https://github.com/PaddlePaddle/Paddle-Lite/releases),选择release版本下载对应的转化工具`opt`
(release/v2.2.0之前的转化工具为model_optimize_tool、release/v2.3.0之后为opt)
3. 可以下载Paddle-Lite源码,从源码编译出opt工具
```bash
git clone https://github.com/PaddlePaddle/Paddle-Lite.git
cd Paddle-Lite
git checkout <release-version-tag>
./lite/tools/build.sh build_optimize_tool
```
编译结果位于`Paddle-Lite/build.opt/lite/api/opt`
**注意**:从源码编译opt前需要先[安装Paddle-Lite的开发环境](../installation/source_compile)
## 使用opt
opt是x86平台上的可执行文件,需要在PC端运行:包括Linux终端和Mac终端。
### 帮助信息
执行opt时不加入任何输入选项,会输出帮助信息,提示当前支持的选项:
```bash
./opt
```
![](https://paddlelite-data.bj.bcebos.com/doc_images/1.png)
### 功能一:转化模型为Paddle-Lite格式
opt可以将PaddlePaddle支持的模型转化为Paddle-Lite支持的模型格式,期间执行的操作包括:将protobuf格式的模型文件转化为naive_buffer格式的模型文件,有效降低模型体积;执行“量化、子图融合、混合调度、Kernel优选”等图优化操作,提升其在Paddle-Lite上的运行速度、内存占用等性能指标。
模型优化过程:
(1)准备待优化的PaddlePaddle模型
PaddlePaddle模型有两种保存格式:
Combined Param:所有参数信息保存在单个文件`params`中,模型的拓扑信息保存在`__model__`文件中。
![opt_combined_model](https://paddlelite-data.bj.bcebos.com/doc_images%2Fcombined_model.png)
Seperated Param:参数信息分开保存在多个参数文件中,模型的拓扑信息保存在`__model__`文件中。
![opt_seperated_model](https://paddlelite-data.bj.bcebos.com/doc_images%2Fseperated_model.png)
(2) 终端中执行`opt`优化模型
**使用示例**:转化`mobilenet_v1`模型
```
./opt --model_dir=./mobilenet_v1 --valid_targets=arm --optimize_out_type=naive_buffer --optimize_out=mobilenet_v1_opt
```
以上命令可以将`mobilenet_v1`模型转化为arm硬件平台、naive_buffer格式的Paddle_Lite支持模型,优化后的模型文件为`mobilenet_v1_opt.nb`,转化结果如下图所示:
![opt_resulted_model](https://paddlelite-data.bj.bcebos.com/doc_images/2.png)
(3) **更详尽的转化命令**总结:
```shell
./opt \
--model_dir=<model_param_dir> \
--model_file=<model_path> \
--param_file=<param_path> \
--optimize_out_type=(protobuf|naive_buffer) \
--optimize_out=<output_optimize_model_dir> \
--valid_targets=(arm|opencl|x86|npu|xpu) \
--prefer_int8_kernel=(true|false) \
--record_tailoring_info =(true|false)
```
| 选项 | 说明 |
| ------------------- | ------------------------------------------------------------ |
| --model_dir | 待优化的PaddlePaddle模型(非combined形式)的路径 |
| --model_file | 待优化的PaddlePaddle模型(combined形式)的网络结构文件路径。 |
| --param_file | 待优化的PaddlePaddle模型(combined形式)的权重文件路径。 |
| --optimize_out_type | 输出模型类型,目前支持两种类型:protobuf和naive_buffer,其中naive_buffer是一种更轻量级的序列化/反序列化实现。若您需要在mobile端执行模型预测,请将此选项设置为naive_buffer。默认为protobuf。 |
| --optimize_out | 优化模型的输出路径。 |
| --valid_targets | 指定模型可执行的backend,默认为arm。目前可支持x86、arm、opencl、npu、xpu,可以同时指定多个backend(以空格分隔),Model Optimize Tool将会自动选择最佳方式。如果需要支持华为NPU(Kirin 810/990 Soc搭载的达芬奇架构NPU),应当设置为npu, arm。 |
| --prefer_int8_kernel | 若待优化模型为int8量化模型(如量化训练得到的量化模型),则设置该选项为true以使用int8内核函数进行推理加速,默认为false。 |
| --record_tailoring_info | 当使用 [根据模型裁剪库文件](./library_tailoring.html) 功能时,则设置该选项为true,以记录优化后模型含有的kernel和OP信息,默认为false。 |
* 如果待优化的fluid模型是非combined形式,请设置`--model_dir`,忽略`--model_file``--param_file`
* 如果待优化的fluid模型是combined形式,请设置`--model_file``--param_file`,忽略`--model_dir`
* 优化后的模型包括__model__.nb和param.nb文件。
### 功能二:统计模型算子信息、判断是否支持
opt可以统计并打印出model中的算子信息、判断Paddle-Lite是否支持该模型。并可以打印出当前Paddle-Lite的算子支持情况。
(1)使用opt统计模型中算子信息
下面命令可以打印出mobilenet_v1模型中包含的所有算子,并判断在硬件平台`valid_targets`下Paddle-Lite是否支持该模型
`./opt --print_model_ops=true --model_dir=mobilenet_v1 --valid_targets=arm`
![opt_print_modelops](https://paddlelite-data.bj.bcebos.com/doc_images/3.png)
(2)使用opt打印当前Paddle-Lite支持的算子信息
`./opt --print_all_ops=true`
以上命令可以打印出当前Paddle-Lite支持的所有算子信息,包括OP的数量和每个OP支持哪些硬件平台:
![opt_print_allops](https://paddlelite-data.bj.bcebos.com/doc_images/4.png)
`./opt ----print_supported_ops=true --valid_targets=x86`
以上命令可以打印出当`valid_targets=x86`时Paddle-Lite支持的所有OP:
![opt_print_supportedops](https://paddlelite-data.bj.bcebos.com/doc_images/5.png)
## 其他功能:合并x2paddle和opt的一键脚本
**背景**:如果想用Paddle-Lite运行第三方来源(tensorflow、caffe、onnx)模型,一般需要经过两次转化。即使用x2paddle工具将第三方模型转化为PaddlePaddle格式,再使用opt将PaddlePaddle模型转化为Padde-Lite可支持格式。
为了简化这一过程,我们提供一键脚本,将x2paddle转化和opt转化合并:
**一键转化脚本**[auto_transform.sh](https://paddlelite-data.bj.bcebos.com/model_optimize_tool/auto_transform.sh)
**环境要求**:使用`auto_transform.sh`脚本转化第三方模型时,需要先安装x2paddle环境,请参考[x2paddle环境安装方法](https://github.com/PaddlePaddle/X2Paddle#环境依赖) 安装x2paddle和其环境依赖项。
**使用方法**
(1)打印帮助帮助信息:` ./auto_transform.sh`
(2)转化模型方法
```bash
USAGE:
auto_transform.sh combines the function of x2paddle and opt, it can
tranform model from tensorflow/caffe/onnx form into paddle-lite naive-buffer form.
----------------------------------------
example:
./auto_transform.sh --framework=tensorflow --model=tf_model.pb --optimize_out=opt_model_result
----------------------------------------
Arguments about x2paddle:
--framework=(tensorflow|caffe|onnx);
--model='model file for tensorflow or onnx';
--prototxt='proto file for caffe' --weight='weight file for caffe'
For TensorFlow:
--framework=tensorflow --model=tf_model.pb
For Caffe:
--framework=caffe --prototxt=deploy.prototxt --weight=deploy.caffemodel
For ONNX
--framework=onnx --model=onnx_model.onnx
Arguments about opt:
--valid_targets=(arm|opencl|x86|npu|xpu); valid targets on Paddle-Lite.
--fluid_save_dir='path to outputed model after x2paddle'
--optimize_out='path to outputed Paddle-Lite model'
----------------------------------------
```
# Lite基于OpenCL的ARM GPU预测
Lite支持在Android系统上运行基于OpenCL的程序,目前支持Ubuntu环境下armv8、armv7的交叉编译。
## 编译
### 编译环境
1. Docker 容器环境;
2. Linux(推荐 Ubuntu 16.04)环境。
详见 **源码编译指南-环境准备** 章节。
### 编译选项
|参数|介绍|值|
|--------|--------|--------|
|--arm_os|代表目标操作系统|目前仅支持且默认为`android`|
|--arm_abi|代表体系结构类型,支持armv8和armv7|默认为`armv8`即arm64-v8a;`armv7`即armeabi-v7a|
|--arm_lang|代表编译目标文件所使用的编译器|默认为gcc,支持 gcc和clang两种|
### 编译Paddle-Lite OpenCL库范例
注:以android-armv8-opencl的目标、Docker容器的编译开发环境为例,CMake3.10,android-ndk-r17c位于`/opt/`目录下。
```bash
# 假设当前位于处于Lite源码根目录下
# 导入NDK_ROOT变量,注意检查您的安装目录若与本示例不同
export NDK_ROOT=/opt/android-ndk-r17c
# 删除上一次CMake自动生成的.h文件
rm ./lite/api/paddle_use_kernels.h
rm ./lite/api/paddle_use_ops.h
# 根据指定编译参数编译
./lite/tools/ci_build.sh \
--arm_os=android \
--arm_abi=armv8 \
--arm_lang=gcc \
build_test_arm_opencl
```
编译产物位于`build.lite.android.armv8.gcc.opencl`下的`inference_lite_lib.android.armv8.opencl`文件夹内,这里仅罗列关键产物:
- `cxx`:该目录是编译目标的C++的头文件和库文件;
- `demo`:该目录包含了两个demo,用来调用使用`libpaddle_api_full_bundled.a``libpaddle_api_light_bundled.a`,分别对应`mobile_full``mobile_light`文件夹。编译对应的demo仅需在`mobile_full``mobile_light`
- `mobile_full`:使用cxx config,可直接加载fluid模型,若使用OpenCL需要在`mobilenetv1_full_api.cc`代码里开启`DEMO_USE_OPENCL`的宏,详细见代码注释;
- `mobile_light`:使用mobile config,只能加载`model_optimize_tool`优化过的模型;
- `opencl`:该目录存放opencl实现的相关kernel。
```bash
.
|-- cxx
| |-- include
| | |-- paddle_api.h
| | |-- paddle_image_preprocess.h
| | |-- paddle_lite_factory_helper.h
| | |-- paddle_place.h
| | |-- paddle_use_kernels.h
| | |-- paddle_use_ops.h
| | `-- paddle_use_passes.h
| `-- lib
| |-- libpaddle_api_full_bundled.a
| |-- libpaddle_api_light_bundled.a
| |-- libpaddle_full_api_shared.so
| `-- libpaddle_light_api_shared.so
|-- demo
| `-- cxx
| |-- Makefile.def
| |-- README.md
| |-- include
| | |-- paddle_api.h
| | |-- paddle_lite_factory_helper.h
| | |-- paddle_place.h
| | |-- paddle_use_kernels.h
| | |-- paddle_use_ops.h
| | `-- paddle_use_passes.h
| |-- mobile_full
| | |-- Makefile
| | `-- mobilenetv1_full_api.cc
| `-- mobile_light
| |-- Makefile
| `-- mobilenetv1_light_api.cc
`-- opencl
`-- cl_kernel
|-- buffer
| |-- depthwise_conv2d_kernel.cl
| |-- elementwise_add_kernel.cl
| |-- fc_kernel.cl
| |-- im2col_kernel.cl
| |-- layout_kernel.cl
| |-- mat_mul_kernel.cl
| |-- pool_kernel.cl
| `-- relu_kernel.cl
|-- cl_common.h
`-- image
|-- channel_add_kernel.cl
|-- elementwise_add_kernel.cl
|-- pool_kernel.cl
`-- relu_kernel.cl
```
调用`libpaddle_api_full_bundled.a``libpaddle_api_light_bundled.a`见下一部分运行示例。
## 运行示例
下面以android、ARMv8、gcc的环境为例,介绍3个示例,分别如何在手机上执行基于OpenCL的ARM GPU推理过程。
**注意:** 以下命令均在Lite源码根目录下运行。在3个示例前,下面这段命令都先要执行用来准备环境:
```bash
# 在/data/local/tmp目录下创建OpenCL文件目录
adb shell mkdir -p /data/local/tmp/opencl
adb shell mkdir -p /data/local/tmp/opencl/cl_kernel/buffer
adb shell mkdir -p /data/local/tmp/opencl/cl_kernel/image
# 将OpenCL的kernels文件推送到/data/local/tmp/opencl目录下
adb push lite/backends/opencl/cl_kernel/cl_common.h /data/local/tmp/opencl/cl_kernel/
adb push lite/backends/opencl/cl_kernel/buffer/* /data/local/tmp/opencl/cl_kernel/buffer/
adb push lite/backends/opencl/cl_kernel/image/* /data/local/tmp/opencl/cl_kernel/image/
```
### 运行示例1: 编译产物demo示例
```bash
######################################################################
# 编译mobile_full的demo #
######################################################################
# 步骤: #
# 0.确保编译Paddle-Lite时编译了OpenCL; #
# 1.编辑`mobilenetv1_full_api.cc`代码, 开启`DEMO_USE_OPENCL`的宏; #
# 2.在产物目录`demo/cxx/mobile_full`下编译`mobile_full`的demo; #
# 3.上传demo, 模型, opencl kernel文件到手机; #
# 4.运行demo得到预期结果. #
######################################################################
adb shell mkdir /data/local/tmp/opencl/mobilenet_v1
chmod +x ./build.lite.android.armv8.gcc.opencl/inference_lite_lib.android.armv8.opencl/demo/cxx/mobile_full/mobilenetv1_full_api
adb push ./build.lite.android.armv8.gcc.opencl/inference_lite_lib.android.armv8.opencl/demo/cxx/mobile_full/mobilenetv1_full_api /data/local/tmp/opencl/
adb push ./build.lite.android.armv8.gcc.opencl/install/mobilenet_v1/* /data/local/tmp/opencl/mobilenet_v1
# use mobile_full run mobilenet_v1
# `GLOG_v` is log level
adb shell "export GLOG_v=0; \
/data/local/tmp/opencl/mobilenetv1_full_api \
--model_dir=/data/local/tmp/opencl/mobilenet_v1 \
--optimized_model_dir=/data/local/tmp/opencl/full_api_opt_model"
######################################################################
# 编译mobile_light的demo #
######################################################################
# 步骤: #
# 0.确保编译Paddle-Lite时编译了OpenCL; #
# 1.编译model_optimize_tool并对模型优化, `targets`参数为`opencl`; #
# 2.在产物目录`demo/cxx/mobile_light`下编译`mobile_light`的demo; #
# 3.上传demo, 模型, opencl kernel文件到手机; #
# 4.运行demo得到预期结果. #
######################################################################
# use model_optimize_tool to optimize model
./build.model_optimize_tool/lite/api/model_optimize_tool \
--model_dir=./build.lite.android.armv8.gcc.opencl/install/mobilenet_v1/ \
--optimize_out_type=naive_buffer \
--optimize_out=./build.lite.android.armv8.gcc.opencl/install/mobilenet_v1/ \
--valid_targets=opencl
adb shell mkdir /data/local/tmp/opencl/mobilenet_v1
chmod +x ./build.lite.android.armv8.gcc.opencl/inference_lite_lib.android.armv8.opencl/demo/cxx/mobile_light/mobilenetv1_light_api
adb push ./build.lite.android.armv8.gcc.opencl/inference_lite_lib.android.armv8.opencl/demo/cxx/mobile_light/mobilenetv1_light_api /data/local/tmp/opencl/
adb push ./build.lite.android.armv8.gcc.opencl/install/mobilenet_v1/* /data/local/tmp/opencl/mobilenet_v1
# use mobile_light run mobilenet_v1
adb shell "export GLOG_v=5; \
/data/local/tmp/opencl/mobilenetv1_light_api \
--model_dir=/data/local/tmp/opencl/"
```
### 运行示例2: test_mobilenetv1单元测试
- **运行文件准备**
```bash
# 将mobilenet_v1的模型文件推送到/data/local/tmp/opencl目录下
adb shell mkdir -p /data/local/tmp/opencl/mobilenet_v1
adb push build.lite.android.armv8.gcc.opencl/third_party/install/mobilenet_v1/* /data/local/tmp/opencl/mobilenet_v1/
# 将OpenCL单元测试程序test_mobilenetv1,推送到/data/local/tmp/opencl目录下
adb push build.lite.android.armv8.gcc.opencl/lite/api/test_mobilenetv1 /data/local/tmp/opencl
```
- **执行OpenCL推理过程**
使用如下命令运行OpenCL程序。其中:
- `--cl_path`指定了OpenCL的kernels文件即cl\_kernel所在目录;
- `--modle_dir`指定了模型文件所在目录。
```bash
adb shell chmod +x /data/local/tmp/opencl/test_mobilenetv1
adb shell /data/local/tmp/opencl/test_mobilenetv1 \
--cl_path=/data/local/tmp/opencl \
--model_dir=/data/local/tmp/opencl/mobilenet_v1 \
--warmup=1 \
--repeats=1
```
**注意:** 因为权重参数均会在Op Kernel第一次运行时进行加载,所以第一次的执行时间会略长。一般将warmup的值设为1,repeats值设为多次。
### 运行示例3: test_layout_opencl单元测试
- **运行文件准备**
```bash
# 将OpenCL单元测试程序test_layout_opencl,推送到/data/local/tmp/opencl目录下
adb push build.lite.android.armv8.gcc.opencl/lite/kernels/opencl/test_layout_opencl /data/local/tmp/opencl/
```
OpenCL推理过程**
```bash
adb shell chmod +x /data/local/tmp/opencl/test_layout_opencl
adb shell /data/local/tmp/opencl/test_layout_opencl
```
# 如何在Code中使用
见运行示例1的demo代码:
1. [./lite/demo/cxx/mobile_light/mobilenetv1_light_api.cc](https://github.com/PaddlePaddle/Paddle-Lite/blob/develop/lite/demo/cxx/mobile_light/mobilenetv1_light_api.cc);
2. [./lite/demo/cxx/mobile_full/mobilenetv1_full_api.cc](https://github.com/PaddlePaddle/Paddle-Lite/blob/develop/lite/demo/cxx/mobile_full/mobilenetv1_full_api.cc).
注:这里给出的链接会跳转到线上最新develop分支的代码,很可能与您本地的代码存在差异,建议参考自己本地位于`lite/demo/cxx/`目录的代码,查看如何使用。
**NOTE:** 对OpenCL的支持还在持续开发中。
......@@ -43,6 +43,7 @@ void Predictor::SaveModel(const std::string &dir,
LOG(FATAL) << "Unknown model type";
}
if (record_info) {
MkDirRecur(dir);
SaveOpKernelInfo(dir);
}
}
......
......@@ -25,6 +25,8 @@ void LightPredictor::Build(const std::string& lite_model_file,
} else {
LoadModelNaiveFromFile(lite_model_file, scope_.get(), &cpp_program_desc_);
}
DequantizeWeight();
BuildRuntimeProgram(cpp_program_desc_);
PrepareFeedFetch();
}
......
......@@ -23,6 +23,10 @@
#include "lite/core/op_registry.h"
DEFINE_string(optimized_model, "", "optimized_model");
DEFINE_int32(N, 1, "input_batch");
DEFINE_int32(C, 3, "input_channel");
DEFINE_int32(H, 224, "input_height");
DEFINE_int32(W, 224, "input_width");
namespace paddle {
namespace lite {
......@@ -37,7 +41,8 @@ void TestModel(const std::vector<Place>& valid_places,
predictor.Build(model_dir, "", "", valid_places);
auto* input_tensor = predictor.GetInput(0);
input_tensor->Resize(DDim(std::vector<DDim::value_type>({1, 3, 224, 224})));
input_tensor->Resize(DDim(
std::vector<DDim::value_type>({FLAGS_N, FLAGS_C, FLAGS_H, FLAGS_W})));
auto* data = input_tensor->mutable_data<float>();
auto item_size = input_tensor->dims().production();
for (int i = 0; i < item_size; i++) {
......@@ -58,6 +63,8 @@ void TestModel(const std::vector<Place>& valid_places,
predictor.SaveModel(FLAGS_optimized_model);
}
LOG(INFO) << "input shape(NCHW):" << FLAGS_N << " " << FLAGS_C << " "
<< FLAGS_H << " " << FLAGS_W;
LOG(INFO) << "================== Speed Report ===================";
LOG(INFO) << "Model: " << model_dir << ", threads num " << FLAGS_threads
<< ", warmup: " << FLAGS_warmup << ", repeats: " << FLAGS_repeats
......@@ -123,10 +130,10 @@ TEST(MobileNetV1, test_arm) {
#ifdef LITE_WITH_OPENCL
TEST(MobileNetV1, test_opencl) {
std::vector<Place> valid_places({
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kNCHW)},
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kNHWC)},
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kImageDefault)},
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNCHW)},
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNHWC)},
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kImageDefault)},
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kNCHW)},
TARGET(kARM), // enable kARM CPU kernel when no opencl kernel
});
......
......@@ -23,6 +23,10 @@
#include "lite/core/op_registry.h"
DEFINE_string(optimized_model, "", "optimized_model");
DEFINE_int32(N, 1, "input_batch");
DEFINE_int32(C, 3, "input_channel");
DEFINE_int32(H, 224, "input_height");
DEFINE_int32(W, 224, "input_width");
namespace paddle {
namespace lite {
......@@ -38,7 +42,8 @@ void TestModel(const std::vector<Place>& valid_places,
predictor.Build(model_dir, "", "", valid_places);
auto* input_tensor = predictor.GetInput(0);
input_tensor->Resize(DDim(std::vector<DDim::value_type>({1, 3, 224, 224})));
input_tensor->Resize(DDim(
std::vector<DDim::value_type>({FLAGS_N, FLAGS_C, FLAGS_H, FLAGS_W})));
auto* data = input_tensor->mutable_data<float>();
auto item_size = input_tensor->dims().production();
for (int i = 0; i < item_size; i++) {
......@@ -59,6 +64,8 @@ void TestModel(const std::vector<Place>& valid_places,
predictor.SaveModel(FLAGS_optimized_model);
}
LOG(INFO) << "input shape(NCHW):" << FLAGS_N << " " << FLAGS_C << " "
<< FLAGS_H << " " << FLAGS_W;
LOG(INFO) << "================== Speed Report ===================";
LOG(INFO) << "Model: " << model_dir << ", threads num " << FLAGS_threads
<< ", warmup: " << FLAGS_warmup << ", repeats: " << FLAGS_repeats
......@@ -123,8 +130,11 @@ TEST(MobileNetV2, test_arm) {
#ifdef LITE_WITH_OPENCL
TEST(MobileNetV2, test_opencl) {
std::vector<Place> valid_places({
Place{TARGET(kOpenCL), PRECISION(kFloat)},
Place{TARGET(kARM), PRECISION(kFloat)},
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kImageDefault)},
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNCHW)},
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kImageDefault)},
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kNCHW)},
TARGET(kARM), // enable kARM CPU kernel when no opencl kernel
});
TestModel(valid_places);
......
......@@ -26,9 +26,11 @@
#include "lite/api/paddle_use_ops.h"
#include "lite/api/paddle_use_passes.h"
#include "lite/core/op_registry.h"
#include "lite/core/version.h"
#include "lite/model_parser/compatible_pb.h"
#include "lite/model_parser/pb/program_desc.h"
#include "lite/utils/cp_logging.h"
#include "lite/utils/io.h"
#include "lite/utils/string.h"
#include "supported_kernel_op_info.h" // NOLINT
......@@ -89,13 +91,13 @@ std::vector<Place> ParserValidPlaces() {
valid_places.emplace_back(TARGET(kARM));
} else if (target_repr == "opencl") {
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kNCHW)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFP16), DATALAYOUT(kNHWC)});
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kImageDefault)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNCHW)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kFloat), DATALAYOUT(kNHWC)});
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kImageDefault)});
valid_places.emplace_back(
Place{TARGET(kOpenCL), PRECISION(kAny), DATALAYOUT(kNCHW)});
valid_places.emplace_back(
TARGET(kARM)); // enable kARM CPU kernel when no opencl kernel
} else if (target_repr == "x86") {
......@@ -239,6 +241,7 @@ void PrintOpsInfo(std::set<std::string> valid_ops = {}) {
/// Print help information
void PrintHelpInfo() {
// at least one argument should be inputed
const std::string opt_version = lite::version();
const char help_info[] =
"At least one argument should be inputed. Valid arguments are listed "
"below:\n"
......@@ -260,7 +263,8 @@ void PrintHelpInfo() {
" `--print_model_ops=true --model_dir=<model_param_dir> "
"--valid_targets=(arm|opencl|x86|npu|xpu)`"
" Display operators in the input model\n";
std::cout << help_info << std::endl;
std::cout << "opt version:" << opt_version << std::endl
<< help_info << std::endl;
exit(1);
}
......@@ -397,6 +401,7 @@ void Main() {
return;
}
lite::MkDirRecur(FLAGS_optimize_out);
auto model_dirs = lite::ListDir(FLAGS_model_set_dir, true);
if (model_dirs.size() == 0) {
LOG(FATAL) << "[" << FLAGS_model_set_dir << "] does not contain any model";
......@@ -451,7 +456,9 @@ int main(int argc, char** argv) {
}
google::ParseCommandLineFlags(&argc, &argv, false);
paddle::lite_api::ParseInputCommand();
paddle::lite_api::CheckIfModelSupported();
if (FLAGS_model_set_dir == "") {
paddle::lite_api::CheckIfModelSupported();
}
paddle::lite_api::Main();
return 0;
}
......@@ -922,7 +922,7 @@ void sgemv_trans(const int M,
/* end */ \
"4: \n" /* end */ \
"fmov s1, %w[alpha] \n" /* mov alpha to s1 */ \
"fcmp s8, #0 \n" /* cmp with zero*/ \
"fcmp s8, #0.0 \n" /* cmp with zero*/ \
"bge 5f \n" /* if ge zero */ \
"fmul s8, s8, s1 \n" /* out * alpha */ \
"5: \n" /* leakey relu label */ \
......
......@@ -142,7 +142,7 @@ __kernel void depth_conv2d_3x3(__private const int global_size_dim0,
#endif
#ifdef RELU
output = activation(output);
output = activation_type4(output);
#endif
......@@ -309,8 +309,8 @@ __kernel void depth_conv2d_3x3s1(__private const int ou_ch_blk,
#endif
#ifdef RELU
output[0] = activation(output[0]);
output[1] = activation(output[1]);
output[0] = activation_type4(output[0]);
output[1] = activation_type4(output[1]);
#endif
WRITE_IMG_TYPE(CL_DTYPE_CHAR, output_image, (int2)(ou_x, ou_nh_id), output[0]);
......
......@@ -142,14 +142,13 @@ void StrideScal(const T* a, const T* x, T* y, int n, int stride);
// remain is the product of dimension shapes after the axis dimension
template <typename T>
void Softmax(const T* x, T* y, int n, int bs, int remain = 1) {
std::vector<T> entities(bs);
for (int i = 0; i < bs; ++i) {
entities[i] = x[i * n];
T entity = x[i * n];
for (int c = 1; c < n; ++c) {
entities[i] = x[i * n + c] > entities[i] ? x[i * n + c] : entities[i];
entity = x[i * n + c] > entity ? x[i * n + c] : entity;
}
for (int c = 0; c < n; ++c) {
y[i * n + c] = x[i * n + c] - entities[i];
y[i * n + c] = x[i * n + c] - entity;
}
}
VExp(y, y, n * bs);
......
......@@ -110,11 +110,7 @@ void set_constant(const lite::Context<Target>& context,
lite::Tensor* tensor,
float value) {
TensorSetConstantWithTarget<Target> func(context, tensor, value);
// #ifdef PADDLE_WITH_CUDA
// tensor->target().apply_visitor(func);
// #else
func();
// #endif
}
template <typename T>
......@@ -123,7 +119,7 @@ struct RowwiseAdd<lite::TargetType::kX86, T> {
const lite::Tensor& input,
const lite::Tensor& vector,
lite::Tensor* output) {
auto in_dims = input.dims();
const auto& in_dims = input.dims();
auto size = input.numel() / in_dims[0];
PADDLE_ENFORCE_EQ(vector.numel(), size);
PADDLE_ENFORCE_EQ(output->dims(), in_dims);
......
......@@ -103,12 +103,6 @@ void DeleteDynamicQuantOpFuser::InsertNewNode(SSAGraph* graph,
// obtain values, save values and relink node
int bit_length = quant_node->stmt()->op_info()->GetAttr<int>("bit_length");
int range = ((1 << (bit_length - 1)) - 1);
auto* scope = quant_node->stmt()->op()->scope();
auto* scale_tensor = scope->FindVar(output_scale_node->arg()->name)
->GetMutable<lite::Tensor>();
float scale_value = scale_tensor->data<float>()[0] / range;
auto outlinks = output_act_node->outlinks;
for (auto* quantized_node : outlinks) {
auto* op_desc = quantized_node->stmt()->mutable_op_info();
......
......@@ -86,7 +86,6 @@ class KernelPlaceCorrectPass : public DebugPass {
<< node_name;
VLOG(4) << "-- input arg_name:" << arg_name << " "
<< "-- node name:" << node_name;
auto type = inst.picked_kernel().GetInputDeclType(arg_name);
if (!x_in->AsArg().type) {
need_correct_place &= false;
} else {
......@@ -108,7 +107,6 @@ class KernelPlaceCorrectPass : public DebugPass {
<< node_name << " in Inst "
<< inst.op_type();
VLOG(4) << "-- output arg_name " << arg_name;
auto type = inst.picked_kernel().GetOutputDeclType(arg_name);
if (!x_out->AsArg().type) {
need_correct_place &= false;
} else {
......
......@@ -203,7 +203,7 @@ TEST(Subgraph, generate_model_and_check_precision) {
valid_places,
input_tensor_shape,
input_tensor_type,
FLAGS_optimized_model_dir + "/ref_opt_model");
FLAGS_optimized_model_dir + "_ref_opt_model");
// Generate and run optimized model on NPU/XPU as the target predictor
#ifdef LITE_WITH_NPU
valid_places.push_back(lite_api::Place{TARGET(kNPU), PRECISION(kFloat)});
......@@ -217,7 +217,7 @@ TEST(Subgraph, generate_model_and_check_precision) {
valid_places,
input_tensor_shape,
input_tensor_type,
FLAGS_optimized_model_dir + "/tar_opt_model");
FLAGS_optimized_model_dir + "_tar_opt_model");
// Check the difference of the output tensors between reference predictor and
// target predictor
CheckOutputTensors(tar_predictor, ref_predictor, output_tensor_type);
......
......@@ -25,7 +25,7 @@ cd Paddle-Lite
--android_stl=c++_static \
--build_extra=ON \
--shutdown_log=OFF \
tiny_publish
full_publish
```
进入编译目录,下载模型和图片的压缩包,编译可执行文件。
......@@ -40,10 +40,11 @@ make
```
# 下载paddlehub以后,通过python执行以下代码
import paddlehub as hub
pyramidbox_lite_mask = hub.Module(name="pyramidbox_lite_mask")
pyramidbox_lite_mobile_mask = hub.Module(name="pyramidbox_lite_mobile_mask")
# 将模型保存在test_program文件夹之中
pyramidbox_lite_mask.processor.save_inference_model(dirname="test_program")
通过以上命令,可以获得人脸检测和口罩佩戴判断模型,分别存储在pyramidbox_lite和mask_detector之中。文件夹中的__model__是模型结构文件,__param__文件是权重文件。
pyramidbox_lite_mobile_mask.processor.save_inference_model(dirname="test_program")
# 通过以上命令,可以获得人脸检测和口罩佩戴判断模型,分别存储在pyramidbox_lite和mask_detector之中。文件夹中的__model__是模型结构文件,__param__文件是权重文件。
# 从PaddleHub下载的是预测模型,需要使用PaddleLite提供的model_optimize_tools对预测模型进行转换,请参考[模型转换文档](https://paddlepaddle.github.io/Paddle-Lite/v2.2.0/model_optimize_tool/)。
```
电脑连接安卓手机,将可执行文件、测试图片、模型文件、预测库push到安卓手机上。
......
......@@ -27,7 +27,10 @@ int64_t ShapeProduction(const shape_t& shape) {
void RunModel(std::string model_dir) {
// 1. Set MobileConfig
MobileConfig config;
config.set_model_dir(model_dir);
config.set_model_from_file(model_dir);
// NOTE: To load model transformed by model_optimize_tool before
// release/v2.3.0, plese use `set_model_dir` API as listed below.
// config.set_model_dir(model_dir);
// 2. Create PaddlePredictor by MobileConfig
std::shared_ptr<PaddlePredictor> predictor =
......
......@@ -46,6 +46,7 @@ class Tensor {
*/
class PaddlePredictor {
public:
void Init();
std::unique_ptr<Tensor> GetTensor(const std::string &id) const;
std::unique_ptr<Tensor> GetMutableTensor(const std::string &id);
......
......@@ -67,21 +67,21 @@ void LookupTableCompute::Run() {
REGISTER_LITE_KERNEL(lookup_table,
kARM,
kFloat,
kAny,
kNCHW,
paddle::lite::kernels::arm::LookupTableCompute,
def)
.BindInput("W", {LiteType::GetTensorTy(TARGET(kARM))})
.BindInput("Ids", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kInt64))})
.BindInput("Ids", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kAny))})
.BindOutput("Out", {LiteType::GetTensorTy(TARGET(kARM))})
.Finalize();
REGISTER_LITE_KERNEL(lookup_table_v2,
kARM,
kFloat,
kAny,
kNCHW,
paddle::lite::kernels::arm::LookupTableCompute,
def)
.BindInput("W", {LiteType::GetTensorTy(TARGET(kARM))})
.BindInput("Ids", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kInt64))})
.BindInput("Ids", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kAny))})
.BindOutput("Out", {LiteType::GetTensorTy(TARGET(kARM))})
.Finalize();
......@@ -21,7 +21,7 @@ namespace lite {
namespace kernels {
namespace arm {
class LookupTableCompute : public KernelLite<TARGET(kARM), PRECISION(kFloat)> {
class LookupTableCompute : public KernelLite<TARGET(kARM), PRECISION(kAny)> {
public:
using param_t = operators::LookupTableParam;
......
......@@ -53,7 +53,7 @@ void lookup_table_compute_ref(const operators::LookupTableParam &param) {
TEST(lookup_table_arm, retrieve_op) {
auto lookup_table =
KernelRegistry::Global().Create<TARGET(kARM), PRECISION(kFloat)>(
KernelRegistry::Global().Create<TARGET(kARM), PRECISION(kAny)>(
"lookup_table");
ASSERT_FALSE(lookup_table.empty());
ASSERT_TRUE(lookup_table.front());
......@@ -61,7 +61,7 @@ TEST(lookup_table_arm, retrieve_op) {
TEST(lookup_table_arm, init) {
LookupTableCompute lookup_table;
ASSERT_EQ(lookup_table.precision(), PRECISION(kFloat));
ASSERT_EQ(lookup_table.precision(), PRECISION(kAny));
ASSERT_EQ(lookup_table.target(), TARGET(kARM));
}
......@@ -112,4 +112,4 @@ TEST(lookup_table_arm, compute) {
} // namespace lite
} // namespace paddle
USE_LITE_KERNEL(lookup_table, kARM, kFloat, kNCHW, def);
USE_LITE_KERNEL(lookup_table, kARM, kAny, kNCHW, def);
......@@ -65,6 +65,6 @@ REGISTER_LITE_KERNEL(write_to_array,
paddle::lite::kernels::arm::WriteToArrayCompute,
def)
.BindInput("X", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kAny))})
.BindInput("I", {LiteType::GetTensorTy(TARGET(kARM))})
.BindInput("I", {LiteType::GetTensorTy(TARGET(kARM), PRECISION(kAny))})
.BindOutput("Out", {LiteType::GetTensorListTy(TARGET(kARM))})
.Finalize();
......@@ -73,8 +73,7 @@ int ReshapeConverter(void* ctx, OpLite* op, KernelBase* kernel) {
auto shape =
std::vector<int>(actual_shape_data,
actual_shape_data + actual_shape_dims.production());
auto out_dims = lite::operators::ValidateShape(shape, x_dims);
auto out_shape = out_dims.Vectorize();
auto out_shape = lite::operators::ValidateShape(shape, x_dims);
if (out_shape.size() > 4) {
LOG(WARNING) << "[NPU] HiAI DDK only supports less than 4 dimensions, "
"but Shape has "
......@@ -88,8 +87,7 @@ int ReshapeConverter(void* ctx, OpLite* op, KernelBase* kernel) {
reshape_op->set_input_w(*actual_shape_node->data());
} else {
auto shape = op_info->GetAttr<std::vector<int>>("shape");
auto out_dims = lite::operators::ValidateShape(shape, x_dims);
auto out_shape = out_dims.Vectorize();
auto out_shape = lite::operators::ValidateShape(shape, x_dims);
if (out_shape.size() > 4) {
LOG(WARNING) << "[NPU] HiAI DDK only supports less than 4 dimensions, "
"but shape has "
......
......@@ -16,7 +16,6 @@ add_kernel(io_copy_compute_opencl OPENCL basic SRCS io_copy_compute.cc DEPS ${te
add_kernel(relu_opencl OPENCL basic SRCS relu_compute.cc DEPS ${cl_kernel_deps})
add_kernel(sigmoid_opencl OPENCL basic SRCS sigmoid_compute.cc DEPS ${cl_kernel_deps})
add_kernel(depthwise_conv2d_opencl OPENCL basic SRCS depthwise_conv2d_compute.cc DEPS ${cl_kernel_deps})
#add_kernel(conv2d_1x1_opencl OPENCL basic SRCS conv2d_1x1_compute.cc DEPS ${cl_kernel_deps})
add_kernel(reshape_opencl OPENCL basic SRCS reshape_compute.cc DEPS ${cl_kernel_deps})
add_kernel(conv_opencl OPENCL basic SRCS conv_compute.cc DEPS ${cl_kernel_deps} cl_image_converter)
add_kernel(layout_opencl OPENCL basic SRCS layout_compute.cc DEPS ${cl_kernel_deps})
......@@ -62,14 +61,10 @@ lite_cc_test(test_depthwise_conv2d_opencl SRCS depthwise_conv2d_compute_test.cc
DEPS depthwise_conv2d_opencl op_registry program context
ARGS --cl_path=${CMAKE_SOURCE_DIR}/lite/backends/opencl)
lite_cc_test(test_depthwise_conv2d_basic_opencl SRCS depthwise_conv2d_basic_compute_test.cc
DEPS depthwise_conv2d_opencl op_registry program context
lite_cc_test(test_depthwise_conv2d_image2d_opencl SRCS depthwise_conv2d_image2d_compute_test.cc
DEPS conv_opencl op_registry program context
ARGS --cl_path=${CMAKE_SOURCE_DIR}/lite/backends/opencl)
#lite_cc_test(test_conv2d_1x1_opencl SRCS conv2d_1x1_compute_test.cc
# DEPS conv2d_1x1_opencl op_registry program context
# ARGS --cl_path=${CMAKE_SOURCE_DIR}/lite/backends/opencl)
lite_cc_test(test_reshape_opencl SRCS reshape_compute_test.cc
DEPS reshape_opencl op_registry program context
ARGS --cl_path=${CMAKE_SOURCE_DIR}/lite/backends/opencl)
......
......@@ -356,17 +356,17 @@ REGISTER_LITE_KERNEL(
DATALAYOUT(kImageDefault))})
.Finalize();
REGISTER_LITE_KERNEL(concat, kOpenCL, kFloat, kNCHW, Concat_buffer, def)
.BindInput("X",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kNCHW))})
.BindInput("AxisTensor",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kInt32),
DATALAYOUT(kNCHW))})
.BindOutput("Out",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kNCHW))})
.Finalize();
// REGISTER_LITE_KERNEL(concat, kOpenCL, kFloat, kNCHW, Concat_buffer, def)
// .BindInput("X",
// {LiteType::GetTensorTy(TARGET(kOpenCL),
// PRECISION(kFloat),
// DATALAYOUT(kNCHW))})
// .BindInput("AxisTensor",
// {LiteType::GetTensorTy(TARGET(kOpenCL),
// PRECISION(kInt32),
// DATALAYOUT(kNCHW))})
// .BindOutput("Out",
// {LiteType::GetTensorTy(TARGET(kOpenCL),
// PRECISION(kFloat),
// DATALAYOUT(kNCHW))})
// .Finalize();
......@@ -73,7 +73,7 @@ void concat_mul_compute_ref(std::vector<const dtype *> ins_data,
}
}
}
#if 1 // concat_buffer
#if 0 // concat_buffer
TEST(opencl_concat_buffer, compute) {
// prepare data
const DDim x0_dim = DDim(std::vector<DDim::value_type>{1, 2, 3, 4});
......@@ -382,7 +382,7 @@ TEST(concat_image2d_fp32, compute) {
} // namespace paddle
// concat buffer
USE_LITE_KERNEL(concat, kOpenCL, kFloat, kNCHW, def);
// USE_LITE_KERNEL(concat, kOpenCL, kFloat, kNCHW, def);
// concat image2d fp32
USE_LITE_KERNEL(layout, kOpenCL, kAny, kImageDefault, NCHW_to_ImageDefault);
......
// Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <vector>
#include "lite/backends/opencl/cl_include.h"
#include "lite/core/kernel.h"
#include "lite/core/op_registry.h"
#include "lite/kernels/opencl/image_helper.h"
#include "lite/operators/op_params.h"
#include "lite/utils/replace_stl/stream.h"
namespace paddle {
namespace lite {
namespace kernels {
namespace opencl {
#define USE_BUFFER_FOR_CONV1x1_BIAS
class Conv2d1x1Image2DCompute : public KernelLite<TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault)> {
public:
using param_t = operators::ConvParam;
void PrepareForRun() override {
const auto& param = *param_.get_mutable<param_t>();
if (param.fuse_relu) {
build_options_ += " -DRELU";
}
const bool has_bias = param.bias != nullptr;
const bool is_element_wise_bias =
has_bias && param.output->dims() == param.bias->dims();
if (has_bias) {
build_options_ += is_element_wise_bias ? " -DBIASE_ELE" : " -DBIASE_CH";
}
auto& context = ctx_->As<OpenCLContext>();
if (param.x->dims()[1] % 4 == 0) {
context.cl_context()->AddKernel(kernel_func_name_simple_,
"image/conv2d_1x1_kernel.cl",
build_options_);
} else {
context.cl_context()->AddKernel(
kernel_func_name_, "image/conv2d_1x1_kernel.cl", build_options_);
}
}
void Run() override {
const auto& param = *param_.get_mutable<param_t>();
auto input_dims = param.x->dims();
auto paddings = *param.paddings;
auto strides = param.strides;
auto* input_image = param.x->data<float, cl::Image2D>();
auto* filter_image = param.filter->data<float, cl::Image2D>();
auto filter_dims = param.filter->dims();
auto output_dims = param.output->dims();
int input_width = input_dims[3];
int input_height = input_dims[2];
int output_width = output_dims[3];
int output_height = output_dims[2];
auto out_image_shape = InitImageDimInfoWith(output_dims);
auto* out_image = param.output->mutable_data<float, cl::Image2D>(
out_image_shape["width"], out_image_shape["height"]);
const bool has_bias = param.bias != nullptr;
const bool is_element_wise_bias =
has_bias && param.output->dims() == param.bias->dims();
int offset = static_cast<int>(param.filter->dims()[2]) / 2 -
static_cast<int>(paddings[0]);
// calc input_c_block
auto input_image_shape = InitImageDimInfoWith(input_dims);
int input_c_block = input_image_shape["width"] / input_dims[3];
int input_c = input_dims[1];
auto dilations = *param.dilations;
const std::vector<size_t>& default_work_size =
DefaultWorkSize(output_dims,
DDim(std::vector<DDim::value_type>{
static_cast<int64_t>(out_image_shape["width"]),
static_cast<int64_t>(out_image_shape["height"])}));
int c_block = default_work_size[0];
int w = default_work_size[1];
int nh = default_work_size[2];
VLOG(4) << "============ conv2d_1x1 params ============";
VLOG(4) << "input_image_shape: " << input_image_shape["width"] << ","
<< input_image_shape["height"];
VLOG(4) << "input_c_block: " << input_c_block;
VLOG(4) << "input_c: " << input_c;
VLOG(4) << "input_image: " << input_image;
VLOG(4) << "filter_dims: " << filter_dims;
VLOG(4) << "filter_image: " << filter_image;
VLOG(4) << "output_dims: " << output_dims;
VLOG(4) << "out_image_shape: " << out_image_shape["width"] << ", "
<< out_image_shape["height"];
VLOG(4) << "paddings: " << paddings[0] << "," << paddings[1];
VLOG(4) << "has bias: " << has_bias;
VLOG(4) << "is_element_wise_bias : " << is_element_wise_bias;
VLOG(4) << "strides: " << strides[0] << "," << strides[1];
VLOG(4) << "offset: " << offset;
VLOG(4) << "dilations.size : " << dilations.size();
VLOG(4) << "dilations: " << dilations[0] << ", " << dilations[1];
VLOG(4) << "default work size{c_block, w, nh}: "
<< "{" << c_block << ", " << w << ", " << nh << ""
<< "}";
CHECK_GE(dilations.size(), 2);
CHECK(dilations[0] == dilations[1]);
CHECK_GE(input_dims.size(), 4);
CHECK_GE(paddings.size(), 2);
CHECK(paddings[0] == paddings[1]);
CHECK_GE(strides.size(), 2);
CHECK(strides[0] == strides[1]);
// handle bias use buffer for channel wise , use image for element wise
const cl::Buffer* bias_buf = nullptr;
const cl::Image2D* bias_image = nullptr;
if (has_bias) {
#ifndef USE_BUFFER_FOR_CONV1x1_BIAS
is_element_wise_bias
? (bias_image = param.bias->data<float, cl::Image2D>())
: (bias_buf = param.bias->data<float, cl::Buffer>());
#else
bias_image = param.bias->data<float, cl::Image2D>();
#endif
}
auto& context = ctx_->As<OpenCLContext>();
CHECK(context.cl_context() != nullptr);
STL::stringstream kernel_key;
if (input_dims[1] % 4 == 0) {
kernel_key << kernel_func_name_simple_ << build_options_;
} else {
kernel_key << kernel_func_name_ << build_options_;
}
auto kernel = context.cl_context()->GetKernel(kernel_key.str());
int maped_w = maptofactor(w, 4);
VLOG(4) << "kernel_key: " << kernel_key.str();
VLOG(4) << "kernel ready ... " << kernel_key.str();
VLOG(4) << "maped_w: " << maped_w;
cl_int status;
int arg_idx = 0;
status = kernel.setArg(arg_idx, c_block);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, maped_w);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, nh);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, *input_image);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, *filter_image);
CL_CHECK_FATAL(status);
if (has_bias) {
#ifndef USE_BUFFER_FOR_CONV1x1_BIAS
if (is_element_wise_bias != 0) {
VLOG(4) << "set bias_image: ";
status = kernel.setArg(++arg_idx, *bias_image);
} else {
VLOG(4) << "set bias_buf: ";
status = kernel.setArg(++arg_idx, *bias_buf);
}
#else
status = kernel.setArg(++arg_idx, *bias_image);
#endif
CL_CHECK_FATAL(status);
}
status = kernel.setArg(++arg_idx, *out_image);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, strides[0]);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, offset);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, input_c_block);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, input_c);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, dilations[0]);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, input_width);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, input_height);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, output_width);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, output_height);
CL_CHECK_FATAL(status);
status = kernel.setArg(++arg_idx, w);
CL_CHECK_FATAL(status);
auto global_work_size =
cl::NDRange{static_cast<size_t>(default_work_size.data()[0]),
static_cast<size_t>(maped_w),
static_cast<size_t>(default_work_size.data()[2])};
VLOG(4) << "out_image: " << out_image;
VLOG(4) << "global_work_size[3D]: {" << global_work_size[0] << ","
<< global_work_size[1] << "," << global_work_size[2] << "}";
status = context.cl_context()->GetCommandQueue().enqueueNDRangeKernel(
kernel,
cl::NullRange,
global_work_size,
cl::NullRange,
nullptr,
event_.get());
CL_CHECK_FATAL(status);
context.cl_wait_list()->emplace(out_image, event_);
}
private:
std::string kernel_func_name_{"conv2d_1x1"};
std::string kernel_func_name_simple_{"conv2d_1x1_simple"};
std::string build_options_{"-DCL_DTYPE_float"};
std::shared_ptr<cl::Event> event_{new cl::Event};
};
} // namespace opencl
} // namespace kernels
} // namespace lite
} // namespace paddle
REGISTER_LITE_KERNEL(conv2d_1x1,
kOpenCL,
kFloat,
kImageDefault,
paddle::lite::kernels::opencl::Conv2d1x1Image2DCompute,
image2d)
.BindInput("Input",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault))})
.BindInput("Bias",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault))})
.BindInput("Filter",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageNW))})
.BindOutput("Output",
{LiteType::GetTensorTy(TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault))})
.Finalize();
// Copyright (c) 2019 PaddlePaddle Authors. All Rights Reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <gtest/gtest.h>
#include <random>
#include "lite/backends/opencl/cl_image_converter.h"
#include "lite/backends/opencl/target_wrapper.h"
#include "lite/core/op_registry.h"
#include "lite/core/tensor.h"
#include "lite/utils/logging.h"
namespace paddle {
namespace lite {
template <typename Dtype1, typename Dtype2>
static void conv_basic(const Dtype1* din,
Dtype2* dout,
int num,
int chout,
int hout,
int wout,
int chin,
int hin,
int win,
const Dtype1* weights,
const Dtype2* bias,
int group,
int kernel_w,
int kernel_h,
int stride_w,
int stride_h,
int dila_w,
int dila_h,
int pad_w,
int pad_h,
bool flag_bias,
bool flag_relu) {
Dtype2 beta = 0;
auto src_data = din;
auto dst_data_ref = dout;
auto weights_data = weights;
auto with_bias = flag_bias;
auto bias_data = bias;
int in_num = num;
int out_channels = chout;
int out_h = hout;
int out_w = wout;
int in_channel = chin;
int in_h = hin;
int in_w = win;
int out_c_group = out_channels / group;
int in_c_group = in_channel / group;
for (int n = 0; n < in_num; ++n) {
for (int g = 0; g < group; ++g) {
for (int oc = 0; oc < out_c_group; ++oc) {
for (int oh = 0; oh < out_h; ++oh) {
for (int ow = 0; ow < out_w; ++ow) {
int out_idx = n * group * out_c_group * out_h * out_w +
g * out_c_group * out_h * out_w + oc * out_h * out_w +
oh * out_w + ow;
Dtype2 bias_d =
with_bias ? (bias_data[g * out_c_group + oc]) : (Dtype2)0;
dst_data_ref[out_idx] = bias_d; // + dst_data_ref[out_idx] * beta;
for (int ic = 0; ic < in_c_group; ++ic) {
for (int kh = 0; kh < kernel_h; ++kh) {
for (int kw = 0; kw < kernel_w; ++kw) {
int iw = ow * stride_w - pad_w + kw * (dila_w);
int ih = oh * stride_h - pad_h + kh * (dila_h);
if (iw < 0 || iw >= in_w) continue;
if (ih < 0 || ih >= in_h) continue;
int iidx = n * in_channel * in_h * in_w +
g * in_c_group * in_h * in_w + ic * in_h * in_w +
ih * in_w + iw;
int widx =
g * out_c_group * in_c_group * kernel_h * kernel_w +
oc * in_c_group * kernel_h * kernel_w +
ic * kernel_h * kernel_w + kh * kernel_w + kw;
dst_data_ref[out_idx] += src_data[iidx] * weights_data[widx];
}
}
}
if (flag_relu) {
dst_data_ref[out_idx] = dst_data_ref[out_idx] > (Dtype2)0
? dst_data_ref[out_idx]
: (Dtype2)0;
}
}
}
}
}
}
}
TEST(conv2d_1x1, compute) {
// conv infos
const int ksize = 1;
const int stride = 1;
const int pad = 0;
const int group = 1;
const int dilation = 0;
// int loop_cnt = 0;
#ifdef LOOP_TEST
for (int batch_size = 1; batch_size < 4; ++batch_size) {
for (int oc = 4; oc < 10; oc += 1) { // oc
for (int ih = 4; ih < 9; ih += 1) { // ih
/*int iw = ih;*/ for (int iw = 4; iw < 10; iw += 1) { // iw
for (int ic = 4; ic < 10; ic += 1) { // ic
for (bool bias_flag : {true, false}) {
for (bool relu_flag : {true, false}) {
#else
const int batch_size = 1;
const int oc = 4;
const int ih = 8;
const int iw = 8;
const int ic = 4;
const bool bias_flag = false;
const bool relu_flag = false;
#endif
const int oh = ih;
const int ow = iw;
VLOG(4) << "to get kernel ...";
auto kernels =
KernelRegistry::Global().Create("conv2d_1x1",
TARGET(kOpenCL),
PRECISION(kFloat),
DATALAYOUT(kImageDefault));
ASSERT_FALSE(kernels.empty());
auto kernel = std::move(kernels.front());
VLOG(4) << "created conv2d_1x1 kernel";
VLOG(4) << "prepare kernel ------";
lite::Tensor input, filter, bias, output;
operators::ConvParam param;
param.x = &input;
param.filter = &filter;
param.output = &output;
if (bias_flag) {
param.bias = &bias;
}
param.fuse_relu = relu_flag;
std::vector<int> paddings = {pad, pad, pad, pad};
std::vector<int> dilations = {dilation, dilation};
param.paddings = std::make_shared<std::vector<int>>(paddings);
param.dilations = std::make_shared<std::vector<int>>(dilations);
param.strides = std::vector<int>{stride, stride};
std::unique_ptr<KernelContext> context(new KernelContext);
context->As<OpenCLContext>().InitOnce();
std::unique_ptr<KernelContext> conv_1x1_context(
new KernelContext);
context->As<OpenCLContext>().CopySharedTo(
&(conv_1x1_context->As<OpenCLContext>()));
kernel->SetContext(std::move(conv_1x1_context));
const DDim& input_dim =
lite::DDim{std::vector<int64_t>({batch_size, ic, ih, iw})};
const DDim& filter_dim =
lite::DDim{std::vector<int64_t>({oc, ic, ksize, ksize})};
const DDim& out_dim =
lite::DDim{std::vector<int64_t>({batch_size, oc, ih, iw})};
// element wise bias
const DDim& bias_dim = lite::DDim{std::vector<int64_t>({oc})};
param.x->Resize(input_dim);
param.filter->Resize(filter_dim);
param.output->Resize(out_dim);
if (bias_flag) {
param.bias->Resize(bias_dim);
}
kernel->SetParam(param);
size_t input_image_width = iw * ((ic + 3) / 4);
size_t input_image_height = ih * batch_size;
size_t out_image_width = ow * ((oc + 3) / 4);
size_t out_image_height = oh * batch_size;
size_t bias_image_width = ow * ((oc + 3) / 4);
size_t bias_image_height = oh * batch_size;
size_t filter_image_width = ksize * ((oc + 3) / 4);
size_t filter_image_height = ic * ksize;
const size_t cl_image2d_row_pitch{0};
const size_t cl_image2d_slice_pitch{0};
std::default_random_engine engine;
std::uniform_real_distribution<float> gen(-5, 5);
std::vector<float> input_v(batch_size * ic * ih * iw);
std::vector<float> filter_v(oc * ic * ksize * ksize);
std::vector<float> output_v(batch_size * oc * ih * iw);
std::vector<float> bias_v(oc);
VLOG(4) << "gen input and filter ...";
for (auto& i : input_v) {
i = gen(engine);
}
for (auto& f : filter_v) {
f = gen(engine);
}
VLOG(4) << "after gen input and filter ...";
VLOG(4) << "input_v.size(): " << input_v.size();
VLOG(4) << "filter_v.size(): " << filter_v.size();
VLOG(4) << "output_v.size(): " << output_v.size();
VLOG(4) << "bias_v.size(): " << bias_v.size();
VLOG(4) << "input_dim.production(): " << input_dim.production();
VLOG(4) << "filter_dim.production(): "
<< filter_dim.production();
VLOG(4) << "out_dim.production(): " << out_dim.production();
VLOG(4) << "bias_dim.production(): " << bias_dim.production();
VLOG(4) << "4 * input_image_height * input_image_width: "
<< 4 * input_image_height * input_image_width;
VLOG(4) << "4 * filter_image_width * filter_image_height: "
<< 4 * filter_image_width * filter_image_height;
CHECK(input_dim.production() == input_v.size());
CHECK_LE(input_dim.production(),
4 * input_image_height * input_image_width);
CHECK(filter_dim.production() == filter_v.size());
CHECK_LE(filter_dim.production(),
4 * filter_image_width * filter_image_height);
paddle::lite::CLImageConverterDefault default_convertor;
VLOG(4) << "set mapped input ...";
std::vector<float> x_image_v(
input_image_width * input_image_height * 4); // 4 : RGBA
std::vector<float> filter_image_v(
filter_image_width * filter_image_height * 4); // 4 : RGBA
std::vector<float> bias_image_v(
bias_image_width * bias_image_height * 4); // 4 : RGBA
std::vector<float> out_image_v(
out_image_width * out_image_height * 4); // 4 : RGBA
default_convertor.NCHWToImage(
input_v.data(), x_image_v.data(), input_dim);
/* for (int j = 0; j < input_v.size(); j += 1) {
// VLOG(4) << "input_v
input[" << j << "]:
// " << input_v.data()[j];
std::cout << j << " " << input_v.data()[j] <<
std::endl;
}
std::cout << std::endl;
for (int j = 0; j < x_image_v.size(); j += 1) {
// VLOG(4) << "x_image_v
input[" << j <<
// "]: " <<
x_image_v.data()[j];
std::cout << j << " " << x_image_v.data()[j]
<< std::endl;
}*/
VLOG(4) << "set mapped filter ...";
paddle::lite::CLImageConverterNWBlock nw_convertor;
nw_convertor.NCHWToImage(
filter_v.data(), filter_image_v.data(), filter_dim);
auto* input_image2d = input.mutable_data<float, cl::Image2D>(
input_image_width, input_image_height, x_image_v.data());
auto* filter_image2d = filter.mutable_data<float, cl::Image2D>(
filter_image_width,
filter_image_height,
filter_image_v.data());
if (bias_flag) {
nw_convertor.NCHWToImage(
filter_v.data(), filter_image_v.data(), filter_dim);
for (int i = 0; i < bias_dim.production(); ++i) {
bias_v[i] = static_cast<int>(gen(engine));
}
CLImageConverterFolder folder_convertor;
folder_convertor.NCHWToImage(
bias_v.data(), bias_image_v.data(), bias_dim);
auto* bias_data = bias.mutable_data<float, cl::Image2D>(
bias_image_width, bias_image_height, bias_image_v.data());
}
VLOG(4) << "resize output ...";
output.Resize(out_dim);
// cpu conv basic calc
lite::Tensor out_ref;
out_ref.Resize(out_dim);
VLOG(4) << "prepare kernel ready";
VLOG(4) << "kernel launch ...";
kernel->Launch();
VLOG(4) << "mutable output ...";
auto* output_image2d = output.mutable_data<float, cl::Image2D>(
out_image_width, out_image_height);
auto* wait_list = context->As<OpenCLContext>().cl_wait_list();
auto* out_ptr = param.output->data<float, cl::Image2D>();
auto it = wait_list->find(out_ptr);
if (it != wait_list->end()) {
VLOG(4) << "--- Find the sync event for the target cl "
"tensor. ---";
auto& event = *(it->second);
event.wait();
} else {
LOG(FATAL) << "Could not find the sync event for the target "
"cl tensor.";
}
TargetWrapperCL::ImgcpySync(out_image_v.data(),
output.data<float, cl::Image2D>(),
out_image_width,
out_image_height,
cl_image2d_row_pitch,
cl_image2d_slice_pitch,
IoDirection::DtoH);
DDim out_image_shape =
default_convertor.InitImageDimInfoWith(output.dims());
default_convertor.ImageToNCHW(out_image_v.data(),
output_v.data(),
out_image_shape,
output.dims());
VLOG(4) << "mutable_data out_ref_data: ";
// run cpu ref
auto* out_ref_data = out_ref.mutable_data<float>(TARGET(kARM));
VLOG(4) << " conv_basic beigin ..... ";
conv_basic<float, float>(input_v.data(),
out_ref_data,
batch_size,
oc,
oh,
ow,
ic,
ih,
iw,
filter_v.data(),
bias_v.data(), // mapped_bias,
group,
ksize,
ksize,
stride,
stride,
dilation,
dilation,
pad,
pad,
bias_flag,
relu_flag);
VLOG(4) << " conv_basic end ..... ";
VLOG(4) << " out_dim: " << out_dim;
const DDim& out_image_dims = lite::DDim{std::vector<int64_t>(
{static_cast<int64_t>(out_image_width),
static_cast<int64_t>(out_image_height)})};
for (int i = 0; i < out_dim.production(); i++) {
EXPECT_NEAR(output_v[i], out_ref_data[i], 1e-2);
if (abs(output_v[i] - out_ref_data[i]) > 1e-2) {
LOG(FATAL) << "error idx:" << i;
}
}
#ifdef LOOP_TEST
}
}
}
}
}
}
}
#else
// nothing to do.
#endif
}
} // namespace lite
} // namespace paddle
USE_LITE_KERNEL(conv2d_1x1, kOpenCL, kFloat, kImageDefault, image2d);
此差异已折叠。
......@@ -74,6 +74,9 @@ class ConvImageCompute : public KernelLite<TARGET(kOpenCL),
void Conv2d3x3();
void Conv2d5x5();
void Conv2d7x7();
void DepthwiseConv2d3x3s1();
void DepthwiseConv2d3x3();
void DepthwiseConv2d();
kernel_t impl_;
std::vector<std::string> kernel_func_names_{};
......
......@@ -166,6 +166,8 @@ void PrintData(std::string name,
}
}
// buffer
#if 0
// #define PRINT_RESULT
#define LOOP_TEST
TEST(conv2d, compute_conv2d_1x1) {
......@@ -623,8 +625,9 @@ TEST(conv2d, compute_conv2d_gemm) {
} // batch_size
#endif
}
#endif
} // namespace lite
} // namespace paddle
USE_LITE_KERNEL(conv2d, kOpenCL, kFloat, kNCHW, def);
// USE_LITE_KERNEL(conv2d, kOpenCL, kFloat, kNCHW, def);
......@@ -559,9 +559,11 @@ TEST(conv2d, compute_image2d_3x3) {
// element wise bias
const DDim& bias_dim = lite::DDim{std::vector<int64_t>({oc})};
LOG(INFO) << "input_dim:" << input_dim
<< " filter_dim:" << filter_dim
<< " out_dim:" << out_dim;
VLOG(2) << "input_dim:" << input_dim
<< " filter_dim:" << filter_dim << " out_dim:" << out_dim
<< " bias_flag:" << bias_flag << " bias_dim:" << bias_dim
<< " group:" << group << " stride:" << stride
<< " pad:" << pad << " dilation:" << dilation;
param.x->Resize(input_dim);
param.filter->Resize(filter_dim);
......@@ -902,6 +904,12 @@ TEST(conv2d, compute_image2d_5x5) {
// element wise bias
const DDim& bias_dim = lite::DDim{std::vector<int64_t>({oc})};
VLOG(2) << "input_dim:" << input_dim
<< " filter_dim:" << filter_dim << " out_dim:" << out_dim
<< " bias_flag:" << bias_flag << " bias_dim:" << bias_dim
<< " group:" << group << " stride:" << stride
<< " pad:" << pad << " dilation:" << dilation;
param.x->Resize(input_dim);
param.filter->Resize(filter_dim);
param.output->Resize(out_dim);
......
......@@ -66,6 +66,8 @@ void PrintData(std::string name, float* a, const int rows, const int cols) {
}
}
// buffer
#if 0 // fc_buffer
// #define PRINT_RESULT
#define LOOP_TEST
TEST(fc, compute) {
......@@ -193,8 +195,9 @@ TEST(fc, compute) {
} // m
#endif
}
#endif // fc_buffer
} // namespace lite
} // namespace paddle
USE_LITE_KERNEL(fc, kOpenCL, kFloat, kNCHW, def);
// USE_LITE_KERNEL(fc, kOpenCL, kFloat, kNCHW, def);
......@@ -229,15 +229,15 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
} // namespace lite
} // namespace paddle
REGISTER_LITE_KERNEL(pool2d,
kOpenCL,
kFloat,
kNCHW,
paddle::lite::kernels::opencl::PoolCompute,
def)
.BindInput("X", {LiteType::GetTensorTy(TARGET(kOpenCL))})
.BindOutput("Out", {LiteType::GetTensorTy(TARGET(kOpenCL))})
.Finalize();
// REGISTER_LITE_KERNEL(pool2d,
// kOpenCL,
// kFloat,
// kNCHW,
// paddle::lite::kernels::opencl::PoolCompute,
// def)
// .BindInput("X", {LiteType::GetTensorTy(TARGET(kOpenCL))})
// .BindOutput("Out", {LiteType::GetTensorTy(TARGET(kOpenCL))})
// .Finalize();
REGISTER_LITE_KERNEL(pool2d,
kOpenCL,
......
......@@ -73,6 +73,8 @@ void pool_avg(const int padding_height,
}
}
// buffer
#if 0 // pool_buffer
TEST(pool2d_buffer_fp32, compute) {
LOG(INFO) << "to get kernel ...";
auto kernels = KernelRegistry::Global().Create(
......@@ -141,6 +143,7 @@ TEST(pool2d_buffer_fp32, compute) {
}
TargetWrapperCL::Unmap(out_data, mapped_out);
}
#endif // pool_buffer
TEST(pool2d_image2d_fp32, compute) {
LOG(INFO) << "to get kernel ...";
......@@ -239,5 +242,5 @@ TEST(pool2d_image2d_fp32, compute) {
} // namespace lite
} // namespace paddle
USE_LITE_KERNEL(pool2d, kOpenCL, kFloat, kNCHW, def);
// USE_LITE_KERNEL(pool2d, kOpenCL, kFloat, kNCHW, def);
USE_LITE_KERNEL(pool2d, kOpenCL, kFloat, kImageDefault, image2d);
此差异已折叠。
此差异已折叠。
此差异已折叠。
此差异已折叠。
......@@ -44,6 +44,7 @@ if(LITE_BUILD_EXTRA)
lite_cc_test(test_kernel_assign_value_compute SRCS assign_value_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${bm_kernels} ${cuda_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_box_clip_compute SRCS box_clip_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_reduce_mean_compute SRCS reduce_mean_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_reduce_sum_compute SRCS reduce_sum_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_reduce_prod_compute SRCS reduce_prod_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_stack_compute SRCS stack_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
lite_cc_test(test_kernel_range_compute SRCS range_compute_test.cc DEPS arena_framework ${xpu_kernels} ${npu_kernels} ${x86_kernels} ${cuda_kernels} ${bm_kernels} ${arm_kernels} ${lite_ops} ${host_kernels})
......
此差异已折叠。
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册