diff --git a/paddle/fluid/operators/detection/target_assign_op.cc b/paddle/fluid/operators/detection/target_assign_op.cc
index 155ec31fa92d712c8308c8353c7bcdc8159d3517..437b46c459ff3ece3751b0c1c4e6d17a8d76e7e9 100644
--- a/paddle/fluid/operators/detection/target_assign_op.cc
+++ b/paddle/fluid/operators/detection/target_assign_op.cc
@@ -186,6 +186,6 @@ REGISTER_OPERATOR(
     ops::TargetAssignOpMaker,
     paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
     paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>);
-REGISTER_OP_CPU_KERNEL(target_assign,
-                       ops::TargetAssignKernel<phi::CPUContext, int, float>,
-                       ops::TargetAssignKernel<phi::CPUContext, float, float>);
+
+PD_REGISTER_STRUCT_KERNEL(
+    target_assign, CPU, ALL_LAYOUT, ops::TargetAssignKernel, int, float) {}
diff --git a/paddle/fluid/operators/detection/target_assign_op.cu b/paddle/fluid/operators/detection/target_assign_op.cu
index 337f55a3ca81fa7a680efb0b708a4790ffb5ad26..951fcdbafae8e1c0c870c491debbd4d854e692ea 100644
--- a/paddle/fluid/operators/detection/target_assign_op.cu
+++ b/paddle/fluid/operators/detection/target_assign_op.cu
@@ -65,6 +65,6 @@ template struct NegTargetAssignFunctor<phi::GPUContext, float, float>;
 }  // namespace paddle
 
 namespace ops = paddle::operators;
-REGISTER_OP_CUDA_KERNEL(target_assign,
-                        ops::TargetAssignKernel<phi::GPUContext, int, float>,
-                        ops::TargetAssignKernel<phi::GPUContext, float, float>);
+
+PD_REGISTER_STRUCT_KERNEL(
+    target_assign, GPU, ALL_LAYOUT, ops::TargetAssignKernel, int, float) {}
diff --git a/paddle/fluid/operators/detection/target_assign_op.h b/paddle/fluid/operators/detection/target_assign_op.h
index 3319dffd226148913b99f60c1ba2ca6847e26edd..484bd8454bae9d8a2dcf5541eddd2d9cc157c747 100644
--- a/paddle/fluid/operators/detection/target_assign_op.h
+++ b/paddle/fluid/operators/detection/target_assign_op.h
@@ -92,7 +92,7 @@ struct NegTargetAssignFunctor {
                   WT* out_wt) const;
 };
 
-template <typename DeviceContext, typename T, typename WT>
+template <typename T, typename DeviceContext, typename WT = float>
 class TargetAssignKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
