diff --git a/src/operators/batchnorm_op.h b/src/operators/batchnorm_op.h
index 2b2795b64fddfbcd1000088dbab18e54a017b459..52c423f1bb90428e867ea6fb992036ab83c683d7 100644
--- a/src/operators/batchnorm_op.h
+++ b/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) {}
 
diff --git a/src/operators/box_coder_op.h b/src/operators/box_coder_op.h
index 5a75cacaf27f20e69b5e427421bd3dd8f43e8556..5d475c98b6859a33b39e6b36419fa055cde7a1d3 100644
--- a/src/operators/box_coder_op.h
+++ b/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;
 
diff --git a/src/operators/concat_op.h b/src/operators/concat_op.h
index bad0015917c2a9d4016def26c8a332d076b39c99..a169c17dc468dd06ed344a0c7a6ef3cb2c977a27 100644
--- a/src/operators/concat_op.h
+++ b/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;
 
diff --git a/src/operators/conv_op.h b/src/operators/conv_op.h
index d36fa47f4a0b37c467eb2101e2e930fe54a0e28b..267abfeb614dc8e19a2cf0cf43e7c5f232a62072 100644
--- a/src/operators/conv_op.h
+++ b/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;
 
diff --git a/src/operators/conv_transpose_op.h b/src/operators/conv_transpose_op.h
index 81fa6916477fc1be99fe3daf5933a1ca7c726363..c9b5e86bef0674b176ba901212a9add2ee2def83 100644
--- a/src/operators/conv_transpose_op.h
+++ b/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) {}
 
diff --git a/src/operators/depthwise_conv_op.h b/src/operators/depthwise_conv_op.h
index 9d7cbcfa2f2924db040cdc5f38ca6bb7ad8074b5..40e87a9b1bf9d2b5102a56ff59821b9d122563c5 100644
--- a/src/operators/depthwise_conv_op.h
+++ b/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;
 
diff --git a/src/operators/dropout_op.h b/src/operators/dropout_op.h
index 89d658dd8c1e11576a2cc0ef9cceae3fcdf26477..7523fd5fa8f21dfce20bce963be4b3bc323948e9 100644
--- a/src/operators/dropout_op.h
+++ b/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;
diff --git a/src/operators/elementwise_add_op.h b/src/operators/elementwise_add_op.h
index 761a5d35459558d1ca5673757fae13147b7f6a6f..14bcd5264d136007e2eb2ffe917697570b32e40b 100644
--- a/src/operators/elementwise_add_op.h
+++ b/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;
 
diff --git a/src/operators/feed_op.h b/src/operators/feed_op.h
index fb219a48196ab584a9c1bfd3832b4e98fde00c99..1edcb8406360becba25d522df0f67836519fd1de 100644
--- a/src/operators/feed_op.h
+++ b/src/operators/feed_op.h
@@ -70,7 +70,7 @@ class FeedOp : public framework::OperatorBase<DeviceType> {
 #endif
 
  protected:
-  FeedParam param_;
+  FeedParam<DeviceType> param_;
 };
 
 }  // namespace operators
diff --git a/src/operators/fetch_op.h b/src/operators/fetch_op.h
index 417637c80086b099395e93227991491309f656fe..1efe0832b1fc4b2ce240ed838e2f4554c29dccd9 100644
--- a/src/operators/fetch_op.h
+++ b/src/operators/fetch_op.h
@@ -41,7 +41,7 @@ class FetchOp : public framework::OperatorBase<DeviceType> {
   }
 
  protected:
-  FetchParam param_;
+  FetchParam<DeviceType> param_;
 };
 
 }  // namespace operators
diff --git a/src/operators/fusion_conv_add.h b/src/operators/fusion_conv_add.h
index 170df9ce33e4ab90297664fbc81d723e7c246f83..ba1ca997662ce67fdcd8f39d2a12e2f535c5b1a7 100644
--- a/src/operators/fusion_conv_add.h
+++ b/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;
 
diff --git a/src/operators/fusion_conv_add_bn_op.h b/src/operators/fusion_conv_add_bn_op.h
index 7a7f6b2bababd3f5d36d7b6faf60069567d45423..ec10787697deb006fe03a35192efb0d80bd00a3c 100644
--- a/src/operators/fusion_conv_add_bn_op.h
+++ b/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) {}
 
diff --git a/src/operators/fusion_conv_add_bn_relu_op.h b/src/operators/fusion_conv_add_bn_relu_op.h
index 54e7e58f8af4111edd0b86c85bb1cffc87f5cd22..4dee4416622e8dee8ca495026843c7506d084617 100644
--- a/src/operators/fusion_conv_add_bn_relu_op.h
+++ b/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;
 
diff --git a/src/operators/fusion_conv_add_relu_op.h b/src/operators/fusion_conv_add_relu_op.h
index e8a9498819cae330abbd4a007a6510d89f167114..926f309403d37fa8ec1f15f7cb955c1c13842405 100644
--- a/src/operators/fusion_conv_add_relu_op.h
+++ b/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;
 
diff --git a/src/operators/fusion_conv_bn_op.h b/src/operators/fusion_conv_bn_op.h
index 620a22b19ece7cd4945c3af20e20c35d02ccbef4..f43e62c9fa5c4b40c07fcb9cbdab4d06ab2c482f 100644
--- a/src/operators/fusion_conv_bn_op.h
+++ b/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) {}
 
diff --git a/src/operators/fusion_conv_bn_relu_op.h b/src/operators/fusion_conv_bn_relu_op.h
index 7ad72acbdc31b3a0527f544ff8521142e58755d7..2b5ff4ea9d3e77ad9449b3968667ecc4558c2147 100644
--- a/src/operators/fusion_conv_bn_relu_op.h
+++ b/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;
 
diff --git a/src/operators/fusion_dwconv_bn_relu_op.h b/src/operators/fusion_dwconv_bn_relu_op.h
index 6f9f03e4936e082de802ced385060fecb9cc27a9..dd1f85688f576106a46cd3070ab2034ec8f55881 100644
--- a/src/operators/fusion_dwconv_bn_relu_op.h
+++ b/src/operators/fusion_dwconv_bn_relu_op.h
@@ -51,21 +51,22 @@ class FusionDWConvBNReluMatcher : public framework::FusionOpMatcher {
 };
 
 template <typename DeviceType, typename T>
