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

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

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

    下面通过一个具体的例子来详细的介绍，一步一步教会您如何实现。下面通过实现relu 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.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.     KeRelu2<T><<<grid, block, 0, dev_ctx.stream()>>>(x, num, y);
    32.   }
    33. };
    34.  
    35. template <typename T>
    36. __global__ void KeRelu2Grad(const T* y, const T* dy, const int num, T* dx) {
    37.   int gid = blockIdx.x * blockDim.x + threadIdx.x;
    38.   for (int i = gid; i < num; i += blockDim.x * gridDim.x) {
    39.     dx[i] = dy[i] * (y[i] > 0 ? 1. : 0.);
    40.   }
    41. }
    42.  
    43. // 反向OP的kernel的GPU实现
    44. template <typename DeviceContext, typename T>
    45. class Relu2GradCUDAKernel : public framework::OpKernel<T> {
    46.  public:
    47.   void Compute(const framework::ExecutionContext& ctx) const override {
    48.     auto* dy_t = ctx.Input<Tensor>(framework::GradVarName("Y"));
    49.     auto* y_t = ctx.Input<Tensor>("Y");
    50.     auto* dx_t = ctx.Output<Tensor>(framework::GradVarName("X"));
    51.  
    52.     auto dy = dy_t->data<T>();
    53.     auto y = y_t->data<T>();
    54.     auto dx = dx_t->mutable_data<T>(ctx.GetPlace());
    55.  
    56.     auto& dev_ctx = ctx.template device_context<DeviceContext>();
    57.  
    58.     int num = dy_t->numel();
    59.     int block = 512;
    60.     int grid = (num + block - 1) / block;
    61.     KeRelu2Grad<T><<<grid, block, 0, dev_ctx.stream()>>>(y, dy, num, dx);
    62.   }
    63. };
    64.  
    65. }  // namespace operators
    66. }  // namespace paddle
    67.  
    68. using CUDA = paddle::platform::CUDADeviceContext;
    69. // 注册前向的GPU Kernel
    70. REGISTER_OP_CUDA_KERNEL(relu2,
    71.                         paddle::operators::Relu2CUDAKernel<CUDA, double>);
    72. // 注册反向的GPU Kernel
    73. REGISTER_OP_CUDA_KERNEL(relu2_grad,
    74.                         paddle::operators::Relu2GradCUDAKernel<CUDA, float>,
    75.                         paddle::operators::Relu2GradCUDAKernel<CUDA, double>);

    注意点:

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

    自定义OP的编译

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

    编译需要include PaddlePaddle的相关头文件，如上面代码 ，需要链接PaddlePaddle的lib库。 可通过下面命令获取到:

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

    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 -DPADDLE_WITH_MKLDNN -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 -DPADDLE_WITH_MKLDNN \
    13.   -I ${include_dir} \
    14.   -I ${include_dir}/third_party \
    15.   -L /usr/local/cuda/lib64 \
    16.   -L ${lib_dir} -lpaddle_framework -lcudart
    • NVCC编译GPU OP的cu文件时，需要加 -DPADDLE_WITH_CUDA -DEIGEN_USE_GPU -DPADDLE_USE_DSO 。
    • 如果安装的PaddlePaddle不包含MKLDNN，则需要去掉编译选项-DPADDLE_WITH_MKLDNN。默认的安装包已包含MKLDNN。
    • 可多个OP编译到同一个动态库中。
    • 通过pip方式安装的PaddlePaddle由GCC 4.8编译得到，由于GCC 4.8和GCC 5以上C++11 ABI不兼容，您编写的自定义OP，需要通过GCC 4.8编译。若是GCC 5及以上的环境上使用自定义OP，推荐使用，使得编Paddle和编译自定义OP的GCC版本相同。

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

    注意点:

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

    单测测试

    可以写个简单的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])

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

    FAQ

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

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