如何在框架外部自定义C++ OP

通常,如果PaddlePaddle的Operator(OP)库中没有您所需要的操作,建议先尝试使用已有的OP组合,如果无法组合出您需要的操作,可以尝试使用fluid.layers.py_func,也可以按照这篇教程自定义C++ OP。当然,如果用若干OP组合出来的OP性能无法满足您的要求,也可以自定义C++ OP。

自定义OP需要以下几个步骤:

  • 实现OP和注册OP,和在框架内部写OP完全相同,遵守"如何写新的C++ OP"的规范和步骤。当然,实现Gradient OP是可选的。
  • 编译出动态库。
  • 封装该OP的Python接口。
  • 写OP的单测。

下面通过一个具体的例子来详细的介绍,一步一步教会您如何实现。下面通过实现relu op来介绍。

自定义OP的实现

OP的实现与"如何写新的C++ OP"的教程相同,简答的说需要: 1). 定义OP的ProtoMaker,即描述OP的输入、输出、属性信息;2). 实现OP的定义和InferShape,以及OP的kernel函数,反向OP类似。3). 注册OP,以及OP的计算函数。

ReLU OP的CPU实现, relu_op.cc 文件:

  1. // relu_op.cc
  2. #include "paddle/fluid/framework/op_registry.h"
  3.  
  4. namespace paddle {
  5. namespace operators {
  6.  
  7. // 前向OP的输入X、输出Y、属性
  8. class Relu2OpMaker : public framework::OpProtoAndCheckerMaker {
  9. public:
  10. void Make() override {
  11. AddInput("X", "The input tensor.");
  12. AddOutput("Y", "Output of relu_op");
  13. AddComment(R"DOC(
  14. Relu Operator.
  15. Y = max(X, 0)
  16. )DOC");
  17. }
  18. };
  19.  
  20. // 前向OP的定义和InferShape实现,设置输出Y的shape
  21. class Relu2Op : public framework::OperatorWithKernel {
  22. public:
  23. using framework::OperatorWithKernel::OperatorWithKernel;
  24.  
  25. void InferShape(framework::InferShapeContext* ctx) const override {
  26. auto in_dims = ctx->GetInputDim("X");
  27. ctx->SetOutputDim("Y", in_dims);
  28. }
  29. };
  30.  
  31. // 实现前向OP的Kernel计算函数: Y = max(0, X)
  32. using Tensor = framework::Tensor;
  33. template <typename DeviceContext, typename T>
  34. class Relu2Kernel : public framework::OpKernel<T> {
  35. public:
  36. void Compute(const framework::ExecutionContext& ctx) const override {
  37. auto* in_t = ctx.Input<Tensor>("X");
  38. auto* out_t = ctx.Output<Tensor>("Y");
  39. auto x = in_t->data<T>();
  40. // mutable_data分配内存、获取指针
  41. auto y = out_t->mutable_data<T>(ctx.GetPlace());
  42. for (int i = 0; i < in_t->numel(); ++i) {
  43. y[i] = std::max(static_cast<T>(0.), x[i]);
  44. }
  45. }
  46. };
  47.  
  48. // 定义反向OP的输入Y和dY、输出dX、属性:
  49. class Relu2GradMaker : public framework::SingleGradOpDescMaker {
  50. public:
  51. using framework::SingleGradOpDescMaker::SingleGradOpDescMaker;
  52.  
  53. std::unique_ptr<framework::OpDesc> Apply() const override {
  54. auto* op = new framework::OpDesc();
  55. op->SetType("relu2_grad");
  56. op->SetInput("Y", Output("Y"));
  57. op->SetInput(framework::GradVarName("Y"), OutputGrad("Y"));
  58. op->SetAttrMap(Attrs());
  59. op->SetOutput(framework::GradVarName("X"), InputGrad("X"));
  60. return std::unique_ptr<framework::OpDesc>(op);
  61. }
  62. };
  63.  
  64. // 定义反向OP和InferShape实现,设置dX的shape
  65. class Relu2GradOp : public framework::OperatorWithKernel {
  66. public:
  67. using framework::OperatorWithKernel::OperatorWithKernel;
  68.  
  69. void InferShape(framework::InferShapeContext* ctx) const override {
  70. auto in_dims = ctx->GetInputDim(framework::GradVarName("Y"));
  71. ctx->SetOutputDim(framework::GradVarName("X"), in_dims);
  72. }
  73. };
  74.  
  75. // 实现反向OP的kernel函数 dx = dy * ( y > 0. ? 1. : 0)
  76. template <typename DeviceContext, typename T>
  77. class Relu2GradKernel : public framework::OpKernel<T> {
  78. public:
  79. void Compute(const framework::ExecutionContext& ctx) const override {
  80. auto* dy_t = ctx.Input<Tensor>(framework::GradVarName("Y"));
  81. auto* y_t = ctx.Input<Tensor>("Y");
  82. auto* dx_t = ctx.Output<Tensor>(framework::GradVarName("X"));
  83.  
  84. auto dy = dy_t->data<T>();
  85. auto y = y_t->data<T>();
  86. auto dx = dx_t->mutable_data<T>(ctx.GetPlace());
  87.  
  88. for (int i = 0; i < y_t->numel(); ++i) {
  89. dx[i] = dy[i] * (y[i] > static_cast<T>(0) ? 1. : 0.);
  90. }
  91. }
  92. };
  93.  
  94. } // namespace operators
  95. } // namespace paddle
  96.  
  97. namespace ops = paddle::operators;
  98. using CPU = paddle::platform::CPUDeviceContext;
  99. // 注册前向和反向op
  100. // 为了和框架内部的relu区分,这里注册的OP type为relu2
  101. REGISTER_OPERATOR(relu2, ops::Relu2Op, ops::Relu2OpMaker, ops::Relu2GradMaker);
  102. REGISTER_OPERATOR(relu2_grad, ops::Relu2GradOp);
  103. // 注册CPU的Kernel
  104. REGISTER_OP_CPU_KERNEL(relu2,
  105. ops::Relu2Kernel<CPU, float>,
  106. ops::Relu2Kernel<CPU, double>);
  107. REGISTER_OP_CPU_KERNEL(relu2_grad,
  108. ops::Relu2GradKernel<CPU, float>,
  109. ops::Relu2GradKernel<CPU, double>);