diff --git a/paddle/fluid/operators/squared_l2_distance_op.cc b/paddle/fluid/operators/squared_l2_distance_op.cc
index f1ed2d3ee6813585e637dc9fefe25a5d7231a5b5..0f2f727dd9135d6e835c99b09ffd6eae684c0e11 100644
--- a/paddle/fluid/operators/squared_l2_distance_op.cc
+++ b/paddle/fluid/operators/squared_l2_distance_op.cc
@@ -221,8 +221,12 @@ REGISTER_OPERATOR(
 REGISTER_OPERATOR(squared_l2_distance_grad,
                   ops::SquaredL2DistanceGradOp,
                   ops::SquaredL2DistanceGradOpNoBufferVarsInferer);
-REGISTER_OP_CPU_KERNEL(squared_l2_distance,
-                       ops::SquaredL2DistanceKernel<phi::CPUContext, float>);
-REGISTER_OP_CPU_KERNEL(
-    squared_l2_distance_grad,
-    ops::SquaredL2DistanceGradKernel<phi::CPUContext, float>);
+
+PD_REGISTER_STRUCT_KERNEL(
+    squared_l2_distance, CPU, ALL_LAYOUT, ops::SquaredL2DistanceKernel, float) {
+}
+PD_REGISTER_STRUCT_KERNEL(squared_l2_distance_grad,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::SquaredL2DistanceGradKernel,
+                          float) {}
diff --git a/paddle/fluid/operators/squared_l2_distance_op.cu b/paddle/fluid/operators/squared_l2_distance_op.cu
index c10cbfb42f1352f2a44c9dd812542acd98ae94bd..4411df4d9ab7f070859a181ae4c3e4dd60e4e3df 100644
--- a/paddle/fluid/operators/squared_l2_distance_op.cu
+++ b/paddle/fluid/operators/squared_l2_distance_op.cu
@@ -14,8 +14,11 @@ limitations under the License. */
 #include "paddle/fluid/operators/squared_l2_distance_op.h"
 
 namespace ops = paddle::operators;
-REGISTER_OP_CUDA_KERNEL(squared_l2_distance,
-                        ops::SquaredL2DistanceKernel<phi::GPUContext, float>);
-REGISTER_OP_CUDA_KERNEL(
-    squared_l2_distance_grad,
-    ops::SquaredL2DistanceGradKernel<phi::GPUContext, float>);
+PD_REGISTER_STRUCT_KERNEL(
+    squared_l2_distance, GPU, ALL_LAYOUT, ops::SquaredL2DistanceKernel, float) {
+}
+PD_REGISTER_STRUCT_KERNEL(squared_l2_distance_grad,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::SquaredL2DistanceGradKernel,
+                          float) {}
diff --git a/paddle/fluid/operators/squared_l2_distance_op.h b/paddle/fluid/operators/squared_l2_distance_op.h
index f0838c4fad2de89e9933633c864c16f45d8dbc80..18039835c55c3fa9aae964d3ef0dab119b57ff24 100644
--- a/paddle/fluid/operators/squared_l2_distance_op.h
+++ b/paddle/fluid/operators/squared_l2_distance_op.h
@@ -19,7 +19,7 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class SquaredL2DistanceKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
@@ -60,7 +60,7 @@ class SquaredL2DistanceKernel : public framework::OpKernel<T> {
   }
 };
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class SquaredL2DistanceGradKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
diff --git a/paddle/fluid/operators/tdm_child_op.cc b/paddle/fluid/operators/tdm_child_op.cc
index 0ec2c1e85bf838ab7eaf2a265033d809305d59ad..266bbb9657db2093ad3d994551c2d02541b21137 100644
--- a/paddle/fluid/operators/tdm_child_op.cc
+++ b/paddle/fluid/operators/tdm_child_op.cc
@@ -119,9 +119,12 @@ REGISTER_OPERATOR(
     ops::TDMChildOpMaker,
     paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
     paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>);