-class FusionDWConvBNReluOp : public framework::OperatorWithKernel<
-                                 DeviceType, FusionDWConvBNReluParam,
-                                 operators::DWConvBNReluKernel<DeviceType, T>> {
+class FusionDWConvBNReluOp
+    : public framework::OperatorWithKernel<
+          DeviceType, FusionDWConvBNReluParam<DeviceType>,
+          operators::DWConvBNReluKernel<DeviceType, T>> {
  public:
   FusionDWConvBNReluOp(const string &type, const VariableNameMap &inputs,
                        const VariableNameMap &outputs,
                        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;
 
diff --git a/src/operators/fusion_elementwise_add_relu_op.h b/src/operators/fusion_elementwise_add_relu_op.h
index 286b1efcc85d480aaae614bf6bd72646f327b048..2a92f1e2471cb9e14d84ef03e4bfb872fc738d68 100644
--- a/src/operators/fusion_elementwise_add_relu_op.h
+++ b/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) {}
 
diff --git a/src/operators/fusion_fc_op.h b/src/operators/fusion_fc_op.h
index a6fefc16a80a8d2c1d645ae0e6d7394366209de3..076a95d745e8d44a417dd95fb75844a67b11e653 100644
--- a/src/operators/fusion_fc_op.h
+++ b/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;
 
diff --git a/src/operators/fusion_fc_relu_op.h b/src/operators/fusion_fc_relu_op.h
index 10bbf5580220ea0bce9c3845e42845e3fac1085e..fa7d4045fc10d6e240d93e129aa736be793f7bbf 100644
--- a/src/operators/fusion_fc_relu_op.h
+++ b/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;
 
diff --git a/src/operators/im2sequence_op.h b/src/operators/im2sequence_op.h
index 0695da9308d33ca2b86a5e052210507beb9a82d3..edb87d0012e5514cb5541f94a965965f3dc02825 100644
--- a/src/operators/im2sequence_op.h
+++ b/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;
 
diff --git a/src/operators/kernel/arm/batchnorm_kernel.cpp b/src/operators/kernel/arm/batchnorm_kernel.cpp
index f78d1fdc95ac9e10619dbf32fdc84d01a370f315..c420727f425092240994ee834117225c72abeec2 100644
--- a/src/operators/kernel/arm/batchnorm_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/box_coder_kernel.cpp b/src/operators/kernel/arm/box_coder_kernel.cpp
index d2a479391fbbb416eea7d19ae64125cac4637ef1..b769d4fbbaa7570ee741476f960d9e5b60c61917 100644
--- a/src/operators/kernel/arm/box_coder_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/concat_kernel.cpp b/src/operators/kernel/arm/concat_kernel.cpp
index b6810bf76946bfb8151f3001b76fcbaa5e99e5fc..33cee24900ed185e62b72895ead83f8170463253 100644
--- a/src/operators/kernel/arm/concat_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/conv_add_bn_relu_kernel.cpp b/src/operators/kernel/arm/conv_add_bn_relu_kernel.cpp
index dbf3745eb15cf56bba32dc8cbae50d242ce2da76..ca53ebea8e4577fdc52fad066691d4351eaf12f9 100644
--- a/src/operators/kernel/arm/conv_add_bn_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/arm/conv_add_kernel.cpp b/src/operators/kernel/arm/conv_add_kernel.cpp
index 88f839f611f1ed7f46c11a1b24feb6e29ff07ec7..1af1c3db1159cd4fed007ebf153ba15b804eee75 100644
--- a/src/operators/kernel/arm/conv_add_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/conv_add_relu_kernel.cpp b/src/operators/kernel/arm/conv_add_relu_kernel.cpp
index 8414b7374dd0ed2b10784563dbac9c1565d66f4c..f50e1e3900bb5fce35a29100d6c2cb6004b4af74 100644
--- a/src/operators/kernel/arm/conv_add_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/arm/conv_bn_relu_kernel.cpp b/src/operators/kernel/arm/conv_bn_relu_kernel.cpp
index 23f06c1f0b8a0ed3f22ca9d23d24ae44c59f3618..6b9ea0428fa496980a234c7c895ef9cbf1245b51 100644
--- a/src/operators/kernel/arm/conv_bn_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/arm/conv_kernel.cpp b/src/operators/kernel/arm/conv_kernel.cpp
index ca8aeff0dd3db5fe7b625bdeb947b2927eb619ce..4e9d3a34f231485685bc8f7b087382cb99a3b036 100644
--- a/src/operators/kernel/arm/conv_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/conv_transpose_kernel.cpp b/src/operators/kernel/arm/conv_transpose_kernel.cpp
index d3c1aa96fd34207d401ef96058713d9b93098347..d695e6144b40d945857d547f7c208f1192481e8f 100644
--- a/src/operators/kernel/arm/conv_transpose_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/depthwise_conv_kernel.cpp b/src/operators/kernel/arm/depthwise_conv_kernel.cpp
index 6ede0e2bef2383df8aa0593a07297f2f6233acaf..fd5e068afb6f7f2a069a7d8fccc459d4c2a6828d 100644
--- a/src/operators/kernel/arm/depthwise_conv_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/dropout_kernel.cpp b/src/operators/kernel/arm/dropout_kernel.cpp
index db942b018d7085ca3986533937328101afb08ff9..348b74cd26c79f5d0f2e65012a5760ab7eb58707 100644
--- a/src/operators/kernel/arm/dropout_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/arm/dwconv_bn_relu_kernel.cpp b/src/operators/kernel/arm/dwconv_bn_relu_kernel.cpp
index 0ec08fcecb9fefaa247e0acbb8a085e752b8dba3..b85701bb936b2ccc0323e4d534424abb726a69be 100644
--- a/src/operators/kernel/arm/dwconv_bn_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/arm/elementwise_add_kernel.cpp b/src/operators/kernel/arm/elementwise_add_kernel.cpp
index fdab1c60a310480d8e59f3f84802001ea592433a..af6067f80c98ab7ce62f14c67021bc67a2bf2a9c 100644
--- a/src/operators/kernel/arm/elementwise_add_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/fusion_fc_kernel.cpp b/src/operators/kernel/arm/fusion_fc_kernel.cpp
index c72960e67f19c601e6f27a3bedf7123c80875e0c..1178c980f97792b9574c16f94528f65a277cda80 100644
--- a/src/operators/kernel/arm/fusion_fc_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/im2sequence_kernel.cpp b/src/operators/kernel/arm/im2sequence_kernel.cpp
index 709fa30a23d4efba3531d9bc567c99f53875bc12..8295fd94a31db2ad1c10d32a8c639b067e422f45 100644
--- a/src/operators/kernel/arm/im2sequence_kernel.cpp
+++ b/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>();
diff --git a/src/operators/kernel/arm/lrn_kernel.cpp b/src/operators/kernel/arm/lrn_kernel.cpp
index 0c20c5167adee5165067cc5ab4935df255751755..3ec1bdd9a0e2ebbce555eef944fe56750505430f 100644
--- a/src/operators/kernel/arm/lrn_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/mul_kernel.cpp b/src/operators/kernel/arm/mul_kernel.cpp
index ac5010ce5492ae1d99e59bfa761e22bb3aa5d1c9..e8ef8fe80cd92cd725be8c0251aa4a4f239dbcee 100644
--- a/src/operators/kernel/arm/mul_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/multiclass_nms_kernel.cpp b/src/operators/kernel/arm/multiclass_nms_kernel.cpp
index 9ed8f1731afe2bab723c66ea1e2e8c5042f6ce28..938f81cf485eb64f408c0fb274eeec673349e306 100644
--- a/src/operators/kernel/arm/multiclass_nms_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/pool_kernel.cpp b/src/operators/kernel/arm/pool_kernel.cpp
index be2189340f480bef80fd00a612cf32e71ea10a1c..60d6f1401876b957649d08889218b88cf1fe5eef 100644
--- a/src/operators/kernel/arm/pool_kernel.cpp
+++ b/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
diff --git a/src/operators/kernel/arm/prelu_kernel.cpp b/src/operators/kernel/arm/prelu_kernel.cpp
index ec8f8e0ab6f95517ccada553ea6bce9837a8b27a..eda944a2ff44a0551ef0342c64e6e955c96b356d 100644
--- a/src/operators/kernel/arm/prelu_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/arm/prior_box_kernel.cpp b/src/operators/kernel/arm/prior_box_kernel.cpp
index 217d4b83cb1156a0e942c5ced5917546250e8bb1..71011fa2112f36d573b5bdc55f1b5bf92318c448 100644
--- a/src/operators/kernel/arm/prior_box_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/relu_kernel.cpp b/src/operators/kernel/arm/relu_kernel.cpp
index 63259a0c303f5e186f9eb90b98f2a8685f8ba5ca..6e04e6013aa8dd5c50dcc22a720b470b08ecd648 100644
--- a/src/operators/kernel/arm/relu_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/reshape_kernel.cpp b/src/operators/kernel/arm/reshape_kernel.cpp
index 5ae8e5e3f945d115215652ded58dc8571868fcd7..235288ae13e2c557e6f7310727f5d8e6e83cedf6 100644
--- a/src/operators/kernel/arm/reshape_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/resize_kernel.cpp b/src/operators/kernel/arm/resize_kernel.cpp
index bef24f1c56f656d25261c9c837e5455244d1ad87..5c0c186554a31454447b1df47a1b7573fd948fb9 100644
--- a/src/operators/kernel/arm/resize_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/arm/scale_kernel.cpp b/src/operators/kernel/arm/scale_kernel.cpp
index c20e70896145958aa91de4f00de7ad2eeba1bb5c..299132ea00f40838249022c45d994e7d88547eaa 100644
--- a/src/operators/kernel/arm/scale_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/arm/sigmoid_kernel.cpp b/src/operators/kernel/arm/sigmoid_kernel.cpp
index 9f5e6a2048d940ddc4592777a773c69d976033bd..7912fd8762b693cd40c632d6b152406ed4b0c568 100644
--- a/src/operators/kernel/arm/sigmoid_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/softmax_kernel.cpp b/src/operators/kernel/arm/softmax_kernel.cpp
index 3ce763be38678319cfc23be83180450e5d3b209c..f86a10601aa3a67300736f2f4c751c05bf41a781 100644
--- a/src/operators/kernel/arm/softmax_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/arm/transpose_kernel.cpp b/src/operators/kernel/arm/transpose_kernel.cpp
index c358edd76e93cee3f8be6086a70c34671c87d383..bb7a881bdc1d2706a25a77833ca38695ede2fec7 100644
--- a/src/operators/kernel/arm/transpose_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/batchnorm_kernel.h b/src/operators/kernel/batchnorm_kernel.h
index 367dd0996c0df5fba7c3570285cf5e2cfd3fac99..beac7399583d074956fa4564fdd9312b2d7985f0 100644
--- a/src/operators/kernel/batchnorm_kernel.h
+++ b/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
diff --git a/src/operators/kernel/box_coder_kernel.h b/src/operators/kernel/box_coder_kernel.h
index 2ad63ecd90a07d955c3e239277ac1bd60f3510bb..58144a87349ed3a6504e0074903594be3aa6fe8f 100644
--- a/src/operators/kernel/box_coder_kernel.h
+++ b/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
diff --git a/src/operators/kernel/central-arm-func/batchnorm_arm_func.h b/src/operators/kernel/central-arm-func/batchnorm_arm_func.h
index cc591035065e4cbbe71ff8f6bd6cbab9c6fe9e79..1723835a6ae54b0c0f0a1e6153bf278efb7d8a12 100644
--- a/src/operators/kernel/central-arm-func/batchnorm_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/box_coder_arm_func.h b/src/operators/kernel/central-arm-func/box_coder_arm_func.h
index 63558dea7b8d64d0fba3daae14e10c565f8feb2d..9cdc22cff0bc52d8ae1ff24d619735accd6dca3e 100644
--- a/src/operators/kernel/central-arm-func/box_coder_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/concat_arm_func.h b/src/operators/kernel/central-arm-func/concat_arm_func.h
index e9926505b33b32ee83a16f882cc0f775797f154a..57a22aafa5e0bc75c1041c379c2229deaa310ffe 100644
--- a/src/operators/kernel/central-arm-func/concat_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/conv_add_arm_func.h b/src/operators/kernel/central-arm-func/conv_add_arm_func.h
index c01a068fb9732b64da4097844736f7484fdfcab9..643ee84529e01aebc33a144b4c7a8181ff39a1c9 100644
--- a/src/operators/kernel/central-arm-func/conv_add_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/conv_add_bn_relu_arm_func.h b/src/operators/kernel/central-arm-func/conv_add_bn_relu_arm_func.h
index 2683b078889059f10b85c1e9fe74374342689418..a7d14fbad1e4b72a8571d13898e55a6cad8bf9a8 100644
--- a/src/operators/kernel/central-arm-func/conv_add_bn_relu_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/conv_add_relu_arm_func.h b/src/operators/kernel/central-arm-func/conv_add_relu_arm_func.h
index 6c619dd2a29ce140c783af0637f51153a1866791..36886b9e2ccfaaa3f557eb7941e294a42b5edb94 100644
--- a/src/operators/kernel/central-arm-func/conv_add_relu_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/conv_arm_func.h b/src/operators/kernel/central-arm-func/conv_arm_func.h
index 33caded3afaaf125bac9108f2fafeda3d3c2049f..a3e21e4b4b702630f7942f2a5171a3401f29a431 100644
--- a/src/operators/kernel/central-arm-func/conv_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/conv_bn_relu_arm_func.h b/src/operators/kernel/central-arm-func/conv_bn_relu_arm_func.h
index 9ec74b1d6ede34cf889c5999df1e902bb9ece4f3..c6300f96e1b999c45538417c7b513068697ad4dd 100644
--- a/src/operators/kernel/central-arm-func/conv_bn_relu_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/conv_transpose_arm_func.h b/src/operators/kernel/central-arm-func/conv_transpose_arm_func.h
index 8c120c60be79b51c2a2e4c8fbccf1d546871f839..343e5f147644cc5bb86c2929d4bd35b44301c4cf 100644
--- a/src/operators/kernel/central-arm-func/conv_transpose_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/depthwise_conv_arm_func.h b/src/operators/kernel/central-arm-func/depthwise_conv_arm_func.h
index 60b09df597f218eefc7f95ba3f342ae0c51c7000..2a1afb3cf6fdbdc0a80cec5558c2b42fec6699f3 100644
--- a/src/operators/kernel/central-arm-func/depthwise_conv_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/dwconv_bn_relu_arm_func.h b/src/operators/kernel/central-arm-func/dwconv_bn_relu_arm_func.h
index 2166fd2a97846dc10733e89a999407afc79e888b..b60bf9b4d6df9d85cc2fbe378a3904c2d13e5e60 100644
--- a/src/operators/kernel/central-arm-func/dwconv_bn_relu_arm_func.h
+++ b/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] &&
diff --git a/src/operators/kernel/central-arm-func/elementwise_add_arm_func.h b/src/operators/kernel/central-arm-func/elementwise_add_arm_func.h
index 91150cfb8c206521fd628c972ab8b35400b5ff53..ace72b6faddb04ee3547f1b2bc01461d8c9f2e98 100644
--- a/src/operators/kernel/central-arm-func/elementwise_add_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/fusion_fc_arm_func.h b/src/operators/kernel/central-arm-func/fusion_fc_arm_func.h
index 4a689dfc18e3b8677faa61b5c90cb46321f3f4c3..425685ba319afee2ce3bba285d69c4edb756a718 100644
--- a/src/operators/kernel/central-arm-func/fusion_fc_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/lrn_arm_func.h b/src/operators/kernel/central-arm-func/lrn_arm_func.h
index 15c279d1605ed7348b766855497411fbe541e2f6..165ad8dd8af1f8a09f16b1737fbaff363cfaf5bd 100644
--- a/src/operators/kernel/central-arm-func/lrn_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/mul_arm_func.h b/src/operators/kernel/central-arm-func/mul_arm_func.h
index d2da67afe1d2eb746971a2443bdb449eb2b66ec4..dd6df54da5a81c2c4d1030103b6bb9811a54246a 100644
--- a/src/operators/kernel/central-arm-func/mul_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/multiclass_nms_arm_func.h b/src/operators/kernel/central-arm-func/multiclass_nms_arm_func.h
index ccc99c13c673fb42d151da6db3372848f5d8bd7c..9de57910540b4c9f7ab807053add9c5af9947ae7 100644
--- a/src/operators/kernel/central-arm-func/multiclass_nms_arm_func.h
+++ b/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();
 
diff --git a/src/operators/kernel/central-arm-func/pool_arm_func.h b/src/operators/kernel/central-arm-func/pool_arm_func.h
index c1beb82da1072d199217d0722eaae6fcb0123490..35fb09f94c49cd91915128260b6426fe0fedf725 100644
--- a/src/operators/kernel/central-arm-func/pool_arm_func.h
+++ b/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 &&
diff --git a/src/operators/kernel/central-arm-func/prior_box_arm_func.h b/src/operators/kernel/central-arm-func/prior_box_arm_func.h
index 892dceb9254ac423d3591a0fc9e9347bc375831b..7129996319aac7c71836d8706eb5c02300e576e6 100644
--- a/src/operators/kernel/central-arm-func/prior_box_arm_func.h
+++ b/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();
 
diff --git a/src/operators/kernel/central-arm-func/relu_arm_func.h b/src/operators/kernel/central-arm-func/relu_arm_func.h
index c1002c1f10e9c2d72d764693384500c4793dda46..d68569c0a5c0730d96a89cd534b2a89c0d3a9bff 100644
--- a/src/operators/kernel/central-arm-func/relu_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/reshape_arm_func.h b/src/operators/kernel/central-arm-func/reshape_arm_func.h
index aa0a3a7c6b5a58938af1c5917c8c2fa0de3f9e90..6e1a29dee6003ec26d58fd61e7445d74eca85edb 100644
--- a/src/operators/kernel/central-arm-func/reshape_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/sigmoid_arm_func.h b/src/operators/kernel/central-arm-func/sigmoid_arm_func.h
index c612c4b092143ef8925f81a6d6fefe9cd9dff25b..1076fa49d555d14da76ff08a67c0943fb9ab115a 100644
--- a/src/operators/kernel/central-arm-func/sigmoid_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/softmax_arm_func.h b/src/operators/kernel/central-arm-func/softmax_arm_func.h
index 896532109409d316346ba4a8d14aaa85a500b007..d311d97984a7207df9075befe71a9806092966e1 100644
--- a/src/operators/kernel/central-arm-func/softmax_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/central-arm-func/transpose_arm_func.h b/src/operators/kernel/central-arm-func/transpose_arm_func.h
index 43a16d0e85ff058764b5a711283e38e15eccae57..1bd2e11a3405abc99c5a33be4ec9b61855f77b08 100644
--- a/src/operators/kernel/central-arm-func/transpose_arm_func.h
+++ b/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();
diff --git a/src/operators/kernel/concat_kernel.h b/src/operators/kernel/concat_kernel.h
index adba64391e3e79569030c95e2d2681a31187f03a..61100bf5f0e9de43bfb6295a0719f1be0954d128 100644
--- a/src/operators/kernel/concat_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_add_bn_kernel.h b/src/operators/kernel/conv_add_bn_kernel.h
index cc11ef1d71f402f32b2da6490877626247884a44..ee73215c4688c3e604de69cda55b05e63844c0b8 100644
--- a/src/operators/kernel/conv_add_bn_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_add_bn_relu_kernel.h b/src/operators/kernel/conv_add_bn_relu_kernel.h
index 73aaf4c900393b9cbee4682fc67147d9ef0853fc..9faaaedcf8d6f825f818ebf5121dc7685185d5d8 100644
--- a/src/operators/kernel/conv_add_bn_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_add_kernel.h b/src/operators/kernel/conv_add_kernel.h
index 5a351f8afcf7b73fb6c56dff48c08d7b5204ca10..360cbb6775168885e9c1a25db1f9ffb9e552324b 100644
--- a/src/operators/kernel/conv_add_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_add_relu_kernel.h b/src/operators/kernel/conv_add_relu_kernel.h
index 931313273d150fa1ad159e7069fbc3812d6e6657..f33b1dc312e1d94be0c23cff55e9e6789a556bc7 100644
--- a/src/operators/kernel/conv_add_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_bn_kernel.h b/src/operators/kernel/conv_bn_kernel.h
index 3bd9f87dcf046313569cfaf7d1142407fa0d493f..f740ca836481c1331ea2e889865b3078d48644a6 100644
--- a/src/operators/kernel/conv_bn_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_bn_relu_kernel.h b/src/operators/kernel/conv_bn_relu_kernel.h
index c9d4df5d8f597deebaf2b53491851b7ce03fc7aa..225976aa5db31096ef691ecefa8b63d4ae3dc277 100644
--- a/src/operators/kernel/conv_bn_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_kernel.h b/src/operators/kernel/conv_kernel.h
index fedbee32a006f263fd3de25064496dad1a23177b..93474adaa97743d1850b53df114ae08f144aebca 100644
--- a/src/operators/kernel/conv_kernel.h
+++ b/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
diff --git a/src/operators/kernel/conv_transpose_kernel.h b/src/operators/kernel/conv_transpose_kernel.h
index bb9751d2462e0f338c8fe90d3de8020c133722e6..9cbd7c8c3bafde8b4f4939e86ceabdd94dbd3bc8 100644
--- a/src/operators/kernel/conv_transpose_kernel.h
+++ b/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
diff --git a/src/operators/kernel/depthwise_conv_kernel.h b/src/operators/kernel/depthwise_conv_kernel.h
index b74a58a649bd9fa27e941e2cd5ea50b30c0218cb..605b81cd6ed4ccd54b1803cf7a603b8f4576982d 100644
--- a/src/operators/kernel/depthwise_conv_kernel.h
+++ b/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
diff --git a/src/operators/kernel/dropout_kernel.h b/src/operators/kernel/dropout_kernel.h
index 3ef6b9dd62d88f012eba3456c676ac0d33bf9e52..b7535095d4fef11ee628aea96a074abcc3562f7f 100644
--- a/src/operators/kernel/dropout_kernel.h
+++ b/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
diff --git a/src/operators/kernel/dwconv_bn_relu_kernel.h b/src/operators/kernel/dwconv_bn_relu_kernel.h
index 91478ae5ecba37472e7e30f774f2c515b6952eee..594c594cb00f8f4ddd8a511f3c992c4efbfcdfc6 100644
--- a/src/operators/kernel/dwconv_bn_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/elementwise_add_kernel.h b/src/operators/kernel/elementwise_add_kernel.h
index 70334c1d3f788f60e974da74133823f82ab05765..67182af2e20e23c40effab6b87eefde1e0ab629d 100644
--- a/src/operators/kernel/elementwise_add_kernel.h
+++ b/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
diff --git a/src/operators/kernel/elementwise_add_relu_kernel.h b/src/operators/kernel/elementwise_add_relu_kernel.h
index 1b3ff25ec983365a7a94cff5b047eba3466fe932..5eda5a0c56c228ad54c888b6faa82ce9417f2dc1 100644
--- a/src/operators/kernel/elementwise_add_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/fc_relu_kernel.h b/src/operators/kernel/fc_relu_kernel.h
index ceff36f8ef49ff996769802b1f39e52e955c45d0..6e9446da37df4ba83db85d416aa87f216816c4a5 100644
--- a/src/operators/kernel/fc_relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/fpga/concat_kernel.cpp b/src/operators/kernel/fpga/concat_kernel.cpp
index c6e04787a58bc437bf0738cf67072426f1cbaa57..8bf7b20a224743f4395cd862d27f1882a847812a 100644
--- a/src/operators/kernel/fpga/concat_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/fpga/conv_add_bn_kernel.cpp b/src/operators/kernel/fpga/conv_add_bn_kernel.cpp
index 152b200cfa88d010bb4c8e8022c01ee3663cc179..7f120fa930334194600103b7310e3e8b50adbe31 100644
--- a/src/operators/kernel/fpga/conv_add_bn_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/fpga/conv_add_bn_relu_kernel.cpp b/src/operators/kernel/fpga/conv_add_bn_relu_kernel.cpp
index caa1e94c6bb9b583efb15e181d46c80f0b66c7ff..7bf5cd3a66c2149079ac213342e8ed7b046cfa99 100644
--- a/src/operators/kernel/fpga/conv_add_bn_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/fpga/conv_add_relu_kernel.cpp b/src/operators/kernel/fpga/conv_add_relu_kernel.cpp
index 33e55773ad0be4f174916f0e5f066b6eeec1d46e..48007e4cb8e90c500d53455d4dd8095827c92831 100644
--- a/src/operators/kernel/fpga/conv_add_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/fpga/conv_bn_kernel.cpp b/src/operators/kernel/fpga/conv_bn_kernel.cpp
index 3ad65a254f95bde431efbd3c5995df6cc2295d3d..49c5f92beb1f12f2b73d31b22c27cd4dd38e115f 100644
--- a/src/operators/kernel/fpga/conv_bn_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/fpga/conv_bn_relu_kernel.cpp b/src/operators/kernel/fpga/conv_bn_relu_kernel.cpp
index 18ef4b4e15e488f01a435d89218992e63873bb14..6bbe1b8763160993da5edb96162e54c8ab688d14 100644
--- a/src/operators/kernel/fpga/conv_bn_relu_kernel.cpp
+++ b/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>;
diff --git a/src/operators/kernel/fpga/dropout_kernel.cpp b/src/operators/kernel/fpga/dropout_kernel.cpp
index bb6ece969d83aba6bb7fe91a3688607df0be8a8b..3a4dd216d481322a9228cfd247bf6f0d0098177e 100644
--- a/src/operators/kernel/fpga/dropout_kernel.cpp
+++ b/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>();
diff --git a/src/operators/kernel/fpga/elementwise_add_relu_kernel.cpp b/src/operators/kernel/fpga/elementwise_add_relu_kernel.cpp
index 88a19beb41f67e5fc9336c8883c8ea75aaa939e0..a71801680a25ae0b21ed5b2a2738c4b593e7c0ac 100644
--- a/src/operators/kernel/fpga/elementwise_add_relu_kernel.cpp
+++ b/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
diff --git a/src/operators/kernel/fpga/fc_relu_kernel.cpp b/src/operators/kernel/fpga/fc_relu_kernel.cpp
index fb6a3e7508bf11f0bba1c3e34c065fa63caa2100..52c8c71537cd3d11eb3bfab43c78a5ad79d0db37 100644
--- a/src/operators/kernel/fpga/fc_relu_kernel.cpp
+++ b/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());
 };
 