ReLU OP的GPU实现, relu_op.cc 文件:

  1. // relu_op.cu
  2. #include "paddle/fluid/framework/op_registry.h"
  3.  
  4. namespace paddle {
  5. namespace operators {
  6.  
  7. using Tensor = framework::Tensor;
  8.  
  9. template <typename T>
  10. __global__ void KeRelu2(const T* x, const int num, T* y) {
  11. int gid = blockIdx.x * blockDim.x + threadIdx.x;
  12. for (int i = gid; i < num; i += blockDim.x * gridDim.x) {
  13. y[i] = max(x[i], static_cast<T>(0.));
  14. }
  15. }
  16.  
  17. // 前向OP的kernel的GPU实现
  18. template <typename DeviceContext, typename T>
  19. class Relu2CUDAKernel : public framework::OpKernel<T> {
  20. public:
  21. void Compute(const framework::ExecutionContext& ctx) const override {
  22. auto* in_t = ctx.Input<Tensor>("X");
  23. auto* out_t = ctx.Output<Tensor>("Y");
  24. auto x = in_t->data<T>();
  25. auto y = out_t->mutable_data<T>(ctx.GetPlace());
  26.  
  27. auto& dev_ctx = ctx.template device_context<DeviceContext>();
  28.  
  29. int num = in_t->numel();
  30. int block = 512;
  31. int grid = (num + block - 1) / block;
  32. KeRelu2<T><<<grid, block, 0, dev_ctx.stream()>>>(x, num, y);
  33. }
  34. };
  35.  
  36. template <typename T>
  37. __global__ void KeRelu2Grad(const T* y, const T* dy, const int num, T* dx) {
  38. int gid = blockIdx.x * blockDim.x + threadIdx.x;
  39. for (int i = gid; i < num; i += blockDim.x * gridDim.x) {
  40. dx[i] = dy[i] * (y[i] > 0 ? 1. : 0.);
  41. }
  42. }
  43.  
  44. // 反向OP的kernel的GPU实现
  45. template <typename DeviceContext, typename T>
  46. class Relu2GradCUDAKernel : public framework::OpKernel<T> {
  47. public:
  48. void Compute(const framework::ExecutionContext& ctx) const override {
  49. auto* dy_t = ctx.Input<Tensor>(framework::GradVarName("Y"));
  50. auto* y_t = ctx.Input<Tensor>("Y");
  51. auto* dx_t = ctx.Output<Tensor>(framework::GradVarName("X"));
  52.  
  53. auto dy = dy_t->data<T>();
  54. auto y = y_t->data<T>();
  55. auto dx = dx_t->mutable_data<T>(ctx.GetPlace());
  56.  
  57. auto& dev_ctx = ctx.template device_context<DeviceContext>();
  58.  
  59. int num = dy_t->numel();
  60. int block = 512;
  61. int grid = (num + block - 1) / block;
  62. KeRelu2Grad<T><<<grid, block, 0, dev_ctx.stream()>>>(y, dy, num, dx);
  63. }
  64. };
  65.  
  66. } // namespace operators
  67. } // namespace paddle
  68.  
  69. using CUDA = paddle::platform::CUDADeviceContext;
  70. // 注册前向的GPU Kernel
  71. REGISTER_OP_CUDA_KERNEL(relu2,
  72. paddle::operators::Relu2CUDAKernel<CUDA, float>,
  73. paddle::operators::Relu2CUDAKernel<CUDA, double>);
  74. // 注册反向的GPU Kernel
  75. REGISTER_OP_CUDA_KERNEL(relu2_grad,
  76. paddle::operators::Relu2GradCUDAKernel<CUDA, float>,
  77. paddle::operators::Relu2GradCUDAKernel<CUDA, double>);