-REGISTER_OP_CPU_KERNEL(
-    tdm_child,
-    ops::TDMChildKernel<paddle::platform::CPUPlace, float>,
-    ops::TDMChildKernel<paddle::platform::CPUPlace, double>,
-    ops::TDMChildKernel<paddle::platform::CPUPlace, int>,
-    ops::TDMChildKernel<paddle::platform::CPUPlace, int64_t>);
+
+PD_REGISTER_STRUCT_KERNEL(tdm_child,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::TDMChildKernel,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/tdm_child_op.h b/paddle/fluid/operators/tdm_child_op.h
index 0064567887eb0f7f5135019e9dc29f2c9fc8705d..afd8e667b82bef29e885285f13f4895f5266c2ec 100644
--- a/paddle/fluid/operators/tdm_child_op.h
+++ b/paddle/fluid/operators/tdm_child_op.h
@@ -105,7 +105,7 @@ void TDMChildInner(const framework::ExecutionContext &context,
   memcpy(leaf_mask_data, &item_mask_vec[0], sizeof(OutT) * output_nums);
 }
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class TDMChildKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &ctx) const override {
diff --git a/paddle/fluid/operators/tdm_sampler_op.cc b/paddle/fluid/operators/tdm_sampler_op.cc
index 66e9728d88f7333dcb897b9979fa44aaf4a50162..cb667342f0784e3b6529165c11fbf664ea8309b6 100644
--- a/paddle/fluid/operators/tdm_sampler_op.cc
+++ b/paddle/fluid/operators/tdm_sampler_op.cc
@@ -136,9 +136,12 @@ REGISTER_OPERATOR(
     ops::TDMSamplerOpMaker,
     paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
     paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>);
-REGISTER_OP_CPU_KERNEL(
-    tdm_sampler,
-    ops::TDMSamplerKernel<paddle::platform::CPUPlace, float>,
-    ops::TDMSamplerKernel<paddle::platform::CPUPlace, double>,
-    ops::TDMSamplerKernel<paddle::platform::CPUPlace, int>,
-    ops::TDMSamplerKernel<paddle::platform::CPUPlace, int64_t>);
+
+PD_REGISTER_STRUCT_KERNEL(tdm_sampler,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::TDMSamplerKernel,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/tdm_sampler_op.h b/paddle/fluid/operators/tdm_sampler_op.h
index e7fd36b8dcb68e91aa16a7550dd7b3f067dad371..4baff820784713afcfbd9ab28d74619feaa233b5 100644
--- a/paddle/fluid/operators/tdm_sampler_op.h
+++ b/paddle/fluid/operators/tdm_sampler_op.h
@@ -251,7 +251,7 @@ void TDMSamplerInner(const framework::ExecutionContext &context,
   }
 }
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class TDMSamplerKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &context) const override {
diff --git a/paddle/fluid/operators/teacher_student_sigmoid_loss_op.cc b/paddle/fluid/operators/teacher_student_sigmoid_loss_op.cc
index fdb78f9da326e359e655f6d75c12603873c57baa..ad54a49f820f924536cc1f0a94b4487273f38632 100644
--- a/paddle/fluid/operators/teacher_student_sigmoid_loss_op.cc
+++ b/paddle/fluid/operators/teacher_student_sigmoid_loss_op.cc
@@ -249,10 +249,15 @@ REGISTER_OPERATOR(
 REGISTER_OPERATOR(teacher_student_sigmoid_loss_grad,
                   ops::TeacherStudentSigmoidLossGradientOp);
 
-REGISTER_OP_CPU_KERNEL(teacher_student_sigmoid_loss,
-                       ops::TeacherStudentSigmoidLossOpKernel<float>,
-                       ops::TeacherStudentSigmoidLossOpKernel<double>);
-
-REGISTER_OP_CPU_KERNEL(teacher_student_sigmoid_loss_grad,
-                       ops::TeacherStudentSigmoidLossGradOpKernel<float>,
-                       ops::TeacherStudentSigmoidLossGradOpKernel<double>);
+PD_REGISTER_STRUCT_KERNEL(teacher_student_sigmoid_loss,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::TeacherStudentSigmoidLossOpKernel,
+                          float,
+                          double) {}
+PD_REGISTER_STRUCT_KERNEL(teacher_student_sigmoid_loss_grad,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::TeacherStudentSigmoidLossGradOpKernel,
+                          float,
+                          double) {}
diff --git a/paddle/fluid/operators/teacher_student_sigmoid_loss_op.h b/paddle/fluid/operators/teacher_student_sigmoid_loss_op.h
index 133d9656284f39d1b554e41c9238047a719ec270..7ccb9438d418868fbbdf9ef27c644efe5d4c71c0 100644
--- a/paddle/fluid/operators/teacher_student_sigmoid_loss_op.h
+++ b/paddle/fluid/operators/teacher_student_sigmoid_loss_op.h
@@ -19,7 +19,7 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class TeacherStudentSigmoidLossOpKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
@@ -63,7 +63,7 @@ class TeacherStudentSigmoidLossOpKernel : public framework::OpKernel<T> {
   }
 };
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class TeacherStudentSigmoidLossGradOpKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
diff --git a/paddle/fluid/operators/temporal_shift_op.cu b/paddle/fluid/operators/temporal_shift_op.cu
index d2583aeb143ec65ab63c56c5f1e958310d1d824b..68f8153e88e600a7ff554ecf266ac19917214829 100644
--- a/paddle/fluid/operators/temporal_shift_op.cu
+++ b/paddle/fluid/operators/temporal_shift_op.cu
@@ -152,7 +152,7 @@ __global__ void KeTemporalShiftBwNHWC(const T* output_grad,
   }
 }
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class TemporalShiftOpCUDAKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
@@ -208,7 +208,7 @@ class TemporalShiftOpCUDAKernel : public framework::OpKernel<T> {
   }
 };
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class TemporalShiftGradOpCUDAKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
@@ -267,13 +267,19 @@ class TemporalShiftGradOpCUDAKernel : public framework::OpKernel<T> {
 }  // namespace paddle
 
 namespace ops = paddle::operators;