diff --git a/src/operators/kernel/fpga/fusion_fc_kernel.cpp b/src/operators/kernel/fpga/fusion_fc_kernel.cpp
index 5479deb6c19cf085dcea03555e4895d4ad98c4e8..199b6b6878ad2c838e3b3d12d8e92a70ea541dad 100644
--- a/src/operators/kernel/fpga/fusion_fc_kernel.cpp
+++ b/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
diff --git a/src/operators/kernel/fpga/pool_kernel.cpp b/src/operators/kernel/fpga/pool_kernel.cpp
index a7ff022c3b8616847c48a71bf94e4018cedcad2e..1565c1519d0048d027fb860f80f313179188c4d5 100644
--- a/src/operators/kernel/fpga/pool_kernel.cpp
+++ b/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
diff --git a/src/operators/kernel/fpga/softmax_kernel.cpp b/src/operators/kernel/fpga/softmax_kernel.cpp
index d58ab0f751eeb584f286a0920d08e9473be38402..1c609431310b1399e0427326d5237368e649c14c 100644
--- a/src/operators/kernel/fpga/softmax_kernel.cpp
+++ b/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);
 }
 
diff --git a/src/operators/kernel/fusion_fc_kernel.h b/src/operators/kernel/fusion_fc_kernel.h
index 0e31134ba5a18405a5855db1e85b3885608c4071..39cfd898a203e742168a775ec892e562bd19f5db 100644
--- a/src/operators/kernel/fusion_fc_kernel.h
+++ b/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
diff --git a/src/operators/kernel/im2sequence_kernel.h b/src/operators/kernel/im2sequence_kernel.h
index aa798fd6af5592a062de207714dc9fee2afb93df..df93ea5abacda1a5291caa53dc5dae7ea2b5d710 100644
--- a/src/operators/kernel/im2sequence_kernel.h
+++ b/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
diff --git a/src/operators/kernel/lrn_kernel.h b/src/operators/kernel/lrn_kernel.h
index ee19d6e40ee0b5b66f62ce6535370a81c28950af..164178f1dcc0ee2523fc9c5fdc4736c14a3e55ce 100644
--- a/src/operators/kernel/lrn_kernel.h
+++ b/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
diff --git a/src/operators/kernel/mali/batchnorm_kernel.cpp b/src/operators/kernel/mali/batchnorm_kernel.cpp
index ad648d615cd8f9134b212d484d7174c95e027551..50f6ef5f566347c089869c30b8f7534a4f8b6779 100755
--- a/src/operators/kernel/mali/batchnorm_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/concat_kernel.cpp b/src/operators/kernel/mali/concat_kernel.cpp
old mode 100755
new mode 100644
index aaa586b6d977bfca96e596261ec090637cf87207..267c0101a8f66de3d508dbe5795c87ee5027a288
--- a/src/operators/kernel/mali/concat_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/conv_add_kernel.cpp b/src/operators/kernel/mali/conv_add_kernel.cpp
index 318db016d12f36981c07627139bcc49d07162d52..74cace00dd2dead7a5d9ddfc76e2d48c67cccf89 100644
--- a/src/operators/kernel/mali/conv_add_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/conv_kernel.cpp b/src/operators/kernel/mali/conv_kernel.cpp
old mode 100755
new mode 100644
index c548977ebaa34fabc1b1fe54d6db9690bcb424f1..7852e64990e5a2cd6f3d7e803e71c23c55aa7a27
--- a/src/operators/kernel/mali/conv_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/elementwise_add_kernel.cpp b/src/operators/kernel/mali/elementwise_add_kernel.cpp
old mode 100755
new mode 100644
index 9748bbbb5454f10ad9ea83e37d599fb1c6cdb53e..5596476e1bb33ecc2b3122bf237090b099307156
--- a/src/operators/kernel/mali/elementwise_add_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/mali/fushion_fc_kernel.cpp b/src/operators/kernel/mali/fushion_fc_kernel.cpp
index 44a7ce2af62a1d27aff8181f6742bebda1d6d066..c3197f38c6c6ee1a4f4f684c824a9a9e43d69d4f 100755
--- a/src/operators/kernel/mali/fushion_fc_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/mali/lrn_kernel.cpp b/src/operators/kernel/mali/lrn_kernel.cpp
index 4fb5fca8695dccc45c7169d8572618965b3d84a3..fc088f735c538bedc4d5c79593aa31c48acc4fc6 100644
--- a/src/operators/kernel/mali/lrn_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/mul_kernel.cpp b/src/operators/kernel/mali/mul_kernel.cpp
old mode 100755
new mode 100644
index 3a9ec4ebb319d9e521240ad987a49549c22c1ff2..a9e54dad2b51c595be4f68df3916a4803047617e
--- a/src/operators/kernel/mali/mul_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/mali/pool_kernel.cpp b/src/operators/kernel/mali/pool_kernel.cpp
old mode 100755
new mode 100644
index 1f49391341d69a0690352c69c9c208550f8e1c24..33b3bd7017739144a519bfb1be247b4751883779
--- a/src/operators/kernel/mali/pool_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/relu_kernel.cpp b/src/operators/kernel/mali/relu_kernel.cpp
old mode 100755
new mode 100644
index 1a8c0f88543199e7a863cc44b5b0a6be3bc6212d..10b270800dee1a0ad8176da1f788100d29b60173
--- a/src/operators/kernel/mali/relu_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mali/reshape_kernel.cpp b/src/operators/kernel/mali/reshape_kernel.cpp
old mode 100755
new mode 100644
index 57837a677033590e92a307bd69a77c076c5ba805..69c077e252162017cb477a000b5f17f5a968fc10
--- a/src/operators/kernel/mali/reshape_kernel.cpp
+++ b/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();
diff --git a/src/operators/kernel/mali/softmax_kernel.cpp b/src/operators/kernel/mali/softmax_kernel.cpp
old mode 100755
new mode 100644
index 37d2f2b6b1656602e5acfecd3ac79733f570844d..d4f25c96cc47d7baa394645d4e0c84e0e3f7ad29
--- a/src/operators/kernel/mali/softmax_kernel.cpp
+++ b/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());
diff --git a/src/operators/kernel/mul_kernel.h b/src/operators/kernel/mul_kernel.h
index f7dcb738b38448fe38eb60dcbbd4a2abda7a858a..e441de4d4495b736aec248c0ef85191b32bfcbf9 100644
--- a/src/operators/kernel/mul_kernel.h
+++ b/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
diff --git a/src/operators/kernel/multiclass_nms_kernel.h b/src/operators/kernel/multiclass_nms_kernel.h
index 9bd00b874a1140373decca582f793febf0e941ec..b1b20ddd81b395ea94ae62b1abf2fe861d9257db 100644
--- a/src/operators/kernel/multiclass_nms_kernel.h
+++ b/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
diff --git a/src/operators/kernel/pool_kernel.h b/src/operators/kernel/pool_kernel.h
index fd9faa3d5a508084924e080f5c5ed7e7b454b5f2..2be254444cc410fb95a94125cccb224ca9505545 100644
--- a/src/operators/kernel/pool_kernel.h
+++ b/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
diff --git a/src/operators/kernel/prelu_kernel.h b/src/operators/kernel/prelu_kernel.h
index 15696174377f04ad9a62366e03ded1f2cdcdee9e..f6c7c3ac7f139cf7eafe8843ef48e53c90292082 100644
--- a/src/operators/kernel/prelu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/prior_box_kernel.h b/src/operators/kernel/prior_box_kernel.h
index d169a01d7f45f7dbdcc02be0e1e71690b8550af8..5640375483d42d52965986dab6795254bbf4b908 100644
--- a/src/operators/kernel/prior_box_kernel.h
+++ b/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
diff --git a/src/operators/kernel/relu_kernel.h b/src/operators/kernel/relu_kernel.h
index 64016656b20b0fdb08f1342f7853e2e727a6bb81..b0c32791d626f14b0840ce1c8f3f12f02b403d97 100644
--- a/src/operators/kernel/relu_kernel.h
+++ b/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
diff --git a/src/operators/kernel/reshape_kernel.h b/src/operators/kernel/reshape_kernel.h
index 47eba531b9f36d83d44588d9cdfb162519c24180..73eb63f797f34ec4eb2baec8c4ab79fafb06f0e2 100644
--- a/src/operators/kernel/reshape_kernel.h
+++ b/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
diff --git a/src/operators/kernel/resize_kernel.h b/src/operators/kernel/resize_kernel.h
index f1b9e64fc00863e83caa1c44606cdb3ec9f44817..4c06429858b9575ffc061c000e4a9343fa7eee26 100644
--- a/src/operators/kernel/resize_kernel.h
+++ b/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
diff --git a/src/operators/kernel/scale_kernel.h b/src/operators/kernel/scale_kernel.h
index 98ac71d0bbad86f595171ad7ac5b2a1cdf5908fa..2da92d8d3c8b0d7867e7e6e628a04a853dd69464 100644
--- a/src/operators/kernel/scale_kernel.h
+++ b/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
diff --git a/src/operators/kernel/sigmoid_kernel.h b/src/operators/kernel/sigmoid_kernel.h
index fc3eb5e1bf158c541b2f00d9e57ddd4699344006..e8cfe6cad9ce2f25b9f38e1784ded9ea0741ff9a 100644
--- a/src/operators/kernel/sigmoid_kernel.h
+++ b/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
diff --git a/src/operators/kernel/slice_kernel.h b/src/operators/kernel/slice_kernel.h
index fd3b8dc767076c5244509f6015c42bee87df100b..17f7fe4a9ebf5b78fc92c41abd4756a7bc6bff45 100644
--- a/src/operators/kernel/slice_kernel.h
+++ b/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
diff --git a/src/operators/kernel/softmax_kernel.h b/src/operators/kernel/softmax_kernel.h
index a500d9c81cce96b0f1db6d45981ad9aa02ea7c0b..67bd9167e8c717355fc326d3025cde410ce66010 100644
--- a/src/operators/kernel/softmax_kernel.h
+++ b/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
diff --git a/src/operators/kernel/transpose_kernel.h b/src/operators/kernel/transpose_kernel.h
index f1a21ebbb28c2acdb905ce9f09c28f0d47e17294..56c41fd221e080a4db3b34fbd4ab208c9986c2a8 100644
--- a/src/operators/kernel/transpose_kernel.h
+++ b/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
diff --git a/src/operators/lrn_op.h b/src/operators/lrn_op.h
index 0d756a14f4d935fd59ac2bfc7c811c674b1587fe..6c609c7654cca022f473dba0aad1f4214a4e43e3 100644
--- a/src/operators/lrn_op.h
+++ b/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;
 
