Merge pull request #814 from NHZlX/add_template_for_op_param

Add template for op param

src/operators/batchnorm_op.h

@@ -26,14 +26,15 @@ namespace operators {
 using std::string;
 template <typename DeviceType, typename T>
 class BatchNormOp
-    : public framework::OperatorWithKernel<DeviceType, BatchNormParam,
+    : public framework::OperatorWithKernel<DeviceType,
+                                           BatchNormParam<DeviceType>,
                                            BatchNormKernel<DeviceType, T>> {
  public:
   BatchNormOp(const string &type, const VariableNameMap &inputs,
               const VariableNameMap &outputs,
               const framework::AttributeMap &attrs,
               std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, BatchNormParam,
+      : framework::OperatorWithKernel<DeviceType, BatchNormParam<DeviceType>,
                                       BatchNormKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 

src/operators/box_coder_op.h

@@ -28,20 +28,20 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class BoxCoderOp
-    : public framework::OperatorWithKernel<
-          DeviceType, BoxCoderParam, operators::BoxCoderKernel<DeviceType, T>> {
+class BoxCoderOp : public framework::OperatorWithKernel<
+                       DeviceType, BoxCoderParam<DeviceType>,
+                       operators::BoxCoderKernel<DeviceType, T>> {
  public:
   BoxCoderOp(const std::string &type, const VariableNameMap &inputs,
              const VariableNameMap &outputs,
              const framework::AttributeMap &attrs,
              std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, BoxCoderParam,
+      : framework::OperatorWithKernel<DeviceType, BoxCoderParam<DeviceType>,
                                       operators::BoxCoderKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, BoxCoderParam,
+      DeviceType, BoxCoderParam<DeviceType>,
       operators::BoxCoderKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/concat_op.h

@@ -24,19 +24,19 @@ namespace paddle_mobile {
 namespace operators {
 using std::string;
 template <typename DeviceType, typename T>
-class ConcatOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ConcatParam, operators::ConcatKernel<DeviceType, T>> {
+class ConcatOp : public framework::OperatorWithKernel<
+                     DeviceType, ConcatParam<DeviceType>,
+                     operators::ConcatKernel<DeviceType, T>> {
  public:
   ConcatOp(const string &type, const VariableNameMap &inputs,
            const VariableNameMap &outputs, const framework::AttributeMap &attrs,
            std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ConcatParam,
+      : framework::OperatorWithKernel<DeviceType, ConcatParam<DeviceType>,
                                       operators::ConcatKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ConcatParam,
+      DeviceType, ConcatParam<DeviceType>,
       operators::ConcatKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/conv_op.h

@@ -24,19 +24,19 @@ namespace paddle_mobile {
 namespace operators {
 using std::string;
 template <typename DeviceType, typename T>
-class ConvOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ConvParam, operators::ConvKernel<DeviceType, T>> {
+class ConvOp : public framework::OperatorWithKernel<
+                   DeviceType, ConvParam<DeviceType>,
+                   operators::ConvKernel<DeviceType, T>> {
  public:
   ConvOp(const std::string &type, const VariableNameMap &inputs,
          const VariableNameMap &outputs, const framework::AttributeMap &attrs,
          std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ConvParam,
+      : framework::OperatorWithKernel<DeviceType, ConvParam<DeviceType>,
                                       operators::ConvKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ConvParam,
+      DeviceType, ConvParam<DeviceType>,
       operators::ConvKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/conv_transpose_op.h

@@ -26,7 +26,7 @@ namespace paddle_mobile {
 namespace operators {
 template <typename DeviceType, typename T>
 class ConvOpTranspose : public framework::OperatorWithKernel<
-                            DeviceType, ConvTransposeParam,
+                            DeviceType, ConvTransposeParam<DeviceType>,
                             operators::ConvTransposeKernel<DeviceType, T>> {
  public:
   ConvOpTranspose(const std::string &type, const VariableNameMap &inputs,
@@ -34,7 +34,7 @@ class ConvOpTranspose : public framework::OperatorWithKernel<
                   const framework::AttributeMap &attrs,
                   std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, ConvTransposeParam,
+            DeviceType, ConvTransposeParam<DeviceType>,
             operators::ConvTransposeKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 

src/operators/depthwise_conv_op.h

@@ -25,7 +25,7 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class DepthwiseConvOp : public framework::OperatorWithKernel<
-                            DeviceType, ConvParam,
+                            DeviceType, ConvParam<DeviceType>,
                             operators::DepthwiseConvKernel<DeviceType, T>> {
  public:
   DepthwiseConvOp(const std::string &type, const VariableNameMap &inputs,
@@ -33,12 +33,12 @@ class DepthwiseConvOp : public framework::OperatorWithKernel<
                   const framework::AttributeMap &attrs,
                   std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, ConvParam,
+            DeviceType, ConvParam<DeviceType>,
             operators::DepthwiseConvKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ConvParam,
+      DeviceType, ConvParam<DeviceType>,
       operators::DepthwiseConvKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/dropout_op.h

@@ -28,18 +28,18 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class DropoutOp
-    : public framework::OperatorWithKernel<
-          DeviceType, DropoutParam, operators::DropoutKernel<DeviceType, T>> {
+class DropoutOp : public framework::OperatorWithKernel<
+                      DeviceType, DropoutParam<DeviceType>,
+                      operators::DropoutKernel<DeviceType, T>> {
  public:
   DropoutOp(const std::string &type, const VariableNameMap &inputs,
             const VariableNameMap &outputs, const framework::AttributeMap attrs,
             std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, DropoutParam,
+      : framework::OperatorWithKernel<DeviceType, DropoutParam<DeviceType>,
                                       operators::DropoutKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
-  // using framework::OperatorWithKernel<DeviceType, DropoutParam,
+  // using framework::OperatorWithKernel<DeviceType, DropoutParam<DeviceType>,
   //                                    operators::DropoutKernel<DeviceType,
   //                                    T>>;
   void InferShape() const override;

src/operators/elementwise_add_op.h

@@ -26,7 +26,7 @@ namespace operators {
 using std::string;
 template <typename DeviceType, typename T>
 class ElementwiseAddOp : public framework::OperatorWithKernel<
-                             DeviceType, ElementwiseAddParam,
+                             DeviceType, ElementwiseAddParam<DeviceType>,
                              operators::ElementwiseAddKernel<DeviceType, T>> {
  public:
   ElementwiseAddOp(const string &type, const VariableNameMap &inputs,
@@ -34,12 +34,12 @@ class ElementwiseAddOp : public framework::OperatorWithKernel<
                    const framework::AttributeMap &attrs,
                    std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, ElementwiseAddParam,
+            DeviceType, ElementwiseAddParam<DeviceType>,
             operators::ElementwiseAddKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ElementwiseAddParam,
+      DeviceType, ElementwiseAddParam<DeviceType>,
       operators::ElementwiseAddKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/feed_op.h

@@ -70,7 +70,7 @@ class FeedOp : public framework::OperatorBase<DeviceType> {
 #endif
 
  protected:
-  FeedParam param_;
+  FeedParam<DeviceType> param_;
 };
 
 }  // namespace operators

src/operators/fetch_op.h

@@ -41,7 +41,7 @@ class FetchOp : public framework::OperatorBase<DeviceType> {
   }
 
  protected:
-  FetchParam param_;
+  FetchParam<DeviceType> param_;
 };
 
 }  // namespace operators

src/operators/fusion_conv_add.h

@@ -45,19 +45,20 @@ class FusionConvAddMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionConvAddOp : public framework::OperatorWithKernel<
-                            DeviceType, FusionConvAddParam,
+                            DeviceType, FusionConvAddParam<DeviceType>,
                             operators::ConvAddKernel<DeviceType, T>> {
  public:
   FusionConvAddOp(const string &type, const VariableNameMap &inputs,
                   const VariableNameMap &outputs,
                   const framework::AttributeMap &attrs,
                   std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, FusionConvAddParam,
+      : framework::OperatorWithKernel<DeviceType,
+                                      FusionConvAddParam<DeviceType>,
                                       operators::ConvAddKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionConvAddParam,
+      DeviceType, FusionConvAddParam<DeviceType>,
       operators::ConvAddKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_conv_add_bn_op.h

@@ -53,7 +53,7 @@ class FusionConvAddBNMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionConvAddBNOp : public framework::OperatorWithKernel<
-                              DeviceType, FusionConvAddBNParam,
+                              DeviceType, FusionConvAddBNParam<DeviceType>,
                               operators::ConvAddBNKernel<DeviceType, T>> {
  public:
   FusionConvAddBNOp(const string &type, const VariableNameMap &inputs,
@@ -61,7 +61,7 @@ class FusionConvAddBNOp : public framework::OperatorWithKernel<
                     const framework::AttributeMap &attrs,
                     std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionConvAddBNParam,
+            DeviceType, FusionConvAddBNParam<DeviceType>,
             operators::ConvAddBNKernel<DeviceType, T>>(type, inputs, outputs,
                                                        attrs, scope) {}
 

src/operators/fusion_conv_add_bn_relu_op.h

@@ -55,7 +55,7 @@ class FusionConvAddBNReluMatcher : public framework::FusionOpMatcher {
 template <typename DeviceType, typename T>
 class FusionConvAddBNReluOp
     : public framework::OperatorWithKernel<
-          DeviceType, FusionConvAddBNReluParam,
+          DeviceType, FusionConvAddBNReluParam<DeviceType>,
           operators::ConvAddBNReluKernel<DeviceType, T>> {
  public:
   FusionConvAddBNReluOp(const string &type, const VariableNameMap &inputs,
@@ -63,12 +63,12 @@ class FusionConvAddBNReluOp
                         const framework::AttributeMap &attrs,
                         std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionConvAddBNReluParam,
+            DeviceType, FusionConvAddBNReluParam<DeviceType>,
             operators::ConvAddBNReluKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionConvAddBNReluParam,
+      DeviceType, FusionConvAddBNReluParam<DeviceType>,
       operators::ConvAddBNReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_conv_add_relu_op.h

@@ -45,7 +45,7 @@ class FusionConvAddReluOpMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionConvAddReluOp : public framework::OperatorWithKernel<
-                                DeviceType, FusionConvAddReluParam,
+                                DeviceType, FusionConvAddReluParam<DeviceType>,
                                 operators::ConvAddReluKernel<DeviceType, T>> {
  public:
   FusionConvAddReluOp(const string &type, const VariableNameMap &inputs,
@@ -53,12 +53,12 @@ class FusionConvAddReluOp : public framework::OperatorWithKernel<
                       const framework::AttributeMap &attrs,
                       std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionConvAddReluParam,
+            DeviceType, FusionConvAddReluParam<DeviceType>,
             operators::ConvAddReluKernel<DeviceType, T>>(type, inputs, outputs,
                                                          attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionConvAddReluParam,
+      DeviceType, FusionConvAddReluParam<DeviceType>,
       operators::ConvAddReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_conv_bn_op.h

@@ -51,14 +51,14 @@ class FusionConvBNMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionConvBNOp : public framework::OperatorWithKernel<
-                           DeviceType, FusionConvBNParam,
+                           DeviceType, FusionConvBNParam<DeviceType>,
                            operators::ConvBNKernel<DeviceType, T>> {
  public:
   FusionConvBNOp(const string &type, const VariableNameMap &inputs,
                  const VariableNameMap &outputs,
                  const framework::AttributeMap &attrs,
                  std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, FusionConvBNParam,
+      : framework::OperatorWithKernel<DeviceType, FusionConvBNParam<DeviceType>,
                                       operators::ConvBNKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 

src/operators/fusion_conv_bn_relu_op.h

@@ -52,7 +52,7 @@ class FusionConvBNReluMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionConvBNReluOp : public framework::OperatorWithKernel<
-                               DeviceType, FusionConvBNReluParam,
+                               DeviceType, FusionConvBNReluParam<DeviceType>,
                                operators::ConvBNReluKernel<DeviceType, T>> {
  public:
   FusionConvBNReluOp(const string &type, const VariableNameMap &inputs,
@@ -60,12 +60,12 @@ class FusionConvBNReluOp : public framework::OperatorWithKernel<
                      const framework::AttributeMap &attrs,
                      std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionConvBNReluParam,
+            DeviceType, FusionConvBNReluParam<DeviceType>,
             operators::ConvBNReluKernel<DeviceType, T>>(type, inputs, outputs,
                                                         attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionConvBNReluParam,
+      DeviceType, FusionConvBNReluParam<DeviceType>,
       operators::ConvBNReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_dwconv_bn_relu_op.h

@@ -51,8 +51,9 @@ class FusionDWConvBNReluMatcher : public framework::FusionOpMatcher {
 };
 
 template <typename DeviceType, typename T>
-class FusionDWConvBNReluOp : public framework::OperatorWithKernel<
-                                 DeviceType, FusionDWConvBNReluParam,
+class FusionDWConvBNReluOp
+    : public framework::OperatorWithKernel<
+          DeviceType, FusionDWConvBNReluParam<DeviceType>,
           operators::DWConvBNReluKernel<DeviceType, T>> {
  public:
   FusionDWConvBNReluOp(const string &type, const VariableNameMap &inputs,
@@ -60,12 +61,12 @@ class FusionDWConvBNReluOp : public framework::OperatorWithKernel<
                        const framework::AttributeMap &attrs,
                        std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionDWConvBNReluParam,
+            DeviceType, FusionDWConvBNReluParam<DeviceType>,
             operators::DWConvBNReluKernel<DeviceType, T>>(type, inputs, outputs,
                                                           attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionDWConvBNReluParam,
+      DeviceType, FusionDWConvBNReluParam<DeviceType>,
       operators::DWConvBNReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_elementwise_add_relu_op.h

@@ -44,7 +44,7 @@ class FusioneElementwiseAddReluMatcher : public framework::FusionOpMatcher {
 template <typename DeviceType, typename T>
 class FusionElementwiseAddReluOp
     : public framework::OperatorWithKernel<
-          DeviceType, ElementwiseAddReluParam,
+          DeviceType, ElementwiseAddReluParam<DeviceType>,
           operators::ElementwiseAddReluKernel<DeviceType, T>> {
  public:
   FusionElementwiseAddReluOp(const string &type, const VariableNameMap &inputs,
@@ -52,7 +52,7 @@ class FusionElementwiseAddReluOp
                              const framework::AttributeMap &attrs,
                              std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, ElementwiseAddReluParam,
+            DeviceType, ElementwiseAddReluParam<DeviceType>,
             operators::ElementwiseAddReluKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 

src/operators/fusion_fc_op.h

@@ -45,20 +45,20 @@ class FusionFcMatcher : public framework::FusionOpMatcher {
 };
 
 template <typename DeviceType, typename T>
-class FusionFcOp
-    : public framework::OperatorWithKernel<
-          DeviceType, FusionFcParam, operators::FusionFcKernel<DeviceType, T>> {
+class FusionFcOp : public framework::OperatorWithKernel<
+                       DeviceType, FusionFcParam<DeviceType>,
+                       operators::FusionFcKernel<DeviceType, T>> {
  public:
   FusionFcOp(const string &type, const VariableNameMap &inputs,
              const VariableNameMap &outputs,
              const framework::AttributeMap &attrs,
              std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, FusionFcParam,
+      : framework::OperatorWithKernel<DeviceType, FusionFcParam<DeviceType>,
                                       operators::FusionFcKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionFcParam,
+      DeviceType, FusionFcParam<DeviceType>,
       operators::FusionFcKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/fusion_fc_relu_op.h

@@ -44,7 +44,7 @@ class FusionFcReluMatcher : public framework::FusionOpMatcher {
 
 template <typename DeviceType, typename T>
 class FusionFcReluOp : public framework::OperatorWithKernel<
-                           DeviceType, FusionFcReluParam,
+                           DeviceType, FusionFcReluParam<DeviceType>,
                            operators::FusionFcReluKernel<DeviceType, T>> {
  public:
   FusionFcReluOp(const string &type, const VariableNameMap &inputs,
@@ -52,12 +52,12 @@ class FusionFcReluOp : public framework::OperatorWithKernel<
                  const framework::AttributeMap &attrs,
                  std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, FusionFcReluParam,
+            DeviceType, FusionFcReluParam<DeviceType>,
             operators::FusionFcReluKernel<DeviceType, T>>(type, inputs, outputs,
                                                           attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, FusionFcReluParam,
+      DeviceType, FusionFcReluParam<DeviceType>,
       operators::FusionFcReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/im2sequence_op.h

@@ -27,7 +27,7 @@ using namespace framework;
 
 template <typename DeviceType, typename T>
 class Im2SequenceOp : public framework::OperatorWithKernel<
-                          DeviceType, Im2SequenceParam,
+                          DeviceType, Im2SequenceParam<DeviceType>,
                           operators::Im2SequenceKernel<DeviceType, T>> {
  public:
   Im2SequenceOp(const std::string &type, const VariableNameMap &inputs,
@@ -35,12 +35,12 @@ class Im2SequenceOp : public framework::OperatorWithKernel<
                 const framework::AttributeMap &attrs,
                 std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, Im2SequenceParam,
+            DeviceType, Im2SequenceParam<DeviceType>,
             operators::Im2SequenceKernel<DeviceType, T>>(type, inputs, outputs,
                                                          attrs, scope) {}
 
   // using framework::OperatorWithKernel<
-  //    DeviceType, Im2SequenceParam,
+  //    DeviceType, Im2SequenceParam<DeviceType>,
   //    operators::Im2SequenceKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/kernel/arm/batchnorm_kernel.cpp

@@ -21,12 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool BatchNormKernel<CPU, float>::Init(BatchNormParam *param) {
+bool BatchNormKernel<CPU, float>::Init(BatchNormParam<CPU> *param) {
   return true;
 }
 
 template <>
-void BatchNormKernel<CPU, float>::Compute(const BatchNormParam &param) const {
+void BatchNormKernel<CPU, float>::Compute(
+    const BatchNormParam<CPU> &param) const {
   BatchnormCompute<float>(param);
 }
 

src/operators/kernel/arm/box_coder_kernel.cpp

@@ -21,12 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool BoxCoderKernel<CPU, float>::Init(BoxCoderParam *param) {
+bool BoxCoderKernel<CPU, float>::Init(BoxCoderParam<CPU> *param) {
   return true;
 }
 
 template <>
-void BoxCoderKernel<CPU, float>::Compute(const BoxCoderParam &param) const {
+void BoxCoderKernel<CPU, float>::Compute(
+    const BoxCoderParam<CPU> &param) const {
   BoxCoderCompute<float>(param);
 }
 

src/operators/kernel/arm/concat_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConcatKernel<CPU, float>::Init(ConcatParam *param) {
+bool ConcatKernel<CPU, float>::Init(ConcatParam<CPU> *param) {
   return true;
 }
 
 template <>
-void ConcatKernel<CPU, float>::Compute(const ConcatParam &param) const {
+void ConcatKernel<CPU, float>::Compute(const ConcatParam<CPU> &param) const {
   ConcatCompute<float>(param);
 }
 

src/operators/kernel/arm/conv_add_bn_relu_kernel.cpp

@@ -21,7 +21,8 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddBNReluKernel<CPU, float>::Init(FusionConvAddBNReluParam *param) {
+bool ConvAddBNReluKernel<CPU, float>::Init(
+    FusionConvAddBNReluParam<CPU> *param) {
   const Tensor *mean = param->InputMean();
   const Tensor *variance = param->InputVariance();
   const Tensor *scale = param->InputScale();
@@ -54,7 +55,7 @@ bool ConvAddBNReluKernel<CPU, float>::Init(FusionConvAddBNReluParam *param) {
 
 template <>
 void ConvAddBNReluKernel<CPU, float>::Compute(
-    const FusionConvAddBNReluParam &param) const {
+    const FusionConvAddBNReluParam<CPU> &param) const {
   ConvAddBNReluCompute<float>(param);
 }
 template class ConvAddBNReluKernel<CPU, float>;

src/operators/kernel/arm/conv_add_kernel.cpp

@@ -20,12 +20,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddKernel<CPU, float>::Init(FusionConvAddParam *param) {
+bool ConvAddKernel<CPU, float>::Init(FusionConvAddParam<CPU> *param) {
   return true;
 }
 
 template <>
-void ConvAddKernel<CPU, float>::Compute(const FusionConvAddParam &param) const {
+void ConvAddKernel<CPU, float>::Compute(
+    const FusionConvAddParam<CPU> &param) const {
   ConvAddCompute<float>(param);
 }
 

src/operators/kernel/arm/conv_add_relu_kernel.cpp

@@ -21,13 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddReluKernel<CPU, float>::Init(FusionConvAddReluParam *param) {
+bool ConvAddReluKernel<CPU, float>::Init(FusionConvAddReluParam<CPU> *param) {
   return true;
 }
 
 template <>
 void ConvAddReluKernel<CPU, float>::Compute(
-    const FusionConvAddReluParam &param) const {
+    const FusionConvAddReluParam<CPU> &param) const {
   ConvAddReluCompute<float>(param);
 }
 template class ConvAddReluKernel<CPU, float>;

src/operators/kernel/arm/conv_bn_relu_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvBNReluKernel<CPU, float>::Init(FusionConvBNReluParam *param) {
+bool ConvBNReluKernel<CPU, float>::Init(FusionConvBNReluParam<CPU> *param) {
   const Tensor *mean = param->InputMean();
   const Tensor *variance = param->InputVariance();
   const Tensor *scale = param->InputScale();
@@ -57,7 +57,7 @@ bool ConvBNReluKernel<CPU, float>::Init(FusionConvBNReluParam *param) {
 
 template <>
 void ConvBNReluKernel<CPU, float>::Compute(
-    const FusionConvBNReluParam &param) const {
+    const FusionConvBNReluParam<CPU> &param) const {
   ConvBNReluCompute<float>(param);
 }
 template class ConvBNReluKernel<CPU, float>;

src/operators/kernel/arm/conv_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvKernel<CPU, float>::Init(ConvParam *param) {
+bool ConvKernel<CPU, float>::Init(ConvParam<CPU> *param) {
   return true;
 }
 
 template <>
-void ConvKernel<CPU, float>::Compute(const ConvParam &param) const {
+void ConvKernel<CPU, float>::Compute(const ConvParam<CPU> &param) const {
   ConvCompute<float>(param);
 }
 

src/operators/kernel/arm/conv_transpose_kernel.cpp

@@ -21,13 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvTransposeKernel<CPU, float>::Init(ConvTransposeParam *param) {
+bool ConvTransposeKernel<CPU, float>::Init(ConvTransposeParam<CPU> *param) {
   return true;
 }
 
 template <>
 void ConvTransposeKernel<CPU, float>::Compute(
-    const ConvTransposeParam &param) const {
+    const ConvTransposeParam<CPU> &param) const {
   ConvTransposeCompute<float>(param);
 }
 

src/operators/kernel/arm/depthwise_conv_kernel.cpp

@@ -21,12 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool DepthwiseConvKernel<CPU, float>::Init(ConvParam *param) {
+bool DepthwiseConvKernel<CPU, float>::Init(ConvParam<CPU> *param) {
   return true;
 }
 
 template <>
-void DepthwiseConvKernel<CPU, float>::Compute(const ConvParam &param) const {
+void DepthwiseConvKernel<CPU, float>::Compute(
+    const ConvParam<CPU> &param) const {
   DepthwiseConvCompute<float>(param);
 }
 

src/operators/kernel/arm/dropout_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool DropoutKernel<CPU, float>::Init(DropoutParam *para) {
+bool DropoutKernel<CPU, float>::Init(DropoutParam<CPU> *para) {
   return true;
 }
 
@@ -31,7 +31,7 @@ struct DropoutFunctor {
 };
 
 template <>
-void DropoutKernel<CPU, float>::Compute(const DropoutParam &param) const {
+void DropoutKernel<CPU, float>::Compute(const DropoutParam<CPU> &param) const {
   const auto *input_x = param.InputX();
   auto *input_x_ptr = input_x->data<float>();
   auto *out = param.Out();

src/operators/kernel/arm/dwconv_bn_relu_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool DWConvBNReluKernel<CPU, float>::Init(FusionDWConvBNReluParam *param) {
+bool DWConvBNReluKernel<CPU, float>::Init(FusionDWConvBNReluParam<CPU> *param) {
   const Tensor *mean = param->InputMean();
   const Tensor *variance = param->InputVariance();
   const Tensor *scale = param->InputScale();
@@ -54,7 +54,7 @@ bool DWConvBNReluKernel<CPU, float>::Init(FusionDWConvBNReluParam *param) {
 
 template <>
 void DWConvBNReluKernel<CPU, float>::Compute(
-    const FusionDWConvBNReluParam &param) const {
+    const FusionDWConvBNReluParam<CPU> &param) const {
   DWConvBNReluCompute<float>(param);
 }
 template class DWConvBNReluKernel<CPU, float>;

src/operators/kernel/arm/elementwise_add_kernel.cpp

@@ -21,13 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ElementwiseAddKernel<CPU, float>::Init(ElementwiseAddParam *param) {
+bool ElementwiseAddKernel<CPU, float>::Init(ElementwiseAddParam<CPU> *param) {
   return true;
 }
 
 template <>
 void ElementwiseAddKernel<CPU, float>::Compute(
-    const ElementwiseAddParam &param) const {
+    const ElementwiseAddParam<CPU> &param) const {
   ElementwiseAddCompute<float>(param);
 }
 

src/operators/kernel/arm/fusion_fc_kernel.cpp

@@ -21,12 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool FusionFcKernel<CPU, float>::Init(FusionFcParam *param) {
+bool FusionFcKernel<CPU, float>::Init(FusionFcParam<CPU> *param) {
   return true;
 }
 
 template <>
-void FusionFcKernel<CPU, float>::Compute(const FusionFcParam &param) const {
+void FusionFcKernel<CPU, float>::Compute(
+    const FusionFcParam<CPU> &param) const {
   FusionFcCompute<float>(param);
 }
 

src/operators/kernel/arm/im2sequence_kernel.cpp

@@ -20,7 +20,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool Im2SequenceKernel<CPU, float>::Init(Im2SequenceParam *para) {
+bool Im2SequenceKernel<CPU, float>::Init(Im2SequenceParam<CPU> *para) {
   return true;
 }
 
@@ -33,7 +33,7 @@ inline int Im2SeqOutputSize(int input_size, int filter_size, int padding_0,
 
 template <>
 void Im2SequenceKernel<CPU, float>::Compute(
-    const Im2SequenceParam &param) const {
+    const Im2SequenceParam<CPU> &param) const {
   const Tensor *in_x = param.Input();
   Tensor *out = param.Output();
   out->mutable_data<float>();

src/operators/kernel/arm/lrn_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool LrnKernel<CPU, float>::Init(LrnParam *param) {
+bool LrnKernel<CPU, float>::Init(LrnParam<CPU> *param) {
   return true;
 }
 
 template <>
-void LrnKernel<CPU, float>::Compute(const LrnParam &param) const {
+void LrnKernel<CPU, float>::Compute(const LrnParam<CPU> &param) const {
   LrnCompute<float>(param);
 }
 

src/operators/kernel/arm/mul_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool MulKernel<CPU, float>::Init(MulParam *param) {
+bool MulKernel<CPU, float>::Init(MulParam<CPU> *param) {
   return true;
 }
 
 template <>
-void MulKernel<CPU, float>::Compute(const MulParam &param) const {
+void MulKernel<CPU, float>::Compute(const MulParam<CPU> &param) const {
   MulCompute<float>(param);
 }
 

src/operators/kernel/arm/multiclass_nms_kernel.cpp

@@ -21,13 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool MultiClassNMSKernel<CPU, float>::Init(MultiClassNMSParam *param) {
+bool MultiClassNMSKernel<CPU, float>::Init(MultiClassNMSParam<CPU> *param) {
   return true;
 }
 
 template <>
 void MultiClassNMSKernel<CPU, float>::Compute(
-    const MultiClassNMSParam &param) const {
+    const MultiClassNMSParam<CPU> &param) const {
   MultiClassNMSCompute<float>(param);
 }
 

src/operators/kernel/arm/pool_kernel.cpp

@@ -20,12 +20,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool PoolKernel<CPU, float>::Init(PoolParam *param) {
+bool PoolKernel<CPU, float>::Init(PoolParam<CPU> *param) {
   return true;
 }
 
 template <>
-void PoolKernel<CPU, float>::Compute(const PoolParam &param) const {
+void PoolKernel<CPU, float>::Compute(const PoolParam<CPU> &param) const {
   PoolCompute<float>(param);
 }
 }  // namespace operators

src/operators/kernel/arm/prelu_kernel.cpp

@@ -32,7 +32,7 @@ struct PReluFunctor {
  * @b 特化到具体平台的实现, param 从 op 层传入
  * */
 template <>
-void PReluKernel<CPU, float>::Compute(const PReluParam &param) const {
+void PReluKernel<CPU, float>::Compute(const PReluParam<CPU> &param) const {
   auto *x = param.InputX();
   auto *alpha = param.InputAlpha();
   auto *out = param.Out();

src/operators/kernel/arm/prior_box_kernel.cpp

@@ -21,12 +21,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool PriorBoxKernel<CPU, float>::Init(PriorBoxParam *param) {
+bool PriorBoxKernel<CPU, float>::Init(PriorBoxParam<CPU> *param) {
   return true;
 }
 
 template <>
-void PriorBoxKernel<CPU, float>::Compute(const PriorBoxParam &param) const {
+void PriorBoxKernel<CPU, float>::Compute(
+    const PriorBoxParam<CPU> &param) const {
   PriorBoxCompute<float>(param);
 }
 

src/operators/kernel/arm/relu_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ReluKernel<CPU, float>::Init(ReluParam *param) {
+bool ReluKernel<CPU, float>::Init(ReluParam<CPU> *param) {
   return true;
 }
 
 template <>
-void ReluKernel<CPU, float>::Compute(const ReluParam &param) const {
+void ReluKernel<CPU, float>::Compute(const ReluParam<CPU> &param) const {
   ReluCompute<float>(param);
 }
 

src/operators/kernel/arm/reshape_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ReshapeKernel<CPU, float>::Init(ReshapeParam *param) {
+bool ReshapeKernel<CPU, float>::Init(ReshapeParam<CPU> *param) {
   return true;
 }
 
 template <>
-void ReshapeKernel<CPU, float>::Compute(const ReshapeParam &param) const {
+void ReshapeKernel<CPU, float>::Compute(const ReshapeParam<CPU> &param) const {
   ReshapeCompute<float>(param);
 }
 

src/operators/kernel/arm/resize_kernel.cpp

@@ -108,7 +108,7 @@ void ResizeTensor(const Tensor* src, Tensor* dst) {
 }
 
 template <>
-void ResizeKernel<CPU, float>::Compute(const ResizeParam& param) const {
+void ResizeKernel<CPU, float>::Compute(const ResizeParam<CPU>& param) const {
   const auto* input_x = param.InputX();
   const auto& input_x_dims = input_x->dims();
   auto* out = param.Out();

src/operators/kernel/arm/scale_kernel.cpp

@@ -23,7 +23,7 @@ namespace operators {
  * @b 特化到具体平台的实现, param 从 op 层传入
  * */
 template <>
-void ScaleKernel<CPU, float>::Compute(const ScaleParam &param) const {
+void ScaleKernel<CPU, float>::Compute(const ScaleParam<CPU> &param) const {
   const auto *input_x = param.InputX();
   auto *input_x_ptr = input_x->data<float>();
   auto *out = param.Out();

src/operators/kernel/arm/sigmoid_kernel.cpp

@@ -27,12 +27,12 @@ using framework::DDim;
 using framework::Tensor;
 
 template <>
-bool SigmoidKernel<CPU, float>::Init(SigmoidParam *param) {
+bool SigmoidKernel<CPU, float>::Init(SigmoidParam<CPU> *param) {
   return true;
 }
 
 template <>
-void SigmoidKernel<CPU, float>::Compute(const SigmoidParam &param) const {
+void SigmoidKernel<CPU, float>::Compute(const SigmoidParam<CPU> &param) const {
   SigmoidCompute<float>(param);
 }
 

src/operators/kernel/arm/softmax_kernel.cpp

@@ -21,12 +21,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool SoftmaxKernel<CPU, float>::Init(SoftmaxParam *param) {
+bool SoftmaxKernel<CPU, float>::Init(SoftmaxParam<CPU> *param) {
   return true;
 }
 
 template <>
-void SoftmaxKernel<CPU, float>::Compute(const SoftmaxParam &param) const {
+void SoftmaxKernel<CPU, float>::Compute(const SoftmaxParam<CPU> &param) const {
   SoftmaxCompute<float>(param);
 }
 

src/operators/kernel/arm/transpose_kernel.cpp

@@ -20,12 +20,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool TransposeKernel<CPU, float>::Init(TransposeParam *param) {
+bool TransposeKernel<CPU, float>::Init(TransposeParam<CPU> *param) {
   return true;
 }
 
 template <>
-void TransposeKernel<CPU, float>::Compute(const TransposeParam &param) const {
+void TransposeKernel<CPU, float>::Compute(
+    const TransposeParam<CPU> &param) const {
   TransposeCompute<float>(param);
 }
 

src/operators/kernel/batchnorm_kernel.h

@@ -26,10 +26,10 @@ using namespace framework;
 
 template <typename DeviceType, typename T>
 class BatchNormKernel
-    : public framework::OpKernelBase<DeviceType, BatchNormParam> {
+    : public framework::OpKernelBase<DeviceType, BatchNormParam<DeviceType>> {
  public:
-  void Compute(const BatchNormParam &param) const;
-  bool Init(BatchNormParam *param);
+  void Compute(const BatchNormParam<DeviceType> &param) const;
+  bool Init(BatchNormParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/box_coder_kernel.h

@@ -27,10 +27,10 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class BoxCoderKernel
-    : public framework::OpKernelBase<DeviceType, BoxCoderParam> {
+    : public framework::OpKernelBase<DeviceType, BoxCoderParam<DeviceType>> {
  public:
-  void Compute(const BoxCoderParam& param) const;
-  bool Init(BoxCoderParam* param);
+  void Compute(const BoxCoderParam<DeviceType>& param) const;
+  bool Init(BoxCoderParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/central-arm-func/batchnorm_arm_func.h

@@ -23,7 +23,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void BatchnormCompute(const BatchNormParam &param) {
+void BatchnormCompute(const BatchNormParam<CPU> &param) {
   const Tensor *input_x = param.InputX();
   auto input_x_ptr = input_x->data<float>();
   const auto &x_dims = input_x->dims();

src/operators/kernel/central-arm-func/box_coder_arm_func.h

@@ -113,7 +113,7 @@ void DecodeCenterSize(const framework::Tensor& target_box,
 }
 
 template <typename P>
-void BoxCoderCompute(const BoxCoderParam& param) {
+void BoxCoderCompute(const BoxCoderParam<CPU>& param) {
   const auto* input_priorbox = param.InputPriorBox();
   const auto* input_priorboxvar = param.InputPriorBoxVar();
   const auto* input_targetbox = param.InputTargetBox();

src/operators/kernel/central-arm-func/concat_arm_func.h

@@ -54,7 +54,7 @@ class ConcatFunctor {
 };
 
 template <typename P>
-void ConcatCompute(const ConcatParam &param) {
+void ConcatCompute(const ConcatParam<CPU> &param) {
   auto inputs = param.Inputs();
   auto *out = param.Out();
   int64_t axis = param.Axis();

src/operators/kernel/central-arm-func/conv_add_arm_func.h

@@ -25,7 +25,7 @@ limitations under the License. */
 
 namespace paddle_mobile {
 namespace operators {
-void ConvAddBasic(const FusionConvAddParam &param) {
+void ConvAddBasic(const FusionConvAddParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor bias = *param.Bias();
@@ -114,7 +114,7 @@ void ConvAddBasic(const FusionConvAddParam &param) {
 }
 
 template <typename P>
-void ConvAddCompute(const FusionConvAddParam &param) {
+void ConvAddCompute(const FusionConvAddParam<CPU> &param) {
   if (param.Groups() == param.Input()->dims()[1] &&
       param.Input()->dims()[1] == param.Output()->dims()[1] &&
       param.Filter()->dims()[2] == param.Filter()->dims()[3] &&

src/operators/kernel/central-arm-func/conv_add_bn_relu_arm_func.h

@@ -25,7 +25,7 @@ limitations under the License. */
 
 namespace paddle_mobile {
 namespace operators {
-void ConvAddBNReluBasic(const FusionConvAddBNReluParam &param) {
+void ConvAddBNReluBasic(const FusionConvAddBNReluParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor new_bias = *param.NewBias();
@@ -112,7 +112,7 @@ void ConvAddBNReluBasic(const FusionConvAddBNReluParam &param) {
   }
 }
 template <typename P>
-void ConvAddBNReluCompute(const FusionConvAddBNReluParam &param) {
+void ConvAddBNReluCompute(const FusionConvAddBNReluParam<CPU> &param) {
   Tensor Bias;
   Bias.mutable_data<float>({param.Groups()});
   if (param.Groups() == param.Input()->dims()[1] &&

src/operators/kernel/central-arm-func/conv_add_relu_arm_func.h

@@ -26,7 +26,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void ConvAddReluCompute(const FusionConvAddReluParam &param) {
+void ConvAddReluCompute(const FusionConvAddReluParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor bias = *param.Bias();

src/operators/kernel/central-arm-func/conv_arm_func.h

@@ -25,7 +25,7 @@ limitations under the License. */
 
 namespace paddle_mobile {
 namespace operators {
-inline void ConvBasic(const ConvParam &param) {
+inline void ConvBasic(const ConvParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor *output = param.Output();
@@ -112,7 +112,7 @@ inline void ConvBasic(const ConvParam &param) {
 }
 
 template <typename P>
-void ConvCompute(const ConvParam &param) {
+void ConvCompute(const ConvParam<CPU> &param) {
   if (param.Groups() == param.Input()->dims()[1] &&
       param.Input()->dims()[1] == param.Output()->dims()[1] &&
       param.Filter()->dims()[2] == param.Filter()->dims()[3] &&

src/operators/kernel/central-arm-func/conv_bn_relu_arm_func.h

@@ -23,7 +23,7 @@ limitations under the License. */
 #include "operators/op_param.h"
 namespace paddle_mobile {
 namespace operators {
-void ConvBNReluBasic(const FusionConvBNReluParam &param) {
+void ConvBNReluBasic(const FusionConvBNReluParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor new_bias = *param.NewBias();
@@ -113,7 +113,7 @@ void ConvBNReluBasic(const FusionConvBNReluParam &param) {
 }
 
 template <typename P>
-void ConvBNReluCompute(const FusionConvBNReluParam &param) {
+void ConvBNReluCompute(const FusionConvBNReluParam<CPU> &param) {
   if (param.Groups() == param.Input()->dims()[1] &&
       param.Input()->dims()[1] == param.Output()->dims()[1] &&
       param.Filter()->dims()[2] == param.Filter()->dims()[3] &&

src/operators/kernel/central-arm-func/conv_transpose_arm_func.h

@@ -28,7 +28,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void ConvTransposeCompute(const ConvTransposeParam &param) {
+void ConvTransposeCompute(const ConvTransposeParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor *output = param.Output();

src/operators/kernel/central-arm-func/depthwise_conv_arm_func.h

@@ -25,7 +25,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void DepthwiseConvCompute(const ConvParam &param) {
+void DepthwiseConvCompute(const ConvParam<CPU> &param) {
   Tensor Bias;
   Bias.mutable_data<float>({param.Groups()});
   if (param.Groups() == param.Input()->dims()[1] &&

src/operators/kernel/central-arm-func/dwconv_bn_relu_arm_func.h

@@ -23,7 +23,7 @@ limitations under the License. */
 #include "operators/op_param.h"
 namespace paddle_mobile {
 namespace operators {
-void DWConvBNReluBasic(const FusionDWConvBNReluParam &param) {
+void DWConvBNReluBasic(const FusionDWConvBNReluParam<CPU> &param) {
   const Tensor *input = param.Input();
   Tensor filter = *param.Filter();
   Tensor new_bias = *param.NewBias();
@@ -111,7 +111,7 @@ void DWConvBNReluBasic(const FusionDWConvBNReluParam &param) {
   }
 }
 template <typename P>
-void DWConvBNReluCompute(const FusionDWConvBNReluParam &param) {
+void DWConvBNReluCompute(const FusionDWConvBNReluParam<CPU> &param) {
   if (param.Groups() == param.Input()->dims()[1] &&
       param.Input()->dims()[1] == param.Output()->dims()[1] &&
       param.Filter()->dims()[2] == param.Filter()->dims()[3] &&

src/operators/kernel/central-arm-func/elementwise_add_arm_func.h

@@ -27,7 +27,7 @@ struct AddFunctor {
 };
 
 template <typename P>
-void ElementwiseAddCompute(const ElementwiseAddParam &param) {
+void ElementwiseAddCompute(const ElementwiseAddParam<CPU> &param) {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   Tensor *Out = param.Out();

src/operators/kernel/central-arm-func/fusion_fc_arm_func.h

@@ -22,7 +22,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void FusionFcCompute(const FusionFcParam &param) {
+void FusionFcCompute(const FusionFcParam<CPU> &param) {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   const Tensor *input_z = param.InputZ();

src/operators/kernel/central-arm-func/lrn_arm_func.h

@@ -20,7 +20,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void LrnCompute(const LrnParam &param) {
+void LrnCompute(const LrnParam<CPU> &param) {
   const Tensor *input_x = param.InputX();
   auto x_dims = input_x->dims();
   Tensor *out = param.Out();

src/operators/kernel/central-arm-func/mul_arm_func.h

@@ -54,7 +54,7 @@ namespace operators {
 // 结果x(6行4列)乘y(4行2列)，按1中矩阵相乘，结果out(6行2列)
 
 template <typename P>
-void MulCompute(const MulParam &param) {
+void MulCompute(const MulParam<CPU> &param) {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   Tensor *out = param.Out();

src/operators/kernel/central-arm-func/multiclass_nms_arm_func.h

@@ -213,7 +213,7 @@ void MultiClassOutput(const framework::Tensor& scores,
 }
 
 template <typename P>
-void MultiClassNMSCompute(const MultiClassNMSParam& param) {
+void MultiClassNMSCompute(const MultiClassNMSParam<CPU>& param) {
   const auto* input_bboxes = param.InputBBoxes();
   const auto& input_bboxes_dims = input_bboxes->dims();
 

src/operators/kernel/central-arm-func/pool_arm_func.h

@@ -38,7 +38,7 @@ inline void PoolBasic(std::string pooling_type, std::vector<int> ksize,
   }
 }
 template <typename P>
-void PoolCompute(const PoolParam &param) {
+void PoolCompute(const PoolParam<CPU> &param) {
   const Tensor *in_x = param.Input();
   Tensor *out = param.Output();
   std::string pooling_type = param.PoolingType();
@@ -59,7 +59,6 @@ void PoolCompute(const PoolParam &param) {
       ksize[i] = static_cast<int>(in_x->dims()[i + 2]);
     }
   }
-
   if (ksize[0] == 3 && ksize[0] == ksize[1]) {
     if (pooling_type == "max") {
       if (strides[0] == strides[1] && strides[0] == 1 &&

src/operators/kernel/central-arm-func/prior_box_arm_func.h

@@ -29,7 +29,7 @@ struct ClipFunctor {
 };
 
 template <typename P>
-void PriorBoxCompute(const PriorBoxParam &param) {
+void PriorBoxCompute(const PriorBoxParam<CPU> &param) {
   const auto *input_ = param.Input();
   const auto &input_dims = input_->dims();
 

src/operators/kernel/central-arm-func/relu_arm_func.h

@@ -30,7 +30,7 @@ struct ReluFunctor {
  * @b 特化到具体平台的实现, param 从 op 层传入
  * */
 template <typename P>
-void ReluCompute(const ReluParam &param) {
+void ReluCompute(const ReluParam<CPU> &param) {
   const auto *input_x = param.InputX();
   auto *input_x_ptr = input_x->data<float>();
   auto *out = param.Out();

src/operators/kernel/central-arm-func/reshape_arm_func.h

@@ -23,7 +23,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename P>
-void ReshapeCompute(const ReshapeParam &param) {
+void ReshapeCompute(const ReshapeParam<CPU> &param) {
   const auto *input_x = param.InputX();
   const auto &input_x_dims = input_x->dims();
   auto *out = param.Out();

src/operators/kernel/central-arm-func/sigmoid_arm_func.h

@@ -73,7 +73,7 @@ void sigmoid(const Tensor *X, Tensor *Y) {
 }
 
 template <typename P>
-void SigmoidCompute(const SigmoidParam &param) {
+void SigmoidCompute(const SigmoidParam<CPU> &param) {
   const Tensor *in_x = param.InputX();
   Tensor *out = param.Out();
   auto x_dims = in_x->dims();

src/operators/kernel/central-arm-func/softmax_arm_func.h

@@ -19,7 +19,7 @@ limitations under the License. */
 namespace paddle_mobile {
 namespace operators {
 template <typename P>
-void SoftmaxCompute(const SoftmaxParam &param) {
+void SoftmaxCompute(const SoftmaxParam<CPU> &param) {
   const Tensor *in_x = param.InputX();
   Tensor *out = param.Out();
   auto x_dims = in_x->dims();

src/operators/kernel/central-arm-func/transpose_arm_func.h

@@ -39,7 +39,7 @@ namespace operators {
 // }
 
 template <typename P>
-void TransposeCompute(const TransposeParam& param) {
+void TransposeCompute(const TransposeParam<CPU>& param) {
   const auto* input_x = param.InputX();
   const auto input_x_dims = input_x->dims();
   auto* out = param.Out();

src/operators/kernel/concat_kernel.h

@@ -24,10 +24,11 @@ namespace operators {
 using namespace framework;
 
 template <typename DeviceType, typename T>
-class ConcatKernel : public framework::OpKernelBase<DeviceType, ConcatParam> {
+class ConcatKernel
+    : public framework::OpKernelBase<DeviceType, ConcatParam<DeviceType>> {
  public:
-  void Compute(const ConcatParam &param) const;
-  bool Init(ConcatParam *param);
+  void Compute(const ConcatParam<DeviceType> &param) const;
+  bool Init(ConcatParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_add_bn_kernel.h

@@ -32,10 +32,11 @@ using framework::DDim;
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class ConvAddBNKernel : public OpKernelBase<DeviceType, FusionConvAddBNParam> {
+class ConvAddBNKernel
+    : public OpKernelBase<DeviceType, FusionConvAddBNParam<DeviceType>> {
  public:
-  void Compute(const FusionConvAddBNParam &param) const;
-  bool Init(FusionConvAddBNParam *param);
+  void Compute(const FusionConvAddBNParam<DeviceType> &param) const;
+  bool Init(FusionConvAddBNParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_add_bn_relu_kernel.h

@@ -33,10 +33,10 @@ using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
 class ConvAddBNReluKernel
-    : public OpKernelBase<DeviceType, FusionConvAddBNReluParam> {
+    : public OpKernelBase<DeviceType, FusionConvAddBNReluParam<DeviceType>> {
  public:
-  void Compute(const FusionConvAddBNReluParam &param) const;
-  bool Init(FusionConvAddBNReluParam *param);
+  void Compute(const FusionConvAddBNReluParam<DeviceType> &param) const;
+  bool Init(FusionConvAddBNReluParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_add_kernel.h

@@ -37,10 +37,11 @@ using framework::DDim;
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class ConvAddKernel : public OpKernelBase<DeviceType, FusionConvAddParam> {
+class ConvAddKernel
+    : public OpKernelBase<DeviceType, FusionConvAddParam<DeviceType>> {
  public:
-  void Compute(const FusionConvAddParam &param) const;
-  bool Init(FusionConvAddParam *param);
+  void Compute(const FusionConvAddParam<DeviceType> &param) const;
+  bool Init(FusionConvAddParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_add_relu_kernel.h

@@ -33,10 +33,10 @@ using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
 class ConvAddReluKernel
-    : public OpKernelBase<DeviceType, FusionConvAddReluParam> {
+    : public OpKernelBase<DeviceType, FusionConvAddReluParam<DeviceType>> {
  public:
-  void Compute(const FusionConvAddReluParam &param) const;
-  bool Init(FusionConvAddReluParam *param);
+  void Compute(const FusionConvAddReluParam<DeviceType> &param) const;
+  bool Init(FusionConvAddReluParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_bn_kernel.h

@@ -32,10 +32,11 @@ using framework::DDim;
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class ConvBNKernel : public OpKernelBase<DeviceType, FusionConvBNParam> {
+class ConvBNKernel
+    : public OpKernelBase<DeviceType, FusionConvBNParam<DeviceType>> {
  public:
-  void Compute(const FusionConvBNParam &param) const;
-  bool Init(FusionConvBNParam *param);
+  void Compute(const FusionConvBNParam<DeviceType> &param) const;
+  bool Init(FusionConvBNParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_bn_relu_kernel.h

@@ -33,10 +33,10 @@ using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
 class ConvBNReluKernel
-    : public OpKernelBase<DeviceType, FusionConvBNReluParam> {
+    : public OpKernelBase<DeviceType, FusionConvBNReluParam<DeviceType>> {
  public:
-  void Compute(const FusionConvBNReluParam &param) const;
-  bool Init(FusionConvBNReluParam *param);
+  void Compute(const FusionConvBNReluParam<DeviceType> &param) const;
+  bool Init(FusionConvBNReluParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_kernel.h

@@ -29,10 +29,10 @@ namespace operators {
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class ConvKernel : public OpKernelBase<DeviceType, ConvParam> {
+class ConvKernel : public OpKernelBase<DeviceType, ConvParam<DeviceType>> {
  public:
-  void Compute(const ConvParam &param) const;
-  bool Init(ConvParam *param);
+  void Compute(const ConvParam<DeviceType> &param) const;
+  bool Init(ConvParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/conv_transpose_kernel.h

@@ -26,11 +26,11 @@ using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
 class ConvTransposeKernel
-    : public OpKernelBase<DeviceType, ConvTransposeParam> {
+    : public OpKernelBase<DeviceType, ConvTransposeParam<DeviceType>> {
  public:
-  void Compute(const ConvTransposeParam &param) const;
+  void Compute(const ConvTransposeParam<DeviceType> &param) const;
 
-  bool Init(ConvTransposeParam *param);
+  bool Init(ConvTransposeParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/depthwise_conv_kernel.h

@@ -28,10 +28,11 @@ namespace operators {
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class DepthwiseConvKernel : public OpKernelBase<DeviceType, ConvParam> {
+class DepthwiseConvKernel
+    : public OpKernelBase<DeviceType, ConvParam<DeviceType>> {
  public:
-  void Compute(const ConvParam &param) const;
-  bool Init(ConvParam *param);
+  void Compute(const ConvParam<DeviceType> &param) const;
+  bool Init(ConvParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/dropout_kernel.h

@@ -23,10 +23,11 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename DeviceType, typename T>
-class DropoutKernel : public framework::OpKernelBase<DeviceType, DropoutParam> {
+class DropoutKernel
+    : public framework::OpKernelBase<DeviceType, DropoutParam<DeviceType>> {
  public:
-  void Compute(const DropoutParam& param) const;
-  bool Init(DropoutParam* para);
+  void Compute(const DropoutParam<DeviceType>& param) const;
+  bool Init(DropoutParam<DeviceType>* para);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/dwconv_bn_relu_kernel.h

@@ -33,10 +33,10 @@ using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
 class DWConvBNReluKernel
-    : public OpKernelBase<DeviceType, FusionDWConvBNReluParam> {
+    : public OpKernelBase<DeviceType, FusionDWConvBNReluParam<DeviceType>> {
  public:
-  void Compute(const FusionDWConvBNReluParam &param) const;
-  bool Init(FusionDWConvBNReluParam *param);
+  void Compute(const FusionDWConvBNReluParam<DeviceType> &param) const;
+  bool Init(FusionDWConvBNReluParam<DeviceType> *param);
 };
 
 }  // namespace operators

src/operators/kernel/elementwise_add_kernel.h

@@ -27,10 +27,11 @@ using namespace framework;
 
 template <typename DeviceType, typename T>
 class ElementwiseAddKernel
-    : public framework::OpKernelBase<DeviceType, ElementwiseAddParam> {
+    : public framework::OpKernelBase<DeviceType,
+                                     ElementwiseAddParam<DeviceType>> {
  public:
-  void Compute(const ElementwiseAddParam &param) const;
-  bool Init(ElementwiseAddParam *param);
+  void Compute(const ElementwiseAddParam<DeviceType> &param) const;
+  bool Init(ElementwiseAddParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/elementwise_add_relu_kernel.h

@@ -26,10 +26,11 @@ using namespace framework;
 
 template <typename DeviceType, typename T>
 class ElementwiseAddReluKernel
-    : public framework::OpKernelBase<DeviceType, ElementwiseAddReluParam> {
+    : public framework::OpKernelBase<DeviceType,
+                                     ElementwiseAddReluParam<DeviceType>> {
  public:
-  void Compute(const ElementwiseAddReluParam &param) const;
-  bool Init(ElementwiseAddReluParam *param);
+  void Compute(const ElementwiseAddReluParam<DeviceType> &param) const;
+  bool Init(ElementwiseAddReluParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/fc_relu_kernel.h

@@ -25,10 +25,11 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class FusionFcReluKernel
-    : public framework::OpKernelBase<DeviceType, FusionFcReluParam> {
+    : public framework::OpKernelBase<DeviceType,
+                                     FusionFcReluParam<DeviceType>> {
  public:
-  void Compute(const FusionFcReluParam& param) const;
-  bool Init(FusionFcReluParam* param);
+  void Compute(const FusionFcReluParam<DeviceType>& param) const;
+  bool Init(FusionFcReluParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/fpga/concat_kernel.cpp

@@ -20,12 +20,12 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConcatKernel<FPGA, float>::Init(ConcatParam *param) {
+bool ConcatKernel<FPGA, float>::Init(ConcatParam<FPGA> *param) {
   return true;
 }
 
 template <>
-void ConcatKernel<FPGA, float>::Compute(const ConcatParam &param) const {
+void ConcatKernel<FPGA, float>::Compute(const ConcatParam<FPGA> &param) const {
   auto inputs = param.Inputs();
   auto *out = param.Out();
   int64_t axis = param.Axis();

src/operators/kernel/fpga/conv_add_bn_kernel.cpp

@@ -22,7 +22,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddBNKernel<FPGA, float>::Init(FusionConvAddBNParam *param) {
+bool ConvAddBNKernel<FPGA, float>::Init(FusionConvAddBNParam<FPGA> *param) {
   bool relu_enabled = false;
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
@@ -89,7 +89,7 @@ bool ConvAddBNKernel<FPGA, float>::Init(FusionConvAddBNParam *param) {
 
 template <>
 void ConvAddBNKernel<FPGA, float>::Compute(
-    const FusionConvAddBNParam &param) const {
+    const FusionConvAddBNParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 template class ConvAddBNKernel<FPGA, float>;

src/operators/kernel/fpga/conv_add_bn_relu_kernel.cpp

@@ -21,7 +21,8 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddBNReluKernel<FPGA, float>::Init(FusionConvAddBNReluParam *param) {
+bool ConvAddBNReluKernel<FPGA, float>::Init(
+    FusionConvAddBNReluParam<FPGA> *param) {
   bool relu_enabled = true;
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
@@ -84,7 +85,7 @@ bool ConvAddBNReluKernel<FPGA, float>::Init(FusionConvAddBNReluParam *param) {
 
 template <>
 void ConvAddBNReluKernel<FPGA, float>::Compute(
-    const FusionConvAddBNReluParam &param) const {
+    const FusionConvAddBNReluParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 template class ConvAddBNReluKernel<FPGA, float>;

src/operators/kernel/fpga/conv_add_relu_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvAddReluKernel<FPGA, float>::Init(FusionConvAddReluParam *param) {
+bool ConvAddReluKernel<FPGA, float>::Init(FusionConvAddReluParam<FPGA> *param) {
   bool relu_enabled = true;
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
@@ -69,7 +69,7 @@ bool ConvAddReluKernel<FPGA, float>::Init(FusionConvAddReluParam *param) {
 
 template <>
 void ConvAddReluKernel<FPGA, float>::Compute(
-    const FusionConvAddReluParam &param) const {
+    const FusionConvAddReluParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 template class ConvAddReluKernel<FPGA, float>;

src/operators/kernel/fpga/conv_bn_kernel.cpp

@@ -22,7 +22,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvBNKernel<FPGA, float>::Init(FusionConvBNParam *param) {
+bool ConvBNKernel<FPGA, float>::Init(FusionConvBNParam<FPGA> *param) {
   bool relu_enabled = false;
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
@@ -83,7 +83,8 @@ bool ConvBNKernel<FPGA, float>::Init(FusionConvBNParam *param) {
 }
 
 template <>
-void ConvBNKernel<FPGA, float>::Compute(const FusionConvBNParam &param) const {
+void ConvBNKernel<FPGA, float>::Compute(
+    const FusionConvBNParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 template class ConvBNKernel<FPGA, float>;

src/operators/kernel/fpga/conv_bn_relu_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ConvBNReluKernel<FPGA, float>::Init(FusionConvBNReluParam *param) {
+bool ConvBNReluKernel<FPGA, float>::Init(FusionConvBNReluParam<FPGA> *param) {
   bool relu_enabled = true;
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
@@ -80,7 +80,7 @@ bool ConvBNReluKernel<FPGA, float>::Init(FusionConvBNReluParam *param) {
 
 template <>
 void ConvBNReluKernel<FPGA, float>::Compute(
-    const FusionConvBNReluParam &param) const {
+    const FusionConvBNReluParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 template class ConvBNReluKernel<FPGA, float>;

src/operators/kernel/fpga/dropout_kernel.cpp

@@ -20,13 +20,14 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool DropoutKernel<FPGA, float>::Init(DropoutParam *param) {
+bool DropoutKernel<FPGA, float>::Init(DropoutParam<FPGA> *param) {
   param->Out()->ShareDataWith(*param->InputX());
   return true;
 }
 
 template <>
-void DropoutKernel<FPGA, float>::Compute(const DropoutParam &param) const {
+void DropoutKernel<FPGA, float>::Compute(
+    const DropoutParam<FPGA> &param) const {
   // auto *input_x = param.InputX();
   // auto *out = param.Out();
   // auto input_x_ptr = input_x->data<float>();

src/operators/kernel/fpga/elementwise_add_relu_kernel.cpp

@@ -20,7 +20,7 @@ namespace operators {
 
 template <>
 bool ElementwiseAddReluKernel<FPGA, float>::Init(
-    ElementwiseAddReluParam *param) {
+    ElementwiseAddReluParam<FPGA> *param) {
   bool relu_enabled = true;
   const Tensor *input_x = param->InputX();
   const Tensor *input_y = param->InputY();
@@ -57,7 +57,7 @@ bool ElementwiseAddReluKernel<FPGA, float>::Init(
 
 template <>
 void ElementwiseAddReluKernel<FPGA, float>::Compute(
-    const ElementwiseAddReluParam &param) const {
+    const ElementwiseAddReluParam<FPGA> &param) const {
   fpga::ComputeFpgaEWAdd(param.FpgaArgs());
 }
 }  // namespace operators

src/operators/kernel/fpga/fc_relu_kernel.cpp

@@ -21,7 +21,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool FusionFcReluKernel<FPGA, float>::Init(FusionFcReluParam *param) {
+bool FusionFcReluKernel<FPGA, float>::Init(FusionFcReluParam<FPGA> *param) {
   bool relu_enabled = true;
   const Tensor *input_x = param->InputX();
   auto input_x_ptr = input_x->data<half>();
@@ -70,7 +70,7 @@ bool FusionFcReluKernel<FPGA, float>::Init(FusionFcReluParam *param) {
 }
 template <>
 void FusionFcReluKernel<FPGA, float>::Compute(
-    const FusionFcReluParam &param) const {
+    const FusionFcReluParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 };
 

src/operators/kernel/fpga/fusion_fc_kernel.cpp

@@ -20,7 +20,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool FusionFcKernel<FPGA, float>::Init(FusionFcParam *param) {
+bool FusionFcKernel<FPGA, float>::Init(FusionFcParam<FPGA> *param) {
   bool relu_enabled = false;
   const Tensor *input_x = param->InputX();
   auto input_x_ptr = input_x->data<half>();
@@ -66,7 +66,8 @@ bool FusionFcKernel<FPGA, float>::Init(FusionFcParam *param) {
 }
 
 template <>
-void FusionFcKernel<FPGA, float>::Compute(const FusionFcParam &param) const {
+void FusionFcKernel<FPGA, float>::Compute(
+    const FusionFcParam<FPGA> &param) const {
   fpga::ComputeFpgaConv(param.FpgaArgs());
 }
 }  // namespace operators

src/operators/kernel/fpga/pool_kernel.cpp

@@ -20,7 +20,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool PoolKernel<FPGA, float>::Init(PoolParam *param) {
+bool PoolKernel<FPGA, float>::Init(PoolParam<FPGA> *param) {
   const Tensor *input = param->Input();
   auto input_ptr = input->data<half>();
   Tensor *output = param->Output();
@@ -46,7 +46,7 @@ bool PoolKernel<FPGA, float>::Init(PoolParam *param) {
 }
 
 template <>
-void PoolKernel<FPGA, float>::Compute(const PoolParam &param) const {
+void PoolKernel<FPGA, float>::Compute(const PoolParam<FPGA> &param) const {
 #ifdef PADDLE_MOBILE_FPGA
   fpga::ComputeFpgaPool(param.FpgaArgs());
 #endif

src/operators/kernel/fpga/softmax_kernel.cpp

@@ -23,7 +23,7 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool SoftmaxKernel<FPGA, float>::Init(SoftmaxParam *param) {
+bool SoftmaxKernel<FPGA, float>::Init(SoftmaxParam<FPGA> *param) {
   const Tensor *input = param->InputX();
   if (input->type() == typeid(half)) {
     auto input_ptr = input->data<half>();
@@ -43,7 +43,8 @@ bool SoftmaxKernel<FPGA, float>::Init(SoftmaxParam *param) {
 }
 
 template <>
-void SoftmaxKernel<FPGA, float>::Compute(const SoftmaxParam &param) const {
+void SoftmaxKernel<FPGA, float>::Compute(
+    const SoftmaxParam<FPGA> &param) const {
   // SoftmaxCompute<float>(param);
 }
 

src/operators/kernel/fusion_fc_kernel.h

@@ -25,10 +25,10 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class FusionFcKernel
-    : public framework::OpKernelBase<DeviceType, FusionFcParam> {
+    : public framework::OpKernelBase<DeviceType, FusionFcParam<DeviceType>> {
  public:
-  void Compute(const FusionFcParam& param) const;
-  bool Init(FusionFcParam* param);
+  void Compute(const FusionFcParam<DeviceType>& param) const;
+  bool Init(FusionFcParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/im2sequence_kernel.h

@@ -27,10 +27,10 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class Im2SequenceKernel
-    : public framework::OpKernelBase<DeviceType, Im2SequenceParam> {
+    : public framework::OpKernelBase<DeviceType, Im2SequenceParam<DeviceType>> {
  public:
-  void Compute(const Im2SequenceParam& param) const;
-  bool Init(Im2SequenceParam* para);
+  void Compute(const Im2SequenceParam<DeviceType>& param) const;
+  bool Init(Im2SequenceParam<DeviceType>* para);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/lrn_kernel.h

@@ -170,10 +170,11 @@ struct LRNFunctor {
 };
 
 template <typename DeviceType, typename T>
-class LrnKernel : public framework::OpKernelBase<DeviceType, LrnParam> {
+class LrnKernel
+    : public framework::OpKernelBase<DeviceType, LrnParam<DeviceType>> {
  public:
-  void Compute(const LrnParam &param) const;
-  bool Init(LrnParam *param);
+  void Compute(const LrnParam<DeviceType> &param) const;
+  bool Init(LrnParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/mali/batchnorm_kernel.cpp

@@ -36,7 +36,7 @@ class AclBatchNormOp : public acl::ACLOperator {
   AclBatchNormOp& operator=(AclBatchNormOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const BatchNormParam& param) {
+  void InitAclLayer(const BatchNormParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape input_shape(args.in_cols, args.in_rows,
                                          args.in_depth, args.batch);
@@ -68,7 +68,7 @@ class AclBatchNormOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const BatchNormParam& param) {
+  bool Bypass_acl(const BatchNormParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -81,7 +81,7 @@ class AclBatchNormOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const BatchNormParam& param) {
+  void AclParametersByContext(const BatchNormParam<DeviceType>& param) {
     const Tensor* in_x = param.InputX();
     Tensor* out = param.OutputY();
     const Tensor* scale = param.InputScale();
@@ -129,7 +129,7 @@ class AclBatchNormOp : public acl::ACLOperator {
 };
 
 template <>
-bool BatchNormKernel<GPU_MALI, float>::Init(BatchNormParam* param) {
+bool BatchNormKernel<GPU_MALI, float>::Init(BatchNormParam<GPU_MALI>* param) {
   AclBatchNormOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclBatchNormOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -145,7 +145,7 @@ bool BatchNormKernel<GPU_MALI, float>::Init(BatchNormParam* param) {
 
 template <>
 void BatchNormKernel<GPU_MALI, float>::Compute(
-    const BatchNormParam& param) const {
+    const BatchNormParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclBatchNormOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclBatchNormOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/concat_kernel.cpp

@@ -38,7 +38,7 @@ class AclConcatOp : public acl::ACLOperator {
 
   acl::AclParameters& getargs() { return args; }
 
-  void InitAclLayer(const ConcatParam& param) {
+  void InitAclLayer(const ConcatParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     const std::vector<framework::LoDTensor*>* input_data = &args.in_tensor;
     arm_compute::TensorShape output_shape(args.out_cols, args.out_rows,
@@ -70,7 +70,7 @@ class AclConcatOp : public acl::ACLOperator {
     T type;
     acl::acl_run(this, input, output, type);
   }
-  bool Bypass_acl(const ConcatParam& param) {
+  bool Bypass_acl(const ConcatParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -82,7 +82,7 @@ class AclConcatOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const ConcatParam& param) {
+  void AclParametersByContext(const ConcatParam<DeviceType>& param) {
     auto inputs = param.Inputs();
     auto* output = param.Out();
     int64_t axis = param.Axis();
@@ -102,7 +102,7 @@ class AclConcatOp : public acl::ACLOperator {
 };
 
 template <>
-bool ConcatKernel<GPU_MALI, float>::Init(ConcatParam* param) {
+bool ConcatKernel<GPU_MALI, float>::Init(ConcatParam<GPU_MALI>* param) {
   AclConcatOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConcatOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -117,7 +117,8 @@ bool ConcatKernel<GPU_MALI, float>::Init(ConcatParam* param) {
 }
 
 template <>
-void ConcatKernel<GPU_MALI, float>::Compute(const ConcatParam& param) const {
+void ConcatKernel<GPU_MALI, float>::Compute(
+    const ConcatParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclConcatOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConcatOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/conv_add_kernel.cpp

@@ -37,7 +37,7 @@ class AclConvAddOp : public acl::ACLOperator {
   AclConvAddOp& operator=(AclConvAddOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const FusionConvAddParam& param) {
+  void InitAclLayer(const FusionConvAddParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape input_shape(args.in_cols, args.in_rows,
                                          args.in_depth, args.batch);
@@ -75,7 +75,7 @@ class AclConvAddOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const FusionConvAddParam& param) {
+  bool Bypass_acl(const FusionConvAddParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -119,7 +119,7 @@ class AclConvAddOp : public acl::ACLOperator {
     }
   }
 
-  void AclParametersByContext(const FusionConvAddParam& param) {
+  void AclParametersByContext(const FusionConvAddParam<DeviceType>& param) {
     const Tensor* input = param.Input();
     Tensor filter = *param.Filter();
     Tensor* output = param.Output();
@@ -196,7 +196,7 @@ class AclConvAddOp : public acl::ACLOperator {
 };
 
 template <>
-bool ConvAddKernel<GPU_MALI, float>::Init(FusionConvAddParam* param) {
+bool ConvAddKernel<GPU_MALI, float>::Init(FusionConvAddParam<GPU_MALI>* param) {
   AclConvAddOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConvAddOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -212,7 +212,7 @@ bool ConvAddKernel<GPU_MALI, float>::Init(FusionConvAddParam* param) {
 
 template <>
 void ConvAddKernel<GPU_MALI, float>::Compute(
-    const FusionConvAddParam& param) const {
+    const FusionConvAddParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclConvAddOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConvAddOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/conv_kernel.cpp

@@ -37,7 +37,7 @@ class AclConvOp : public acl::ACLOperator {
   AclConvOp& operator=(AclConvOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const ConvParam& param) {
+  void InitAclLayer(const ConvParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape input_shape(args.in_cols, args.in_rows,
                                          args.in_depth, args.batch);
@@ -76,7 +76,7 @@ class AclConvOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const ConvParam& param) {
+  bool Bypass_acl(const ConvParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -120,7 +120,7 @@ class AclConvOp : public acl::ACLOperator {
     }
   }
 
-  void AclParametersByContext(const ConvParam& param) {
+  void AclParametersByContext(const ConvParam<DeviceType>& param) {
     const Tensor* input = param.Input();
     Tensor filter = *param.Filter();
     Tensor* output = param.Output();
@@ -196,7 +196,7 @@ class AclConvOp : public acl::ACLOperator {
 };
 
 template <>
-bool ConvKernel<GPU_MALI, float>::Init(ConvParam* param) {
+bool ConvKernel<GPU_MALI, float>::Init(ConvParam<GPU_MALI>* param) {
   AclConvOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConvOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -211,7 +211,8 @@ bool ConvKernel<GPU_MALI, float>::Init(ConvParam* param) {
 }
 
 template <>
-void ConvKernel<GPU_MALI, float>::Compute(const ConvParam& param) const {
+void ConvKernel<GPU_MALI, float>::Compute(
+    const ConvParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclConvOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclConvOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/elementwise_add_kernel.cpp

@@ -27,13 +27,14 @@ struct AddFunctor {
 };
 
 template <>
-bool ElementwiseAddKernel<GPU_MALI, float>::Init(ElementwiseAddParam *param) {
+bool ElementwiseAddKernel<GPU_MALI, float>::Init(
+    ElementwiseAddParam<GPU_MALI> *param) {
   return true;
 }
 
 template <>
 void ElementwiseAddKernel<GPU_MALI, float>::Compute(
-    const ElementwiseAddParam &param) const {
+    const ElementwiseAddParam<GPU_MALI> &param) const {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   Tensor *Out = param.Out();

src/operators/kernel/mali/fushion_fc_kernel.cpp

@@ -20,13 +20,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool FusionFcKernel<GPU_MALI, float>::Init(FusionFcParam *param) {
+bool FusionFcKernel<GPU_MALI, float>::Init(FusionFcParam<GPU_MALI> *param) {
   return true;
 }
 
 template <>
 void FusionFcKernel<GPU_MALI, float>::Compute(
-    const FusionFcParam &param) const {
+    const FusionFcParam<GPU_MALI> &param) const {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   const Tensor *input_z = param.InputZ();

src/operators/kernel/mali/lrn_kernel.cpp

@@ -40,7 +40,7 @@ class AclLrnOp : public acl::ACLOperator {
   AclLrnOp& operator=(AclLrnOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const LrnParam& param) {
+  void InitAclLayer(const LrnParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape shape(args.in_cols, args.in_rows, args.in_depth);
 
@@ -65,7 +65,7 @@ class AclLrnOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const LrnParam& param) {
+  bool Bypass_acl(const LrnParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -78,7 +78,7 @@ class AclLrnOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const LrnParam& param) {
+  void AclParametersByContext(const LrnParam<DeviceType>& param) {
     const Tensor* in_x = param.InputX();
     Tensor* out = param.Out();
 
@@ -111,7 +111,7 @@ class AclLrnOp : public acl::ACLOperator {
 };
 
 template <>
-bool LrnKernel<GPU_MALI, float>::Init(LrnParam* param) {
+bool LrnKernel<GPU_MALI, float>::Init(LrnParam<GPU_MALI>* param) {
   AclLrnOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclLrnOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -127,7 +127,8 @@ bool LrnKernel<GPU_MALI, float>::Init(LrnParam* param) {
 }
 
 template <>
-void LrnKernel<GPU_MALI, float>::Compute(const LrnParam& param) const {
+void LrnKernel<GPU_MALI, float>::Compute(
+    const LrnParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclLrnOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclLrnOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/mul_kernel.cpp

@@ -22,12 +22,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool MulKernel<GPU_MALI, float>::Init(MulParam *param) {
+bool MulKernel<GPU_MALI, float>::Init(MulParam<GPU_MALI> *param) {
   return true;
 }
 
 template <>
-void MulKernel<GPU_MALI, float>::Compute(const MulParam &param) const {
+void MulKernel<GPU_MALI, float>::Compute(
+    const MulParam<GPU_MALI> &param) const {
   const Tensor *input_x = param.InputX();
   const Tensor *input_y = param.InputY();
   Tensor *out = param.Out();

src/operators/kernel/mali/pool_kernel.cpp

@@ -39,7 +39,7 @@ class AclPoolOp : public acl::ACLOperator {
   AclPoolOp& operator=(AclPoolOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const PoolParam& param) {
+  void InitAclLayer(const PoolParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape input_shape(args.in_cols, args.in_rows,
                                          args.in_depth);
@@ -79,7 +79,7 @@ class AclPoolOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const PoolParam& param) {
+  bool Bypass_acl(const PoolParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -100,7 +100,7 @@ class AclPoolOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const PoolParam& param) {
+  void AclParametersByContext(const PoolParam<DeviceType>& param) {
     const Tensor* in_x = param.Input();
     Tensor* out = param.Output();
     std::string pooling_type = param.PoolingType();
@@ -180,7 +180,7 @@ class AclPoolOp : public acl::ACLOperator {
 };
 
 template <>
-bool PoolKernel<GPU_MALI, float>::Init(PoolParam* param) {
+bool PoolKernel<GPU_MALI, float>::Init(PoolParam<GPU_MALI>* param) {
   AclPoolOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclPoolOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -195,7 +195,8 @@ bool PoolKernel<GPU_MALI, float>::Init(PoolParam* param) {
 }
 
 template <>
-void PoolKernel<GPU_MALI, float>::Compute(const PoolParam& param) const {
+void PoolKernel<GPU_MALI, float>::Compute(
+    const PoolParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclPoolOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclPoolOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/relu_kernel.cpp

@@ -39,7 +39,7 @@ class AclReluOp : public acl::ACLOperator {
   AclReluOp& operator=(AclReluOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const ReluParam& param) {
+  void InitAclLayer(const ReluParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape input_shape(args.in_cols, args.in_rows,
                                          args.in_depth, args.batch);
@@ -68,7 +68,7 @@ class AclReluOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const ReluParam& param) {
+  bool Bypass_acl(const ReluParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -80,7 +80,7 @@ class AclReluOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const ReluParam& param) {
+  void AclParametersByContext(const ReluParam<DeviceType>& param) {
     const auto* input_x = param.InputX();
     auto* out = param.Out();
 
@@ -100,7 +100,7 @@ class AclReluOp : public acl::ACLOperator {
 };
 
 template <>
-bool ReluKernel<GPU_MALI, float>::Init(ReluParam* param) {
+bool ReluKernel<GPU_MALI, float>::Init(ReluParam<GPU_MALI>* param) {
   AclReluOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclReluOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -115,7 +115,8 @@ bool ReluKernel<GPU_MALI, float>::Init(ReluParam* param) {
 }
 
 template <>
-void ReluKernel<GPU_MALI, float>::Compute(const ReluParam& param) const {
+void ReluKernel<GPU_MALI, float>::Compute(
+    const ReluParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclReluOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclReluOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mali/reshape_kernel.cpp

@@ -22,12 +22,13 @@ namespace paddle_mobile {
 namespace operators {
 
 template <>
-bool ReshapeKernel<GPU_MALI, float>::Init(ReshapeParam *param) {
+bool ReshapeKernel<GPU_MALI, float>::Init(ReshapeParam<GPU_MALI> *param) {
   return true;
 }
 
 template <>
-void ReshapeKernel<GPU_MALI, float>::Compute(const ReshapeParam &param) const {
+void ReshapeKernel<GPU_MALI, float>::Compute(
+    const ReshapeParam<GPU_MALI> &param) const {
   const auto *input_x = param.InputX();
   const auto &input_x_dims = input_x->dims();
   auto *out = param.Out();

src/operators/kernel/mali/softmax_kernel.cpp

@@ -39,7 +39,7 @@ class AclSoftmaxOp : public acl::ACLOperator {
   AclSoftmaxOp& operator=(AclSoftmaxOp&&) = delete;
 
   acl::AclParameters& getargs() { return args; }
-  void InitAclLayer(const SoftmaxParam& param) {
+  void InitAclLayer(const SoftmaxParam<DeviceType>& param) {
     setTargetHint(acl::TargetHint::OPENCL);
     arm_compute::TensorShape shape(args.in_depth, args.batch);
 
@@ -58,7 +58,7 @@ class AclSoftmaxOp : public acl::ACLOperator {
   void RunAcl(void* input, void* output) {
     acl::ACLOperator::acl_run(input, output);
   }
-  bool Bypass_acl(const SoftmaxParam& param) {
+  bool Bypass_acl(const SoftmaxParam<DeviceType>& param) {
     bool bypass_acl = false;
     AclParametersByContext(param);
     InitAclLayer(param);
@@ -71,7 +71,7 @@ class AclSoftmaxOp : public acl::ACLOperator {
   }
 
  private:
-  void AclParametersByContext(const SoftmaxParam& param) {
+  void AclParametersByContext(const SoftmaxParam<DeviceType>& param) {
     const framework::Tensor* in_x = param.InputX();
     framework::Tensor* out = param.Out();
     auto x_dims = in_x->dims();
@@ -97,7 +97,7 @@ class AclSoftmaxOp : public acl::ACLOperator {
 };
 
 template <>
-bool SoftmaxKernel<GPU_MALI, float>::Init(SoftmaxParam* param) {
+bool SoftmaxKernel<GPU_MALI, float>::Init(SoftmaxParam<GPU_MALI>* param) {
   AclSoftmaxOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclSoftmaxOp<GPU_MALI, float>*>(this->GetAclOp());
   if (acl_op == nullptr) {
@@ -112,7 +112,8 @@ bool SoftmaxKernel<GPU_MALI, float>::Init(SoftmaxParam* param) {
 }
 
 template <>
-void SoftmaxKernel<GPU_MALI, float>::Compute(const SoftmaxParam& param) const {
+void SoftmaxKernel<GPU_MALI, float>::Compute(
+    const SoftmaxParam<GPU_MALI>& param) const {
   std::cout << "init acl" << std::endl;
   AclSoftmaxOp<GPU_MALI, float>* acl_op =
       reinterpret_cast<AclSoftmaxOp<GPU_MALI, float>*>(this->GetAclOp());

src/operators/kernel/mul_kernel.h

@@ -26,10 +26,11 @@ namespace operators {
 using namespace framework;
 
 template <typename DeviceType, typename T>
-class MulKernel : public framework::OpKernelBase<DeviceType, MulParam> {
+class MulKernel
+    : public framework::OpKernelBase<DeviceType, MulParam<DeviceType>> {
  public:
-  void Compute(const MulParam &param) const;
-  bool Init(MulParam *param);
+  void Compute(const MulParam<DeviceType> &param) const;
+  bool Init(MulParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/multiclass_nms_kernel.h

@@ -25,10 +25,11 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class MultiClassNMSKernel
-    : public framework::OpKernelBase<DeviceType, MultiClassNMSParam> {
+    : public framework::OpKernelBase<DeviceType,
+                                     MultiClassNMSParam<DeviceType>> {
  public:
-  void Compute(const MultiClassNMSParam& param) const;
-  bool Init(MultiClassNMSParam* param);
+  void Compute(const MultiClassNMSParam<DeviceType>& param) const;
+  bool Init(MultiClassNMSParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/pool_kernel.h

@@ -24,10 +24,10 @@ namespace operators {
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class PoolKernel : public OpKernelBase<DeviceType, PoolParam> {
+class PoolKernel : public OpKernelBase<DeviceType, PoolParam<DeviceType>> {
  public:
-  void Compute(const PoolParam &param) const override;
-  bool Init(PoolParam *param);
+  void Compute(const PoolParam<DeviceType> &param) const override;
+  bool Init(PoolParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/prelu_kernel.h

@@ -21,9 +21,10 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename DeviceType, typename T>
-class PReluKernel : public framework::OpKernelBase<DeviceType, PReluParam> {
+class PReluKernel
+    : public framework::OpKernelBase<DeviceType, PReluParam<DeviceType>> {
  public:
-  void Compute(const PReluParam& param) const;
+  void Compute(const PReluParam<DeviceType>& param) const;
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/prior_box_kernel.h

@@ -52,10 +52,10 @@ inline void ExpandAspectRatios(const std::vector<float>& input_aspect_ratior,
 
 template <typename DeviceType, typename T>
 class PriorBoxKernel
-    : public framework::OpKernelBase<DeviceType, PriorBoxParam> {
+    : public framework::OpKernelBase<DeviceType, PriorBoxParam<DeviceType>> {
  public:
-  void Compute(const PriorBoxParam& param) const;
-  bool Init(PriorBoxParam* param);
+  void Compute(const PriorBoxParam<DeviceType>& param) const;
+  bool Init(PriorBoxParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/relu_kernel.h

@@ -24,10 +24,11 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename DeviceType, typename T>
-class ReluKernel : public framework::OpKernelBase<DeviceType, ReluParam> {
+class ReluKernel
+    : public framework::OpKernelBase<DeviceType, ReluParam<DeviceType>> {
  public:
-  void Compute(const ReluParam& param) const;
-  bool Init(ReluParam* param);
+  void Compute(const ReluParam<DeviceType>& param) const;
+  bool Init(ReluParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/reshape_kernel.h

@@ -68,10 +68,11 @@ inline framework::DDim ValidateShape(const std::vector<int> shape,
 }
 
 template <typename DeviceType, typename T>
-class ReshapeKernel : public framework::OpKernelBase<DeviceType, ReshapeParam> {
+class ReshapeKernel
+    : public framework::OpKernelBase<DeviceType, ReshapeParam<DeviceType>> {
  public:
-  void Compute(const ReshapeParam& param) const;
-  bool Init(ReshapeParam* param);
+  void Compute(const ReshapeParam<DeviceType>& param) const;
+  bool Init(ReshapeParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/resize_kernel.h

@@ -23,8 +23,8 @@ limitations under the License. */
 
 namespace paddle_mobile {
 namespace operators {
-
-inline framework::DDim CalOutputShape(const ResizeParam &param) {
+template <typename DeviceType>
+inline framework::DDim CalOutputShape(const ResizeParam<DeviceType> &param) {
   const auto *input_x = param.InputX();
   const auto &input_x_dims = input_x->dims();
   auto *out = param.Out();
@@ -32,7 +32,8 @@ inline framework::DDim CalOutputShape(const ResizeParam &param) {
   const auto *input_shape = param.InputShape();
 
   if (input_shape) {
-    auto *shape_data = input_shape->data<int>();
+    input_x->dims()[0];
+    auto *shape_data = input_shape->template data<int>();
     framework::Tensor cpu_shape_tensor;
     auto shape =
         std::vector<int>(shape_data, shape_data + input_shape->numel());
@@ -69,9 +70,10 @@ inline framework::DDim CalOutputShape(const ResizeParam &param) {
 }
 
 template <typename DeviceType, typename T>
-class ResizeKernel : public framework::OpKernelBase<DeviceType, ResizeParam> {
+class ResizeKernel
+    : public framework::OpKernelBase<DeviceType, ResizeParam<DeviceType>> {
  public:
-  void Compute(const ResizeParam &param) const;
+  void Compute(const ResizeParam<DeviceType> &param) const;
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/scale_kernel.h

@@ -21,9 +21,10 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename DeviceType, typename T>
-class ScaleKernel : public framework::OpKernelBase<DeviceType, ScaleParam> {
+class ScaleKernel
+    : public framework::OpKernelBase<DeviceType, ScaleParam<DeviceType>> {
  public:
-  void Compute(const ScaleParam& param) const;
+  void Compute(const ScaleParam<DeviceType>& param) const;
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/sigmoid_kernel.h

@@ -23,10 +23,11 @@ namespace operators {
 using framework::OpKernelBase;
 void sigmoid(const Tensor* X, Tensor* Y);
 template <typename DeviceType, typename T>
-class SigmoidKernel : public OpKernelBase<DeviceType, SigmoidParam> {
+class SigmoidKernel
+    : public OpKernelBase<DeviceType, SigmoidParam<DeviceType>> {
  public:
-  void Compute(const SigmoidParam& param) const override;
-  bool Init(SigmoidParam* param);
+  void Compute(const SigmoidParam<DeviceType>& param) const override;
+  bool Init(SigmoidParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/slice_kernel.h

@@ -21,9 +21,10 @@ namespace paddle_mobile {
 namespace operators {
 
 template <typename DeviceType, typename T>
-class SliceKernel : public framework::OpKernelBase<DeviceType, SliceParam> {
+class SliceKernel
+    : public framework::OpKernelBase<DeviceType, SliceParam<DeviceType>> {
  public:
-  void Compute(const SliceParam& param) const {}
+  void Compute(const SliceParam<DeviceType>& param) const {}
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/softmax_kernel.h

@@ -24,10 +24,11 @@ namespace operators {
 using framework::OpKernelBase;
 
 template <typename DeviceType, typename T>
-class SoftmaxKernel : public OpKernelBase<DeviceType, SoftmaxParam> {
+class SoftmaxKernel
+    : public OpKernelBase<DeviceType, SoftmaxParam<DeviceType>> {
  public:
-  void Compute(const SoftmaxParam &param) const override;
-  bool Init(SoftmaxParam *param);
+  void Compute(const SoftmaxParam<DeviceType> &param) const override;
+  bool Init(SoftmaxParam<DeviceType> *param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/kernel/transpose_kernel.h

@@ -26,10 +26,10 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class TransposeKernel
-    : public framework::OpKernelBase<DeviceType, TransposeParam> {
+    : public framework::OpKernelBase<DeviceType, TransposeParam<DeviceType>> {
  public:
-  void Compute(const TransposeParam& param) const;
-  bool Init(TransposeParam* param);
+  void Compute(const TransposeParam<DeviceType>& param) const;
+  bool Init(TransposeParam<DeviceType>* param);
 };
 }  // namespace operators
 }  // namespace paddle_mobile

src/operators/lrn_op.h

@@ -26,17 +26,18 @@ namespace operators {
 using std::string;
 template <typename DeviceType, typename T>
 class LrnOp : public framework::OperatorWithKernel<
-                  DeviceType, LrnParam, operators::LrnKernel<DeviceType, T>> {
+                  DeviceType, LrnParam<DeviceType>,
+                  operators::LrnKernel<DeviceType, T>> {
  public:
   LrnOp(const string &type, const VariableNameMap &inputs,
         const VariableNameMap &outputs, const framework::AttributeMap &attrs,
         std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, LrnParam,
+      : framework::OperatorWithKernel<DeviceType, LrnParam<DeviceType>,
                                       operators::LrnKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, LrnParam,
+      DeviceType, LrnParam<DeviceType>,
       operators::LrnKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/math/pooling.cpp

@@ -58,7 +58,6 @@ class PoolFunctor<CPU, PoolProcess, T> {
 
     const T *input_data = input.data<T>();
     T *output_data = output->mutable_data<T>();
-
     for (int i = 0; i < batch_size; i++) {
       for (int c = 0; c < output_channels; ++c) {
 #pragma omp parallel for

src/operators/mul_op.h

@@ -26,17 +26,18 @@ namespace operators {
 
 template <typename DeviceType, typename T>
 class MulOp : public framework::OperatorWithKernel<
-                  DeviceType, MulParam, operators::MulKernel<DeviceType, T>> {
+                  DeviceType, MulParam<DeviceType>,
+                  operators::MulKernel<DeviceType, T>> {
  public:
   MulOp(const std::string &type, const VariableNameMap &inputs,
         const VariableNameMap &outputs, const framework::AttributeMap &attrs,
         std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, MulParam,
+      : framework::OperatorWithKernel<DeviceType, MulParam<DeviceType>,
                                       operators::MulKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, MulParam,
+      DeviceType, MulParam<DeviceType>,
       operators::MulKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/multiclass_nms_op.h

@@ -29,7 +29,7 @@ using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
 class MultiClassNMSOp : public framework::OperatorWithKernel<
-                            DeviceType, MultiClassNMSParam,
+                            DeviceType, MultiClassNMSParam<DeviceType>,
                             operators::MultiClassNMSKernel<DeviceType, T>> {
  public:
   MultiClassNMSOp(const std::string &type, const VariableNameMap &inputs,
@@ -37,12 +37,12 @@ class MultiClassNMSOp : public framework::OperatorWithKernel<
                   const framework::AttributeMap &attrs,
                   std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, MultiClassNMSParam,
+            DeviceType, MultiClassNMSParam<DeviceType>,
             operators::MultiClassNMSKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, MultiClassNMSParam,
+      DeviceType, MultiClassNMSParam<DeviceType>,
       operators::MultiClassNMSKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/op_param.cpp

@@ -17,7 +17,8 @@ namespace paddle_mobile {
 namespace operators {
 
 #ifdef CONV_OP
-Print &operator<<(Print &printer, const ConvParam &conv_param) {
+template <>
+Print &operator<<(Print &printer, const ConvParam<CPU> &conv_param) {
   printer << "parameter of conv: "
           << "\n";
   printer << "  stride: "
@@ -37,11 +38,37 @@ Print &operator<<(Print &printer, const ConvParam &conv_param) {
   printer << "  output dims: " << conv_param.Output()->dims();
   return printer;
 }
+
+template class ConvParam<CPU>;
+template class ConvParam<FPGA>;
+template class ConvParam<GPU_MALI>;
+#endif
+
+template class ElementwiseAddParam<CPU>;
+template class ElementwiseAddParam<FPGA>;
+template class ElementwiseAddParam<GPU_MALI>;
+
+#ifdef MUL_OP
+template class MulParam<CPU>;
+template class MulParam<FPGA>;
+template class MulParam<GPU_MALI>;
+#endif
+
+#ifdef CONCAT_OP
+template class ConcatParam<CPU>;
+template class ConcatParam<FPGA>;
+template class ConcatParam<GPU_MALI>;
+#endif
+
+#ifdef LRN_OP
+template class LrnParam<CPU>;
+template class LrnParam<FPGA>;
+template class LrnParam<GPU_MALI>;
 #endif
 
 #ifdef FUSION_CONVADD_OP
 
-Print &operator<<(Print &printer, const FusionConvAddParam &conv_param) {
+Print &operator<<(Print &printer, const FusionConvAddParam<CPU> &conv_param) {
   printer << "parameter of conv_add: "
           << "\n";
   printer << "  stride: "

src/operators/op_param.h

@@ -18,6 +18,7 @@ limitations under the License. */
 #include <vector>
 #include "common/log.h"
 #include "common/type_define.h"
+#include "common/types.h"
 #include "framework/lod_tensor.h"
 #include "framework/scope.h"
 #include "framework/tensor.h"
@@ -37,6 +38,39 @@ using framework::Tensor;
 using std::string;
 using std::vector;
 
+template <typename Dtype>
+struct DtypeTensorTrait {
+  typedef void ptype;
+  typedef void rtype;
+};
+
+template <>
+struct DtypeTensorTrait<CPU> {
+  // This is the type we obtained in variable.
+  typedef framework::LoDTensor gtype;
+  // This type will be the parent class type
+  // or the same type.
+  typedef framework::Tensor rtype;
+};
+
+template <>
+struct DtypeTensorTrait<FPGA> {
+  // This is the type we obtained in variable.
+  typedef framework::LoDTensor gtype;
+  // This type will be the parent class type
+  // or the same type.
+  typedef framework::Tensor rtype;
+};
+
+template <>
+struct DtypeTensorTrait<GPU_MALI> {
+  // This is the type we obtained in variable.
+  typedef framework::LoDTensor gtype;
+  // This type will be the parent class type
+  // or the same type.
+  typedef framework::Tensor rtype;
+};
+
 class OpParam {
  protected:
   template <typename T>
@@ -200,24 +234,28 @@ class OpParam {
 };
 
 #ifdef CONV_OP
+template <typename Dtype>
 class ConvParam : OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ConvParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
             const AttributeMap &attrs, const Scope &scope) {
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutputFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutputFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -228,41 +266,45 @@ class ConvParam : OpParam {
   const int &Groups() const { return groups; }
 
  private:
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
 };
-
-Print &operator<<(Print &printer, const ConvParam &conv_param);
+template <typename Dtype>
+Print &operator<<(Print &printer, const ConvParam<Dtype> &conv_param);
 #endif
 
+template <typename Dtype>
 class ElementwiseAddParam : OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ElementwiseAddParam(const VariableNameMap &inputs,
                       const VariableNameMap &outputs, const AttributeMap &attrs,
                       const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_y_ = InputYFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_y_ = InputYFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputY() const { return input_y_; }
+  const RType *InputY() const { return input_y_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &Axis() const { return axis_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_y_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_y_;
+  RType *out_;
   int axis_;
 #ifdef PADDLE_MOBILE_FPGA
 
@@ -276,71 +318,84 @@ class ElementwiseAddParam : OpParam {
 };
 
 #ifdef FUSION_ELEMENTWISEADDRELU_OP
-using ElementwiseAddReluParam = ElementwiseAddParam;
+template <typename Dtype>
+using ElementwiseAddReluParam = ElementwiseAddParam<Dtype>;
 #endif
 
 #ifdef MUL_OP
+template <typename Dtype>
 class MulParam : OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   MulParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
            const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_y_ = InputYFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_y_ = InputYFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     x_num_col_dims_ = GetAttr<int>("x_num_col_dims", attrs);
     y_num_col_dims_ = GetAttr<int>("y_num_col_dims", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputY() const { return input_y_; }
+  const RType *InputY() const { return input_y_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &XNumColDims() const { return x_num_col_dims_; }
 
   const int &YNumColDims() const { return y_num_col_dims_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_y_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_y_;
+  RType *out_;
   int x_num_col_dims_;
   int y_num_col_dims_;
 };
 #endif
 
 #ifdef CONCAT_OP
+template <typename Dtype>
 class ConcatParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ConcatParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
               const AttributeMap &attrs, const Scope &scope) {
-    inputs_ = InputMultiFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    inputs_ = InputMultiFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
   }
 
-  vector<LoDTensor *> Inputs() const { return inputs_; }
+  vector<GType *> Inputs() const { return inputs_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &Axis() const { return axis_; }
 
  private:
-  vector<LoDTensor *> inputs_;
-  Tensor *out_;
+  vector<GType *> inputs_;
+  RType *out_;
   int axis_;
 };
 #endif
 
 #ifdef LRN_OP
+template <typename Dtype>
 class LrnParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   LrnParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
            const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
-    mid_out_ = MidOutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
+    mid_out_ = MidOutFrom<GType>(outputs, scope);
     n_ = GetAttr<int>("n", attrs);
     alpha_ = GetAttr<float>("alpha", attrs);
     beta_ = GetAttr<float>("beta", attrs);
@@ -348,11 +403,11 @@ class LrnParam : public OpParam {
     data_format_ = GetAttr<string>("data_format", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
-  Tensor *MidOut() const { return mid_out_; }
+  RType *MidOut() const { return mid_out_; }
 
   const int &N() const { return n_; }
 
@@ -365,9 +420,9 @@ class LrnParam : public OpParam {
   const string &DataFormat() const { return data_format_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
-  Tensor *mid_out_;
+  RType *input_x_;
+  RType *out_;
+  RType *mid_out_;
   int n_;
   float alpha_;
   float beta_;
@@ -377,32 +432,36 @@ class LrnParam : public OpParam {
 #endif
 
 #ifdef BATCHNORM_OP
+template <typename Dtype>
 class BatchNormParam : OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   BatchNormParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                  const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    output_y_ = OutputYFrom<LoDTensor>(outputs, scope);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    output_y_ = OutputYFrom<GType>(outputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  Tensor *OutputY() const { return output_y_; }
+  RType *OutputY() const { return output_y_; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -413,12 +472,12 @@ class BatchNormParam : OpParam {
   const string &DataFormat() const { return data_format_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *output_y_;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_x_;
+  RType *output_y_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
@@ -427,13 +486,17 @@ class BatchNormParam : OpParam {
 #endif
 
 #ifdef POOL_OP
+template <typename Dtype>
 class PoolParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   PoolParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
             const AttributeMap &attrs, const Scope &scope) {
-    input_ = InputXFrom<LoDTensor>(inputs, scope);
+    input_ = InputXFrom<GType>(inputs, scope);
 
-    output_ = OutFrom<LoDTensor>(outputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
     pooling_type_ = GetAttr<string>("pooling_type", attrs);
     ksize_ = GetAttr<vector<int>>("ksize", attrs);
     strides_ = GetAttr<vector<int>>("strides", attrs);
@@ -442,9 +505,9 @@ class PoolParam : public OpParam {
     global_pooling_ = GetAttr<bool>("global_pooling", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const string &PoolingType() const { return pooling_type_; }
 
@@ -459,8 +522,8 @@ class PoolParam : public OpParam {
   bool isGlobalPooling() const { return global_pooling_; }
 
  private:
-  Tensor *input_;
-  Tensor *output_;
+  RType *input_;
+  RType *output_;
   string pooling_type_;
   vector<int> ksize_;
   vector<int> strides_;
@@ -480,14 +543,18 @@ class PoolParam : public OpParam {
 #endif
 
 #ifdef PRIORBOX_OP
+template <typename Dtype>
 class PriorBoxParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   PriorBoxParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                 const AttributeMap &attrs, const Scope &scope) {
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    input_image_ = InputImageFrom<LoDTensor>(inputs, scope);
-    output_boxes_ = OutputBoxesFrom<LoDTensor>(outputs, scope);
-    output_variances_ = OutputVariancesFrom<LoDTensor>(outputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    input_image_ = InputImageFrom<GType>(inputs, scope);
+    output_boxes_ = OutputBoxesFrom<GType>(outputs, scope);
+    output_variances_ = OutputVariancesFrom<GType>(outputs, scope);
     min_sizes_ = GetAttr<vector<float>>("min_sizes", attrs);
     max_sizes_ = GetAttr<vector<float>>("max_sizes", attrs);
     aspect_ratios_ = GetAttr<vector<float>>("aspect_ratios", attrs);
@@ -498,13 +565,13 @@ class PriorBoxParam : public OpParam {
     step_h_ = GetAttr<float>("step_h", attrs);
     offset_ = GetAttr<float>("offset", attrs);
   }
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
-  const Tensor *InputImage() const { return input_image_; }
+  const RType *InputImage() const { return input_image_; }
 
-  Tensor *OutputBoxes() const { return output_boxes_; }
+  RType *OutputBoxes() const { return output_boxes_; }
 
-  Tensor *OutputVariances() const { return output_variances_; }
+  RType *OutputVariances() const { return output_variances_; }
 
   const vector<float> &MinSizes() const { return min_sizes_; }
 
@@ -525,10 +592,10 @@ class PriorBoxParam : public OpParam {
   const float &Offset() const { return offset_; }
 
  private:
-  Tensor *input_;
-  Tensor *input_image_;
-  Tensor *output_boxes_;
-  Tensor *output_variances_;
+  RType *input_;
+  RType *input_image_;
+  RType *output_boxes_;
+  RType *output_variances_;
   vector<float> min_sizes_;
   vector<float> max_sizes_;
   vector<float> aspect_ratios_;
@@ -542,58 +609,66 @@ class PriorBoxParam : public OpParam {
 #endif
 
 #ifdef BOXCODER_OP
+template <typename Dtype>
 class BoxCoderParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   BoxCoderParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                 const AttributeMap &attrs, const Scope &scope) {
-    input_priorbox_ = InputPriorBoxFrom<LoDTensor>(inputs, scope);
-    input_priorboxvar_ = InputPriorBoxVarFrom<LoDTensor>(inputs, scope);
-    input_targetbox_ = InputTargetBoxFrom<LoDTensor>(inputs, scope);
-    output_box_ = OutputBoxFrom<LoDTensor>(outputs, scope);
+    input_priorbox_ = InputPriorBoxFrom<GType>(inputs, scope);
+    input_priorboxvar_ = InputPriorBoxVarFrom<GType>(inputs, scope);
+    input_targetbox_ = InputTargetBoxFrom<GType>(inputs, scope);
+    output_box_ = OutputBoxFrom<GType>(outputs, scope);
     code_type_ = GetAttr<std::string>("code_type", attrs);
   }
-  const Tensor *InputPriorBox() const { return input_priorbox_; }
+  const RType *InputPriorBox() const { return input_priorbox_; }
 
-  const Tensor *InputPriorBoxVar() const { return input_priorboxvar_; }
+  const RType *InputPriorBoxVar() const { return input_priorboxvar_; }
 
-  const Tensor *InputTargetBox() const { return input_targetbox_; }
+  const RType *InputTargetBox() const { return input_targetbox_; }
 
-  Tensor *OutputBox() const { return output_box_; }
+  RType *OutputBox() const { return output_box_; }
 
   const std::string &CodeType() const { return code_type_; }
 
  private:
-  Tensor *input_priorbox_;
-  Tensor *input_priorboxvar_;
-  Tensor *input_targetbox_;
-  Tensor *output_box_;
+  RType *input_priorbox_;
+  RType *input_priorboxvar_;
+  RType *input_targetbox_;
+  RType *output_box_;
   std::string code_type_;
 };
 #endif
 
 #ifdef SOFTMAX_OP
+template <typename Dtype>
 class SoftmaxParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   SoftmaxParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
   }
-  const Tensor *InputX() const { return input_x_; }
-  Tensor *Out() const { return out_; }
+  const RType *InputX() const { return input_x_; }
+  RType *Out() const { return out_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
 
 #ifdef PADDLE_MOBILE_FPGA
 
  private:
-  std::shared_ptr<Tensor> float_input_x_;
+  std::shared_ptr<RType> float_input_x_;
   fpga::BypassArgs fpga_bypass_args;
 
  public:
-  Tensor *FloatInput() {
+  RType *FloatInput() {
     return float_input_x_ == nullptr ? input_x_ : float_input_x_.get();
   }
   void SetFloatInput(Tensor *input) { float_input_x_.reset(input); }
@@ -604,31 +679,39 @@ class SoftmaxParam : public OpParam {
 #endif
 
 #ifdef SIGMOID_OP
+template <typename Dtype>
 class SigmoidParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   SigmoidParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
   }
-  const Tensor *InputX() const { return input_x_; }
-  Tensor *Out() const { return out_; }
+  const RType *InputX() const { return input_x_; }
+  RType *Out() const { return out_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
 };
 #endif
 
 #ifdef MULTICLASSNMS_OP
+template <typename Dtype>
 class MultiClassNMSParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   MultiClassNMSParam(const VariableNameMap &inputs,
                      const VariableNameMap &outputs, const AttributeMap &attrs,
                      const Scope &scope) {
-    input_bboxes_ = InputBBoxesFrom<LoDTensor>(inputs, scope);
-    input_scores_ = InputScoresFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_bboxes_ = InputBBoxesFrom<GType>(inputs, scope);
+    input_scores_ = InputScoresFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     background_label_ = GetAttr<int>("background_label", attrs);
     nms_top_k_ = GetAttr<int>("nms_top_k", attrs);
     keep_top_k_ = GetAttr<int>("keep_top_k", attrs);
@@ -637,11 +720,11 @@ class MultiClassNMSParam : public OpParam {
     score_threshold_ = GetAttr<float>("score_threshold", attrs);
   }
 
-  const Tensor *InputBBoxes() const { return input_bboxes_; }
+  const RType *InputBBoxes() const { return input_bboxes_; }
 
-  const Tensor *InputScores() const { return input_scores_; }
+  const RType *InputScores() const { return input_scores_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &BackGroundLabel() const { return background_label_; }
 
@@ -656,9 +739,9 @@ class MultiClassNMSParam : public OpParam {
   const float &ScoreThreshold() const { return score_threshold_; }
 
  private:
-  Tensor *input_bboxes_;
-  Tensor *input_scores_;
-  Tensor *out_;
+  RType *input_bboxes_;
+  RType *input_scores_;
+  RType *out_;
   int background_label_;
   int nms_top_k_;
   int keep_top_k_;
@@ -668,113 +751,133 @@ class MultiClassNMSParam : public OpParam {
 };
 #endif
 
+template <typename Dtype>
 class FeedParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FeedParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
             const AttributeMap &attrs, Scope *scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, *scope);
-    out_ = OutFrom<LoDTensor>(outputs, *scope);
+    input_x_ = InputXFrom<GType>(inputs, *scope);
+    out_ = OutFrom<GType>(outputs, *scope);
     auto var = scope->Var("batch_size");
     batch_size = var->GetValue<int>();
   }
-  const Tensor *InputX() const { return input_x_; }
-  Tensor *Out() const { return out_; }
+  const RType *InputX() const { return input_x_; }
+  RType *Out() const { return out_; }
   const int BatchSize() const { return batch_size; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
   int batch_size;
 };
 
+template <typename Dtype>
 class FetchParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FetchParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
              const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
   }
-  const Tensor *InputX() const { return input_x_; }
-  Tensor *Out() const { return out_; }
+  const RType *InputX() const { return input_x_; }
+  RType *Out() const { return out_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
 };
 
 #ifdef TRANSPOSE_OP
+template <typename Dtype>
 class TransposeParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   TransposeParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                  const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     axis_ = GetAttr<vector<int>>("axis", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const vector<int> &Axis() const { return axis_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
   vector<int> axis_;
 };
 #endif
 
 #ifdef RESHAPE_OP
+template <typename Dtype>
 class ReshapeParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ReshapeParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_shape_ = InputShapeFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_shape_ = InputShapeFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     shape_ = GetAttr<vector<int>>("shape", attrs);
     inplace_ = GetAttr<bool>("inplace", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputShape() const { return input_shape_; }
+  const RType *InputShape() const { return input_shape_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const vector<int> &Shape() const { return shape_; }
 
   const bool &Inplace() const { return inplace_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_shape_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_shape_;
+  RType *out_;
   vector<int> shape_;
   bool inplace_;
 };
 #endif
 
 #ifdef SCALE_OP
+template <typename Dtype>
 class ScaleParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ScaleParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
              const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_bias_ = InputBiasFrom<framework::LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     inplace_ = GetAttr<bool>("inplace", attrs);
     has_bias_ = GetAttr<bool>("has_bias", attrs);
     scales_ = GetAttr<vector<float>>("scales", attrs);
     biases_ = GetAttr<vector<float>>("biases", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const bool &Inplace() const { return inplace_; }
 
@@ -785,9 +888,9 @@ class ScaleParam : public OpParam {
   const vector<float> &Biases() const { return biases_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_bias_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_bias_;
+  RType *out_;
   bool inplace_;
   bool has_bias_;
   vector<float> scales_;
@@ -796,23 +899,27 @@ class ScaleParam : public OpParam {
 #endif
 
 #ifdef SLICE_OP
+template <typename Dtype>
 class SliceParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   SliceParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
              const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_shape_ = InputShapeFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_shape_ = InputShapeFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
     slice_points_ = GetAttr<vector<int>>("slice_points", attrs);
     inplace_ = GetAttr<bool>("inplace", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputShape() const { return input_shape_; }
+  const RType *InputShape() const { return input_shape_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &Axis() const { return axis_; }
 
@@ -821,9 +928,9 @@ class SliceParam : public OpParam {
   const bool &Inplace() const { return inplace_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_shape_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_shape_;
+  RType *out_;
   int axis_;
   vector<int> slice_points_;
   bool inplace_;
@@ -831,13 +938,17 @@ class SliceParam : public OpParam {
 #endif
 
 #ifdef RESIZE_OP
+template <typename Dtype>
 class ResizeParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ResizeParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
               const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_shape_ = InputShapeFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_shape_ = InputShapeFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     is_pyramid_test_ = GetAttr<bool>("is_pyramid_test", attrs);
     height_ = GetAttr<int>("height", attrs);
     width_ = GetAttr<int>("width", attrs);
@@ -845,11 +956,11 @@ class ResizeParam : public OpParam {
     out_width_scale_ = GetAttr<float>("out_width_scale", attrs);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  const Tensor *InputShape() const { return input_shape_; }
+  const RType *InputShape() const { return input_shape_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const bool &IsPyramidTest() const { return is_pyramid_test_; }
 
@@ -862,9 +973,9 @@ class ResizeParam : public OpParam {
   const float &OutWidthScale() const { return out_width_scale_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_shape_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_shape_;
+  RType *out_;
   bool is_pyramid_test_;
   int height_;
   int width_;
@@ -877,73 +988,85 @@ class ResizeParam : public OpParam {
 /*
  * @b op 层实例化好这个 param 传递给 kernel 层使用
  * */
+template <typename Dtype>
 class ReluParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ReluParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
             const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
 };
 #endif
 
 #ifdef PRELU_OP
+template <typename Dtype>
 class PReluParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   PReluParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
              const AttributeMap &attrs, const Scope &scope) {
     DLOG << "PReluParam inputs before";
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    alpha_ = InputAlphaFrom<LoDTensor>(inputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    alpha_ = InputAlphaFrom<GType>(inputs, scope);
     framework::DDim dims = alpha_->dims();
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     mode_ = GetAttr<std::string>("mode", attrs);
     DLOG << "PReluParam mode after" << mode_;
   }
-  const Tensor *InputX() const { return input_x_; }
-  const Tensor *InputAlpha() const { return alpha_; }
-  Tensor *Out() const { return out_; }
+  const RType *InputX() const { return input_x_; }
+  const RType *InputAlpha() const { return alpha_; }
+  RType *Out() const { return out_; }
   const std::string &Mode() const { return mode_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
-  Tensor *alpha_;
+  RType *input_x_;
+  RType *out_;
+  RType *alpha_;
   std::string mode_;
 };
 #endif
 
+template <typename Dtype>
 class FusionFcParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionFcParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                 const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    input_y_ = InputYFrom<LoDTensor>(inputs, scope);
-    input_z_ = InputZFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_y_ = InputYFrom<GType>(inputs, scope);
+    input_z_ = InputZFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     x_num_col_dims_ = GetAttr<int>("x_num_col_dims", attrs);
     y_num_col_dims_ = GetAttr<int>("y_num_col_dims", attrs);
     axis_ = GetAttr<int>("axis", attrs);
   }
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *InputY() const { return input_y_; }
+  RType *InputY() const { return input_y_; }
 #else
-  const Tensor *InputY() const { return input_y_; }
+  const RType *InputY() const { return input_y_; }
 #endif
 
-  const Tensor *InputZ() const { return input_z_; }
+  const RType *InputZ() const { return input_z_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
   const int &XNumColDims() const { return x_num_col_dims_; }
 
@@ -952,10 +1075,10 @@ class FusionFcParam : public OpParam {
   const int &Axis() const { return axis_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *input_y_;
-  Tensor *input_z_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *input_y_;
+  RType *input_z_;
+  RType *out_;
   int x_num_col_dims_;
   int y_num_col_dims_;
   int axis_;
@@ -971,37 +1094,42 @@ class FusionFcParam : public OpParam {
 };
 
 #ifdef FUSION_FCRELU_OP
-using FusionFcReluParam = FusionFcParam;
+template <typename DeviceType>
+using FusionFcReluParam = FusionFcParam<DeviceType>;
 #endif
 
+template <typename Dtype>
 class FusionConvAddParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionConvAddParam(const VariableNameMap &inputs,
                      const VariableNameMap &outputs, const AttributeMap &attrs,
                      const Scope &scope) {
-    bias_ = InputYFrom<LoDTensor>(inputs, scope);
+    bias_ = InputYFrom<GType>(inputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
   }
-  Tensor *Bias() const { return bias_; }
+  RType *Bias() const { return bias_; }
 
   const int &Axis() const { return axis_; }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 #else
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 #endif
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1012,11 +1140,11 @@ class FusionConvAddParam : public OpParam {
   const int &Groups() const { return groups; }
 
  protected:
-  Tensor *bias_;
+  RType *bias_;
   int axis_;
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
@@ -1032,54 +1160,60 @@ class FusionConvAddParam : public OpParam {
 #endif
 };
 
-Print &operator<<(Print &printer, const FusionConvAddParam &conv_param);
+template <typename Dtype>
+Print &operator<<(Print &printer, const FusionConvAddParam<Dtype> &conv_param);
 
 #ifdef FUSION_CONVADDRELU_OP
-class FusionConvAddReluParam : public FusionConvAddParam {
+template <typename DeviceType>
+class FusionConvAddReluParam : public FusionConvAddParam<DeviceType> {
  public:
   FusionConvAddReluParam(const VariableNameMap &inputs,
                          const VariableNameMap &outputs,
                          const AttributeMap &attrs, const Scope &scope)
-      : FusionConvAddParam(inputs, outputs, attrs, scope) {}
+      : FusionConvAddParam<DeviceType>(inputs, outputs, attrs, scope) {}
 };
 #endif
 
 #ifdef FUSION_CONVADDBNRELU_OP
+template <typename Dtype>
 class FusionConvAddBNReluParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionConvAddBNReluParam(const VariableNameMap &inputs,
                            const VariableNameMap &outputs,
                            const AttributeMap &attrs, const Scope &scope) {
-    bias_ = InputYFrom<LoDTensor>(inputs, scope);
+    bias_ = InputYFrom<GType>(inputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
-  Tensor *Bias() const { return bias_; }
+  RType *Bias() const { return bias_; }
 
   const int &Axis() const { return axis_; }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 #else
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 #endif
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1089,13 +1223,13 @@ class FusionConvAddBNReluParam : public OpParam {
 
   const int &Groups() const { return groups; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -1103,33 +1237,33 @@ class FusionConvAddBNReluParam : public OpParam {
 
   const bool &IsTest() const { return is_test_; }
 
-  void SetNewScale(Tensor *new_scale) { new_scale_ = new_scale; }
+  void SetNewScale(RType *new_scale) { new_scale_ = new_scale; }
 
-  void SetNewBias(Tensor *new_bias) { new_bias_ = new_bias; }
+  void SetNewBias(RType *new_bias) { new_bias_ = new_bias; }
 
-  const Tensor *NewScale() const { return new_scale_; }
+  const RType *NewScale() const { return new_scale_; }
 
-  const Tensor *NewBias() const { return new_bias_; }
+  const RType *NewBias() const { return new_bias_; }
 
  protected:
-  Tensor *bias_;
+  RType *bias_;
   int axis_;
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
-  Tensor *new_bias_;
-  Tensor *new_scale_;
+  RType *new_bias_;
+  RType *new_scale_;
 #ifdef PADDLE_MOBILE_FPGA
 
  private:
@@ -1143,35 +1277,39 @@ class FusionConvAddBNReluParam : public OpParam {
 #endif
 
 #ifdef FUSION_CONVBN_OP
+template <typename Dtype>
 class FusionConvBNParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionConvBNParam(const VariableNameMap &inputs,
                     const VariableNameMap &outputs, const AttributeMap &attrs,
                     const Scope &scope) {
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_y_ = OutputYFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_y_ = OutputYFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 #else
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 #endif
-  Tensor *Output() const { return output_y_; }
+  RType *Output() const { return output_y_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1181,13 +1319,13 @@ class FusionConvBNParam : public OpParam {
 
   const int &Groups() const { return groups; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -1195,31 +1333,31 @@ class FusionConvBNParam : public OpParam {
 
   const bool &IsTest() const { return is_test_; }
 
-  void SetNewScale(Tensor *new_scale) { new_scale_ = new_scale; }
+  void SetNewScale(RType *new_scale) { new_scale_ = new_scale; }
 
-  void SetNewBias(Tensor *new_bias) { new_bias_ = new_bias; }
+  void SetNewBias(RType *new_bias) { new_bias_ = new_bias; }
 
-  const Tensor *NewScale() const { return new_scale_; }
+  const RType *NewScale() const { return new_scale_; }
 
-  const Tensor *NewBias() const { return new_bias_; }
+  const RType *NewBias() const { return new_bias_; }
 
  protected:
-  Tensor *input_;
-  Tensor *output_y_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_y_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
-  Tensor *new_bias_;
-  Tensor *new_scale_;
+  RType *new_bias_;
+  RType *new_scale_;
 #ifdef PADDLE_MOBILE_FPGA
 
  private:
@@ -1233,40 +1371,44 @@ class FusionConvBNParam : public OpParam {
 #endif
 
 #ifdef FUSION_CONVADDBN_OP
+template <typename Dtype>
 class FusionConvAddBNParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionConvAddBNParam(const VariableNameMap &inputs,
                        const VariableNameMap &outputs,
                        const AttributeMap &attrs, const Scope &scope) {
-    bias_ = InputYFrom<LoDTensor>(inputs, scope);
+    bias_ = InputYFrom<GType>(inputs, scope);
     axis_ = GetAttr<int>("axis", attrs);
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_y_ = OutputYFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_y_ = OutputYFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
-  Tensor *Bias() const { return bias_; }
+  RType *Bias() const { return bias_; }
 
   const int &Axis() const { return axis_; }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 #else
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 #endif
-  Tensor *Output() const { return output_y_; }
+  RType *Output() const { return output_y_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1276,13 +1418,13 @@ class FusionConvAddBNParam : public OpParam {
 
   const int &Groups() const { return groups; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -1290,33 +1432,33 @@ class FusionConvAddBNParam : public OpParam {
 
   const bool &IsTest() const { return is_test_; }
 
-  void SetNewScale(Tensor *new_scale) { new_scale_ = new_scale; }
+  void SetNewScale(RType *new_scale) { new_scale_ = new_scale; }
 
-  void SetNewBias(Tensor *new_bias) { new_bias_ = new_bias; }
+  void SetNewBias(RType *new_bias) { new_bias_ = new_bias; }
 
-  const Tensor *NewScale() const { return new_scale_; }
+  const RType *NewScale() const { return new_scale_; }
 
-  const Tensor *NewBias() const { return new_bias_; }
+  const RType *NewBias() const { return new_bias_; }
 
  protected:
-  Tensor *bias_;
+  RType *bias_;
   int axis_;
-  Tensor *input_;
-  Tensor *output_y_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_y_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
-  Tensor *new_bias_;
-  Tensor *new_scale_;
+  RType *new_bias_;
+  RType *new_scale_;
 #ifdef PADDLE_MOBILE_FPGA
 
  private:
@@ -1330,32 +1472,36 @@ class FusionConvAddBNParam : public OpParam {
 #endif
 
 #ifdef FUSION_DWCONVBNRELU_OP
+template <typename Dtype>
 class FusionDWConvBNReluParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionDWConvBNReluParam(const VariableNameMap &inputs,
                           const VariableNameMap &outputs,
                           const AttributeMap &attrs, const Scope &scope) {
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1365,13 +1511,13 @@ class FusionDWConvBNReluParam : public OpParam {
 
   const int &Groups() const { return groups; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -1379,68 +1525,71 @@ class FusionDWConvBNReluParam : public OpParam {
 
   const bool &IsTest() const { return is_test_; }
 
-  void SetNewScale(Tensor *new_scale) { new_scale_ = new_scale; }
+  void SetNewScale(RType *new_scale) { new_scale_ = new_scale; }
 
-  void SetNewBias(Tensor *new_bias) { new_bias_ = new_bias; }
+  void SetNewBias(RType *new_bias) { new_bias_ = new_bias; }
 
-  const Tensor *NewScale() const { return new_scale_; }
+  const RType *NewScale() const { return new_scale_; }
 
-  const Tensor *NewBias() const { return new_bias_; }
+  const RType *NewBias() const { return new_bias_; }
 
  protected:
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
-  Tensor *new_bias_;
-  Tensor *new_scale_;
+  RType *new_bias_;
+  RType *new_scale_;
 };
 
-Print &operator<<(Print &printer, const FusionConvAddParam &conv_param);
 #endif
 
 #ifdef FUSION_CONVBNRELU_OP
+template <typename Dtype>
 class FusionConvBNReluParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   FusionConvBNReluParam(const VariableNameMap &inputs,
                         const VariableNameMap &outputs,
                         const AttributeMap &attrs, const Scope &scope) {
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
 
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
-    input_bias_ = InputBiasFrom<LoDTensor>(inputs, scope);
-    input_mean_ = InputMeanFrom<LoDTensor>(inputs, scope);
-    input_scale_ = InputScaleFrom<LoDTensor>(inputs, scope);
-    input_variance_ = InputVarianceFrom<LoDTensor>(inputs, scope);
+    input_bias_ = InputBiasFrom<GType>(inputs, scope);
+    input_mean_ = InputMeanFrom<GType>(inputs, scope);
+    input_scale_ = InputScaleFrom<GType>(inputs, scope);
+    input_variance_ = InputVarianceFrom<GType>(inputs, scope);
     epsilon_ = GetAttr<float>("epsilon", attrs);
     momentum_ = GetAttr<float>("momentum", attrs);
     //    is_test_ = GetAttr<bool>("is_test", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
 #ifdef PADDLE_MOBILE_FPGA
-  Tensor *Filter() const { return filter_; }
+  RType *Filter() const { return filter_; }
 #else
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 #endif
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1450,13 +1599,13 @@ class FusionConvBNReluParam : public OpParam {
 
   const int &Groups() const { return groups; }
 
-  const Tensor *InputBias() const { return input_bias_; }
+  const RType *InputBias() const { return input_bias_; }
 
-  const Tensor *InputMean() const { return input_mean_; }
+  const RType *InputMean() const { return input_mean_; }
 
-  const Tensor *InputScale() const { return input_scale_; }
+  const RType *InputScale() const { return input_scale_; }
 
-  const Tensor *InputVariance() const { return input_variance_; }
+  const RType *InputVariance() const { return input_variance_; }
 
   const float &Epsilon() const { return epsilon_; }
 
@@ -1464,31 +1613,31 @@ class FusionConvBNReluParam : public OpParam {
 
   const bool &IsTest() const { return is_test_; }
 
-  void SetNewScale(Tensor *new_scale) { new_scale_ = new_scale; }
+  void SetNewScale(RType *new_scale) { new_scale_ = new_scale; }
 
-  void SetNewBias(Tensor *new_bias) { new_bias_ = new_bias; }
+  void SetNewBias(RType *new_bias) { new_bias_ = new_bias; }
 
-  const Tensor *NewScale() const { return new_scale_; }
+  const RType *NewScale() const { return new_scale_; }
 
-  const Tensor *NewBias() const { return new_bias_; }
+  const RType *NewBias() const { return new_bias_; }
 
  protected:
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;
   int groups;
-  Tensor *input_bias_;
-  Tensor *input_mean_;
-  Tensor *input_scale_;
-  Tensor *input_variance_;
+  RType *input_bias_;
+  RType *input_mean_;
+  RType *input_scale_;
+  RType *input_variance_;
   float epsilon_;
   float momentum_;
   bool is_test_;
-  Tensor *new_bias_;
-  Tensor *new_scale_;
+  RType *new_bias_;
+  RType *new_scale_;
 #ifdef PADDLE_MOBILE_FPGA
 
  private:
@@ -1502,21 +1651,25 @@ class FusionConvBNReluParam : public OpParam {
 #endif
 
 #ifdef IM2SEQUENCE_OP
+template <typename Dtype>
 class Im2SequenceParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   Im2SequenceParam(const VariableNameMap &inputs,
                    const VariableNameMap &outputs, const AttributeMap &attrs,
                    const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
     kernels_ = GetAttr<vector<int>>("kernels", attrs);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
   }
 
-  const Tensor *Input() const { return input_x_; }
+  const RType *Input() const { return input_x_; }
 
-  Tensor *Output() const { return out_; }
+  RType *Output() const { return out_; }
 
   const vector<int> &Kernels() const { return kernels_; }
 
@@ -1525,8 +1678,8 @@ class Im2SequenceParam : public OpParam {
   const vector<int> &Paddings() const { return paddings_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
   vector<int> kernels_;
   vector<int> strides_;
   vector<int> paddings_;
@@ -1534,44 +1687,52 @@ class Im2SequenceParam : public OpParam {
 #endif
 
 #ifdef DROPOUT_OP
+template <typename Dtype>
 class DropoutParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   DropoutParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
                const AttributeMap &attrs, const Scope &scope) {
-    input_x_ = InputXFrom<LoDTensor>(inputs, scope);
-    out_ = OutFrom<LoDTensor>(outputs, scope);
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    out_ = OutFrom<GType>(outputs, scope);
   }
 
-  const Tensor *InputX() const { return input_x_; }
+  const RType *InputX() const { return input_x_; }
 
-  Tensor *Out() const { return out_; }
+  RType *Out() const { return out_; }
 
  private:
-  Tensor *input_x_;
-  Tensor *out_;
+  RType *input_x_;
+  RType *out_;
 };
 #endif
 
 #ifdef CONV_TRANSPOSE
+template <typename Dtype>
 class ConvTransposeParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
  public:
   ConvTransposeParam(const VariableNameMap &inputs,
                      const VariableNameMap &outputs, const AttributeMap &attrs,
                      const Scope &scope) {
-    filter_ = FilterFrom<LoDTensor>(inputs, scope);
-    input_ = InputFrom<LoDTensor>(inputs, scope);
-    output_ = OutputFrom<LoDTensor>(outputs, scope);
+    filter_ = FilterFrom<GType>(inputs, scope);
+    input_ = InputFrom<GType>(inputs, scope);
+    output_ = OutputFrom<GType>(outputs, scope);
     strides_ = GetAttr<vector<int>>("strides", attrs);
     paddings_ = GetAttr<vector<int>>("paddings", attrs);
     dilations_ = GetAttr<vector<int>>("dilations", attrs);
     groups = GetAttr<int>("groups", attrs);
   }
 
-  const Tensor *Input() const { return input_; }
+  const RType *Input() const { return input_; }
 
-  const Tensor *Filter() const { return filter_; }
+  const RType *Filter() const { return filter_; }
 
-  Tensor *Output() const { return output_; }
+  RType *Output() const { return output_; }
 
   const vector<int> &Strides() const { return strides_; }
 
@@ -1582,9 +1743,9 @@ class ConvTransposeParam : public OpParam {
   const int &Groups() const { return groups; }
 
  private:
-  Tensor *input_;
-  Tensor *output_;
-  Tensor *filter_;
+  RType *input_;
+  RType *output_;
+  RType *filter_;
   vector<int> strides_;
   vector<int> paddings_;
   vector<int> dilations_;

src/operators/pool_op.h

@@ -29,17 +29,17 @@ using framework::OperatorWithKernel;
 using framework::Scope;
 using std::string;
 template <typename DeviceType, typename T>
-class PoolOp : public OperatorWithKernel<DeviceType, PoolParam,
+class PoolOp : public OperatorWithKernel<DeviceType, PoolParam<DeviceType>,
                                          operators::PoolKernel<DeviceType, T>> {
  public:
   PoolOp(const string &type, const VariableNameMap &inputs,
          const VariableNameMap &outputs, const AttributeMap &attrs,
          std::shared_ptr<Scope> scope)
-      : OperatorWithKernel<DeviceType, PoolParam,
+      : OperatorWithKernel<DeviceType, PoolParam<DeviceType>,
                            operators::PoolKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
   using OperatorWithKernel<
-      DeviceType, PoolParam,
+      DeviceType, PoolParam<DeviceType>,
       operators::PoolKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/prelu_op.h

@@ -28,19 +28,19 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class PReluOp
-    : public framework::OperatorWithKernel<
-          DeviceType, PReluParam, operators::PReluKernel<DeviceType, T>> {
+class PReluOp : public framework::OperatorWithKernel<
+                    DeviceType, PReluParam<DeviceType>,
+                    operators::PReluKernel<DeviceType, T>> {
  public:
   PReluOp(const std::string &type, const VariableNameMap &inputs,
           const VariableNameMap &outputs, const framework::AttributeMap &attrs,
           std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, PReluParam,
+      : framework::OperatorWithKernel<DeviceType, PReluParam<DeviceType>,
                                       operators::PReluKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, PReluParam,
+      DeviceType, PReluParam<DeviceType>,
       operators::PReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/prior_box_op.h

@@ -28,20 +28,20 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class PriorBoxOp
-    : public framework::OperatorWithKernel<
-          DeviceType, PriorBoxParam, operators::PriorBoxKernel<DeviceType, T>> {
+class PriorBoxOp : public framework::OperatorWithKernel<
+                       DeviceType, PriorBoxParam<DeviceType>,
+                       operators::PriorBoxKernel<DeviceType, T>> {
  public:
   PriorBoxOp(const std::string &type, const VariableNameMap &inputs,
              const VariableNameMap &outputs,
              const framework::AttributeMap &attrs,
              std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, PriorBoxParam,
+      : framework::OperatorWithKernel<DeviceType, PriorBoxParam<DeviceType>,
                                       operators::PriorBoxKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, PriorBoxParam,
+      DeviceType, PriorBoxParam<DeviceType>,
       operators::PriorBoxKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/relu_op.h

@@ -28,9 +28,9 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class ReluOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ReluParam, operators::ReluKernel<DeviceType, T>> {
+class ReluOp : public framework::OperatorWithKernel<
+                   DeviceType, ReluParam<DeviceType>,
+                   operators::ReluKernel<DeviceType, T>> {
  public:
   /*
    * @b op 的实例化方法, 需要调用父类的实例化方法, 以及实例化自己的参数结构体
@@ -38,12 +38,12 @@ class ReluOp
   ReluOp(const std::string &type, const VariableNameMap &inputs,
          const VariableNameMap &outputs, const framework::AttributeMap &attrs,
          std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ReluParam,
+      : framework::OperatorWithKernel<DeviceType, ReluParam<DeviceType>,
                                       operators::ReluKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ReluParam,
+      DeviceType, ReluParam<DeviceType>,
       operators::ReluKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/reshape_op.h

@@ -28,20 +28,20 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class ReshapeOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ReshapeParam, operators::ReshapeKernel<DeviceType, T>> {
+class ReshapeOp : public framework::OperatorWithKernel<
+                      DeviceType, ReshapeParam<DeviceType>,
+                      operators::ReshapeKernel<DeviceType, T>> {
  public:
   ReshapeOp(const std::string &type, const VariableNameMap &inputs,
             const VariableNameMap &outputs,
             const framework::AttributeMap &attrs,
             std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ReshapeParam,
+      : framework::OperatorWithKernel<DeviceType, ReshapeParam<DeviceType>,
                                       operators::ReshapeKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ReshapeParam,
+      DeviceType, ReshapeParam<DeviceType>,
       operators::ReshapeKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/resize_op.h

@@ -28,19 +28,19 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class ResizeOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ResizeParam, operators::ResizeKernel<DeviceType, T>> {
+class ResizeOp : public framework::OperatorWithKernel<
+                     DeviceType, ResizeParam<DeviceType>,
+                     operators::ResizeKernel<DeviceType, T>> {
  public:
   ResizeOp(const std::string &type, const VariableNameMap &inputs,
            const VariableNameMap &outputs, const framework::AttributeMap &attrs,
            std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ResizeParam,
+      : framework::OperatorWithKernel<DeviceType, ResizeParam<DeviceType>,
                                       operators::ResizeKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ResizeParam,
+      DeviceType, ResizeParam<DeviceType>,
       operators::ResizeKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/scale_op.h

@@ -28,19 +28,19 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class ScaleOp
-    : public framework::OperatorWithKernel<
-          DeviceType, ScaleParam, operators::ScaleKernel<DeviceType, T>> {
+class ScaleOp : public framework::OperatorWithKernel<
+                    DeviceType, ScaleParam<DeviceType>,
+                    operators::ScaleKernel<DeviceType, T>> {
  public:
   ScaleOp(const std::string &type, const VariableNameMap &inputs,
           const VariableNameMap &outputs, const framework::AttributeMap &attrs,
           std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, ScaleParam,
+      : framework::OperatorWithKernel<DeviceType, ScaleParam<DeviceType>,
                                       operators::ScaleKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, ScaleParam,
+      DeviceType, ScaleParam<DeviceType>,
       operators::ScaleKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/sigmoid_op.h

@@ -25,20 +25,20 @@ limitations under the License. */
 namespace paddle_mobile {
 namespace operators {
 template <typename DeviceType, typename T>
-class SigmoidOp
-    : public framework::OperatorWithKernel<
-          DeviceType, SigmoidParam, operators::SigmoidKernel<DeviceType, T>> {
+class SigmoidOp : public framework::OperatorWithKernel<
+                      DeviceType, SigmoidParam<DeviceType>,
+                      operators::SigmoidKernel<DeviceType, T>> {
  public:
   SigmoidOp(const std::string &type, const VariableNameMap &inputs,
             const VariableNameMap &outputs,
             const framework::AttributeMap &attrs,
             std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, SigmoidParam,
+      : framework::OperatorWithKernel<DeviceType, SigmoidParam<DeviceType>,
                                       operators::SigmoidKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, SigmoidParam,
+      DeviceType, SigmoidParam<DeviceType>,
       operators::SigmoidKernel<DeviceType, T>>::OperatorWithKernel;
 
   void InferShape() const override;

src/operators/slice_op.h

@@ -28,19 +28,19 @@ namespace operators {
 using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
-class SliceOp
-    : public framework::OperatorWithKernel<
-          DeviceType, SliceParam, operators::SliceKernel<DeviceType, T>> {
+class SliceOp : public framework::OperatorWithKernel<
+                    DeviceType, SliceParam<DeviceType>,
+                    operators::SliceKernel<DeviceType, T>> {
  public:
   SliceOp(const std::string &type, const VariableNameMap &inputs,
           const VariableNameMap &outputs, const framework::AttributeMap &attrs,
           std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, SliceParam,
+      : framework::OperatorWithKernel<DeviceType, SliceParam<DeviceType>,
                                       operators::SliceKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, SliceParam,
+      DeviceType, SliceParam<DeviceType>,
       operators::SliceKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 

src/operators/softmax_op.h

@@ -25,20 +25,20 @@ limitations under the License. */
 namespace paddle_mobile {
 namespace operators {
 template <typename DeviceType, typename T>
-class SoftmaxOp
-    : public framework::OperatorWithKernel<
-          DeviceType, SoftmaxParam, operators::SoftmaxKernel<DeviceType, T>> {
+class SoftmaxOp : public framework::OperatorWithKernel<
+                      DeviceType, SoftmaxParam<DeviceType>,
+                      operators::SoftmaxKernel<DeviceType, T>> {
  public:
   SoftmaxOp(const std::string &type, const VariableNameMap &inputs,
             const VariableNameMap &outputs,
             const framework::AttributeMap &attrs,
             std::shared_ptr<framework::Scope> scope)
-      : framework::OperatorWithKernel<DeviceType, SoftmaxParam,
+      : framework::OperatorWithKernel<DeviceType, SoftmaxParam<DeviceType>,
                                       operators::SoftmaxKernel<DeviceType, T>>(
             type, inputs, outputs, attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, SoftmaxParam,
+      DeviceType, SoftmaxParam<DeviceType>,
       operators::SoftmaxKernel<DeviceType, T>>::OperatorWithKernel;
 
   void InferShape() const override;

src/operators/transpose_op.h

@@ -29,7 +29,7 @@ using paddle_mobile::framework::Tensor;
 
 template <typename DeviceType, typename T>
 class TransposeOp : public framework::OperatorWithKernel<
-                        DeviceType, TransposeParam,
+                        DeviceType, TransposeParam<DeviceType>,
                         operators::TransposeKernel<DeviceType, T>> {
  public:
   TransposeOp(const std::string &type, const VariableNameMap &inputs,
@@ -37,12 +37,12 @@ class TransposeOp : public framework::OperatorWithKernel<
               const framework::AttributeMap &attrs,
               std::shared_ptr<framework::Scope> scope)
       : framework::OperatorWithKernel<
-            DeviceType, TransposeParam,
+            DeviceType, TransposeParam<DeviceType>,
             operators::TransposeKernel<DeviceType, T>>(type, inputs, outputs,
                                                        attrs, scope) {}
 
   using framework::OperatorWithKernel<
-      DeviceType, TransposeParam,
+      DeviceType, TransposeParam<DeviceType>,
       operators::TransposeKernel<DeviceType, T>>::OperatorWithKernel;
   void InferShape() const override;
 };