-REGISTER_OP_CUDA_KERNEL(
-    temporal_shift,
-    ops::TemporalShiftOpCUDAKernel<float>,
-    ops::TemporalShiftOpCUDAKernel<double>,
-    ops::TemporalShiftOpCUDAKernel<paddle::platform::float16>);
-REGISTER_OP_CUDA_KERNEL(
-    temporal_shift_grad,
-    ops::TemporalShiftGradOpCUDAKernel<float>,
-    ops::TemporalShiftGradOpCUDAKernel<double>,
-    ops::TemporalShiftGradOpCUDAKernel<paddle::platform::float16>);
+namespace plat = paddle::platform;
+
+PD_REGISTER_STRUCT_KERNEL(temporal_shift,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::TemporalShiftOpCUDAKernel,
+                          float,
+                          double,
+                          plat::float16) {}
+PD_REGISTER_STRUCT_KERNEL(temporal_shift_grad,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::TemporalShiftGradOpCUDAKernel,
+                          float,
+                          double,
+                          plat::float16) {}
diff --git a/paddle/fluid/operators/tree_conv_op.cc b/paddle/fluid/operators/tree_conv_op.cc
index 0e78aa20faa323ebae8c6fabcf7b38da15a81481..7265d966b9e2a4a00f047290b04949456294c174 100644
--- a/paddle/fluid/operators/tree_conv_op.cc
+++ b/paddle/fluid/operators/tree_conv_op.cc
@@ -234,10 +234,7 @@ REGISTER_OPERATOR(tree_conv,
 
 REGISTER_OPERATOR(tree_conv_grad, ops::TreeConvGradOp);
 
-REGISTER_OP_CPU_KERNEL(tree_conv,
-                       ops::TreeConvKernel<phi::CPUContext, float>,
-                       ops::TreeConvKernel<phi::CPUContext, double>);
-
-REGISTER_OP_CPU_KERNEL(tree_conv_grad,
-                       ops::TreeConvGradKernel<phi::CPUContext, float>,
-                       ops::TreeConvGradKernel<phi::CPUContext, double>);
+PD_REGISTER_STRUCT_KERNEL(
+    tree_conv, CPU, ALL_LAYOUT, ops::TreeConvKernel, float, double) {}
+PD_REGISTER_STRUCT_KERNEL(
+    tree_conv_grad, CPU, ALL_LAYOUT, ops::TreeConvGradKernel, float, double) {}
diff --git a/paddle/fluid/operators/tree_conv_op.cu b/paddle/fluid/operators/tree_conv_op.cu
index 1e4ca7bb838cef4ce90cda0d3afc067af62a56f9..1bfcb94013c2bd4e60f3fafea96cbd5e32a0a003 100644
--- a/paddle/fluid/operators/tree_conv_op.cu
+++ b/paddle/fluid/operators/tree_conv_op.cu
@@ -15,9 +15,8 @@
 #include "paddle/fluid/operators/tree_conv_op.h"
 
 namespace ops = paddle::operators;
-REGISTER_OP_CUDA_KERNEL(tree_conv,
-                        ops::TreeConvKernel<phi::GPUContext, float>,
-                        ops::TreeConvKernel<phi::GPUContext, double>);
-REGISTER_OP_CUDA_KERNEL(tree_conv_grad,
-                        ops::TreeConvGradKernel<phi::GPUContext, float>,
-                        ops::TreeConvGradKernel<phi::GPUContext, double>);
+
+PD_REGISTER_STRUCT_KERNEL(
+    tree_conv, GPU, ALL_LAYOUT, ops::TreeConvKernel, float, double) {}
+PD_REGISTER_STRUCT_KERNEL(
+    tree_conv_grad, GPU, ALL_LAYOUT, ops::TreeConvGradKernel, float, double) {}
diff --git a/paddle/fluid/operators/tree_conv_op.h b/paddle/fluid/operators/tree_conv_op.h
index cab0796a71019c32633a95177b6e2d78f351e431..18fd5bea29d308f351aba8fa586ef04b07414cfe 100644
--- a/paddle/fluid/operators/tree_conv_op.h
+++ b/paddle/fluid/operators/tree_conv_op.h
@@ -23,7 +23,7 @@
 namespace paddle {
 namespace operators {
 using DDim = framework::DDim;
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class TreeConvKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &ctx) const override {
@@ -73,7 +73,7 @@ class TreeConvKernel : public framework::OpKernel<T> {
     }
   }
 };
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class TreeConvGradKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &ctx) const override {
diff --git a/paddle/fluid/operators/unique_with_counts_op.cc b/paddle/fluid/operators/unique_with_counts_op.cc
index 3726fd978bd935ffe4ebcf69dda909f762e532d1..5272158805d7112ec6675e07db004bae825ac3bc 100644
--- a/paddle/fluid/operators/unique_with_counts_op.cc
+++ b/paddle/fluid/operators/unique_with_counts_op.cc
@@ -75,8 +75,11 @@ namespace ops = paddle::operators;
 REGISTER_OP_WITHOUT_GRADIENT(unique_with_counts,
                              ops::UniqueWithCountsOp,
                              ops::UniqueWithCountsOpMaker);
-REGISTER_OP_CPU_KERNEL(unique_with_counts,
-                       ops::UniqueWithCountsKernel<float>,
-                       ops::UniqueWithCountsKernel<double>,
-                       ops::UniqueWithCountsKernel<int32_t>,
-                       ops::UniqueWithCountsKernel<int64_t>);
+PD_REGISTER_STRUCT_KERNEL(unique_with_counts,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::UniqueWithCountsKernel,
+                          float,
+                          double,
+                          int32_t,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/unique_with_counts_op.h b/paddle/fluid/operators/unique_with_counts_op.h
index eb3cc2d4731df10dcb5a22f78cec900d18f75c1a..4b1fef5e22447d7cd6062b3e5c3876e8dfae1466 100644
--- a/paddle/fluid/operators/unique_with_counts_op.h
+++ b/paddle/fluid/operators/unique_with_counts_op.h
@@ -25,7 +25,7 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class UniqueWithCountsKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
diff --git a/paddle/fluid/operators/var_conv_2d_op.cc b/paddle/fluid/operators/var_conv_2d_op.cc
index f60190f00cb5561ac2a2ad2cd1b36735bb382c66..112aefbe7f22ada2f1c903ad7a04ba03b1c0fbf3 100644
--- a/paddle/fluid/operators/var_conv_2d_op.cc
+++ b/paddle/fluid/operators/var_conv_2d_op.cc
@@ -176,7 +176,7 @@ void VarConv2dOP::InferShape(framework::InferShapeContext* ctx) const {
   }
 }
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class CPUVarConv2dOPKernel : public framework::OpKernel<T> {
  public:
   void Im2Col(const framework::ExecutionContext& ctx,
@@ -392,7 +392,7 @@ void VarConv2dOpGrad::InferShape(framework::InferShapeContext* ctx) const {
   }
 }
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class CPUVarConv2dOPGradKernel : public framework::OpKernel<T> {
  public:
   void Im2ColGrad(const framework::ExecutionContext& ctx, T* top_diff) const {
@@ -532,11 +532,7 @@ REGISTER_OPERATOR(var_conv_2d,
                   ops::VarConv2dGradMaker<paddle::imperative::OpBase>);
 REGISTER_OPERATOR(var_conv_2d_grad, ops::VarConv2dOpGrad);
 
-REGISTER_OP_CPU_KERNEL(var_conv_2d,
-                       ops::CPUVarConv2dOPKernel<phi::CPUContext, float>);
-//     ops::CPUVarConv2dOPKernel<phi::CPUContext,
-//                                       double>
-REGISTER_OP_CPU_KERNEL(var_conv_2d_grad,
-                       ops::CPUVarConv2dOPGradKernel<phi::CPUContext, float>);
-//     ops::CPUVarConv2dOPGradKernel<phi::CPUContext,
-//                                           double>
+PD_REGISTER_STRUCT_KERNEL(
+    var_conv_2d, CPU, ALL_LAYOUT, ops::CPUVarConv2dOPKernel, float) {}
+PD_REGISTER_STRUCT_KERNEL(
+    var_conv_2d_grad, CPU, ALL_LAYOUT, ops::CPUVarConv2dOPGradKernel, float) {}