diff --git a/src/operators/math/pooling.cpp b/src/operators/math/pooling.cpp
index 24db2e272e3124a223c22c6f687d868d42126f6b..f5bcdf7fdb6b9245eda7d3557b293395bce23b24 100644
--- a/src/operators/math/pooling.cpp
+++ b/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
diff --git a/src/operators/mul_op.h b/src/operators/mul_op.h
index 64b811b01091418c9febdfb8d03bacd77421dcf5..127048efbacf2da87de9371cd8e54875f8554d61 100644
--- a/src/operators/mul_op.h
+++ b/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;
 
diff --git a/src/operators/multiclass_nms_op.h b/src/operators/multiclass_nms_op.h
index 425f7d33e35e0864b5f5a7739dbfa18bc8eb0c30..b40ef5ee009f6c16c685479ffcf58186958bb4cc 100644
--- a/src/operators/multiclass_nms_op.h
+++ b/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;
 
diff --git a/src/operators/op_param.cpp b/src/operators/op_param.cpp
index 4ad7685731bdb44794c235c639be7eed4a0c812b..4d1689911686198612eb4df4dfe8f99450ba503d 100644
--- a/src/operators/op_param.cpp
+++ b/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: "
diff --git a/src/operators/op_param.h b/src/operators/op_param.h
index 06da537e419f3a54ffc9986b12274f9853f12774..6fe851b31d2c50a6c37388905129a414b00b5a98 100644
--- a/src/operators/op_param.h
+++ b/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_;
diff --git a/src/operators/pool_op.h b/src/operators/pool_op.h
index bc823e70c5a8e7b229a0101726316d1e825b7b54..4f76fb8f800dea43432b48562cca563505a1af76 100644
--- a/src/operators/pool_op.h
+++ b/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;
 