注意点:

  • OP的type不能和PaddlePaddle已有的OP type相同,否则在Python中使用时会报错。

自定义OP的编译

需要将实现的C++、CUDA代码编译成动态库,下面通过g++/nvcc编译,当然您也可以写Makefile或者CMake。

编译需要include PaddlePaddle的相关头文件,如上面代码 paddle/fluid/framework/op_registry.h ,需要链接PaddlePaddle的lib库。 可通过下面命令获取到:

  1. # python
  2. >>> import paddle
  3. >>> print(paddle.sysconfig.get_include())
  4. /paddle/pyenv/local/lib/python2.7/site-packages/paddle/include
  5. >>> print(paddle.sysconfig.get_lib())
  6. /paddle/pyenv/local/lib/python2.7/site-packages/paddle/libs

下面命令可编译出动态库:

  1. include_dir=$( python -c 'import paddle; print(paddle.sysconfig.get_include())' )
  2. lib_dir=$( python -c 'import paddle; print(paddle.sysconfig.get_lib())' )
  3.  
  4. echo $include_dir
  5. echo $lib_dir
  6.  
  7. # PaddlePaddel >=1.6.1, 仅需要include ${include_dir} 和 ${include_dir}/third_party
  8. nvcc relu_op.cu -c -o relu_op.cu.o -ccbin cc -DPADDLE_WITH_CUDA -DEIGEN_USE_GPU -DPADDLE_USE_DSO -Xcompiler -fPIC -std=c++11 -Xcompiler -fPIC -w --expt-relaxed-constexpr -O3 -DNVCC \
  9. -I ${include_dir} \
  10. -I ${include_dir}/third_party \
  11.  
  12. g++ relu_op.cc relu_op.cu.o -o relu2_op.so -shared -fPIC -std=c++11 -O3 \
  13. -I ${include_dir} \
  14. -I ${include_dir}/third_party \
  15. -L /usr/local/cuda/lib64 \
  16. -L ${lib_dir} -lpaddle_framework -lcudart
  17.  
  18. # PaddlePaddel 1.6.0, 需要include的third_party如下:
  19. # -I ${include_dir}/third_party/install/protobuf/include \
  20. # -I ${include_dir}/third_party/install/glog/include \
  21. # -I ${include_dir}/third_party/install/gflags/include \
  22. # -I ${include_dir}/third_party/install/xxhash/include \
  23. # -I ${include_dir}/third_party/boost \
  24. # -I ${include_dir}/third_party/eigen3 \
  25. # -I ${include_dir}/third_party/dlpack/include \

注意点:

  • NVCC编译GPU OP的cu文件时,需要加 -DPADDLE_WITH_CUDA -DEIGEN_USE_GPU -DPADDLE_USE_DSO
  • 可多个OP编译到同一个动态库中。

封装Python Layer接口

需要使用 fluid.load_op_library 接口调用加载动态库,使得PaddlePaddle的主进程中可以使用用户自定义的OP。

  1. # custom_op.py
  2. import paddle.fluid as fluid
  3. # 调用load_op_library加载动态库
  4. fluid.load_op_library('relu2_op.so')
  5.  
  6. from paddle.fluid.layer_helper import LayerHelper
  7.  
  8. def relu2(x, name=None):
  9. # relu2的type和在OP中定义的type相同
  10. helper = LayerHelper("relu2", **locals())
  11. # 创建输出Variable
  12. out = helper.create_variable_for_type_inference(dtype=x.dtype)
  13. helper.append_op(type="relu2", inputs={"X": x}, outputs={"Y": out})
  14. return out

注意点:

  • 一个动态库只需使用fluid.load_op_librarypaddle.fluid import之后加载一次即可。
  • Python接口的封装和PaddlePaddle框架内部的封装相同,更多的示例也可以阅读源码中 python/paddle/fluid/layers/nn.py的代码示例。

单测测试

可以写个简单的Python程序测试计算的正确性:

  1. import numpy as np
  2. import paddle.fluid as fluid
  3. from custom_op import relu2
  4.  
  5. data = fluid.layers.data(name='data', shape=[32], dtype='float32')
  6. relu = relu2(data)
  7. use_gpu = True # or False
  8. place = fluid.CUDAPlace(0) if use_gpu else fluid.CPUPlace()
  9. exe = fluid.Executor(place)
  10.  
  11. x = np.random.uniform(-1, 1, [4, 32]).astype('float32')
  12. out, = exe.run(feed={'data': x}, fetch_list=[relu])
  13. np.allclose(out, np.maximum(x,0.))

接下来可以在模型中使用您自定义的OP了!

如何在C++预测库中使用

暂时不支持在C++预测库中使用,后续会补充在C++预测库中的使用示例。

FAQ

  • Q:如果出现类似错误: cannot open shared object file: No such file or directory.

A: 需要设置动态库的路径到环境变量LD_LIBRARY_PATH中。