diff --git a/src/operators/prelu_op.h b/src/operators/prelu_op.h
index 330d1d4a567b10bddf5879b4897a945350a34ac5..af33476b7298a5728a6ef944506d55f422a2fa8c 100644
--- a/src/operators/prelu_op.h
+++ b/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;
 
diff --git a/src/operators/prior_box_op.h b/src/operators/prior_box_op.h
index 6fcaa07c74f0e005fd5b91ae04ec7219e0394064..00fc8e039c9958e4b43653d6360c0f54c78648a1 100644
--- a/src/operators/prior_box_op.h
+++ b/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;
 
diff --git a/src/operators/relu_op.h b/src/operators/relu_op.h
index 204ec3d29c147d0d52b9b05d16de6807211a5e57..0364dd7f8ec4b3861200380597e18ede0819e8b6 100644
--- a/src/operators/relu_op.h
+++ b/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;
 
diff --git a/src/operators/reshape_op.h b/src/operators/reshape_op.h
index da2328ec3570359ccdb45ce1511c02f322498aa1..9284e94f346ed0f225d6dabe16077b1fb2034c64 100644
--- a/src/operators/reshape_op.h
+++ b/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;
 
diff --git a/src/operators/resize_op.h b/src/operators/resize_op.h
index 6cab048dea350d668c92fda56f6b6b197c38093d..c0b38bb1cf4048af4b07d05f28a88a5ac8056ea3 100644
--- a/src/operators/resize_op.h
+++ b/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;
 
diff --git a/src/operators/scale_op.h b/src/operators/scale_op.h
index e6a900d646a53cb8aa2a0ab624a31781f420db06..4c5f5e620f25bef88533e80cdd78b243fef9bc70 100644
--- a/src/operators/scale_op.h
+++ b/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;
 
diff --git a/src/operators/sigmoid_op.h b/src/operators/sigmoid_op.h
index bffef7880b2bb6057f5d489eaac6dea7a3fb3ab5..533ea587958e8766b1469c73b909cfa2fcb60696 100644
--- a/src/operators/sigmoid_op.h
+++ b/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;
diff --git a/src/operators/slice_op.h b/src/operators/slice_op.h
index 4c9d074b24c1aaf1bf28c862f3731ca130f3c462..6bcb6fa0b9e88cefb3c88dfc096e1073ad261c1b 100644
--- a/src/operators/slice_op.h
+++ b/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;
 
diff --git a/src/operators/softmax_op.h b/src/operators/softmax_op.h
index bacae23b522daf1cc689a2d7af6b14cd2bc794bb..579a2ed605cb3f3c8c4a3d0c2f1ccc7bd9595fc2 100644
--- a/src/operators/softmax_op.h
+++ b/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;
diff --git a/src/operators/transpose_op.h b/src/operators/transpose_op.h
index 25cf07c4c9253736d513505e5f8eba6147f3740c..b96ce4e17ca4b0d0e321cefb3175b973cd7df307 100644
--- a/src/operators/transpose_op.h
+++ b/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;
 };