diff --git a/paddle/fluid/operators/custom_device_common_op_registry.cc b/paddle/fluid/operators/custom_device_common_op_registry.cc
index e4c927183f687e10b344f41d13bce52bf56b8a3d..aa5bd9dc07b61e21f1db8100f04fe102f99db28e 100644
--- a/paddle/fluid/operators/custom_device_common_op_registry.cc
+++ b/paddle/fluid/operators/custom_device_common_op_registry.cc
@@ -521,15 +521,15 @@ void RegisterCustomDeviceCommonKernel(const std::string& dev_type) {
       load_combine,
       device_type,
       paddle::operators::
-          LoadCombineOpKernel<paddle::platform::CustomDeviceContext, float>,
+          LoadCombineOpKernel<float, paddle::platform::CustomDeviceContext>,
       paddle::operators::
-          LoadCombineOpKernel<paddle::platform::CustomDeviceContext, double>,
+          LoadCombineOpKernel<double, paddle::platform::CustomDeviceContext>,
       paddle::operators::
-          LoadCombineOpKernel<paddle::platform::CustomDeviceContext, int>,
+          LoadCombineOpKernel<int, paddle::platform::CustomDeviceContext>,
       paddle::operators::
-          LoadCombineOpKernel<paddle::platform::CustomDeviceContext, int8_t>,
+          LoadCombineOpKernel<int8_t, paddle::platform::CustomDeviceContext>,
       paddle::operators::
-          LoadCombineOpKernel<paddle::platform::CustomDeviceContext, int64_t>);
+          LoadCombineOpKernel<int64_t, paddle::platform::CustomDeviceContext>);
   REGISTER_OP_CUSTOM_DEVICE_KERNEL(
       c_concat,
       device_type,
diff --git a/paddle/fluid/operators/load_combine_op.cc b/paddle/fluid/operators/load_combine_op.cc
index 5f03e1304b69d00a38e1558b699ece7a9080634d..3f5b85ecc434baac360d69a5de6e1606cd1e2f80 100644
--- a/paddle/fluid/operators/load_combine_op.cc
+++ b/paddle/fluid/operators/load_combine_op.cc
@@ -77,16 +77,19 @@ that were saved using the SaveCombine operator.
 }  // namespace paddle
 
 namespace ops = paddle::operators;
+namespace plat = paddle::platform;
 
 REGISTER_OPERATOR(load_combine,
                   ops::LoadCombineOp,
                   ops::LoadCombineOpProtoMaker);
 
-REGISTER_OP_CPU_KERNEL(
-    load_combine,
-    ops::LoadCombineOpKernel<phi::CPUContext, float>,
-    ops::LoadCombineOpKernel<phi::CPUContext, double>,
-    ops::LoadCombineOpKernel<phi::CPUContext, paddle::platform::bfloat16>,
-    ops::LoadCombineOpKernel<phi::CPUContext, int>,
-    ops::LoadCombineOpKernel<phi::CPUContext, int8_t>,
-    ops::LoadCombineOpKernel<phi::CPUContext, int64_t>);
+PD_REGISTER_STRUCT_KERNEL(load_combine,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::LoadCombineOpKernel,
+                          float,
+                          double,
+                          plat::bfloat16,
+                          int,
+                          int8_t,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/load_combine_op.cu b/paddle/fluid/operators/load_combine_op.cu
index 9405b3564b9626ed2878ecbc2b70ddbf85efae53..379834daec22e8eb99a766018f6e658222c5c9ee 100644
--- a/paddle/fluid/operators/load_combine_op.cu
+++ b/paddle/fluid/operators/load_combine_op.cu
@@ -15,10 +15,12 @@ limitations under the License. */
 #include "paddle/fluid/operators/load_combine_op.h"
 
 namespace ops = paddle::operators;
-
-REGISTER_OP_CUDA_KERNEL(load_combine,
-                        ops::LoadCombineOpKernel<phi::GPUContext, float>,
-                        ops::LoadCombineOpKernel<phi::GPUContext, double>,
-                        ops::LoadCombineOpKernel<phi::GPUContext, int>,
-                        ops::LoadCombineOpKernel<phi::GPUContext, int8_t>,
-                        ops::LoadCombineOpKernel<phi::GPUContext, int64_t>);
+PD_REGISTER_STRUCT_KERNEL(load_combine,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::LoadCombineOpKernel,
+                          float,
+                          double,
+                          int,
+                          int8_t,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/load_combine_op.h b/paddle/fluid/operators/load_combine_op.h
index 258275f403dcaf6dd2bc5f1c31537008d1d7994b..9f15523ce01297fd93ab937a05dcbd06cb307515 100644
--- a/paddle/fluid/operators/load_combine_op.h
+++ b/paddle/fluid/operators/load_combine_op.h
@@ -28,7 +28,7 @@ limitations under the License. */
 
 namespace paddle {
 namespace operators {
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class LoadCombineOpKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &ctx) const override {
diff --git a/paddle/fluid/operators/load_combine_op_xpu.cc b/paddle/fluid/operators/load_combine_op_xpu.cc
index 9fa7ba3f7526a8fbdb50930821b9c957ca8c1ecf..307c9042c543d1a9b14418723fc780eb65f085d9 100644
--- a/paddle/fluid/operators/load_combine_op_xpu.cc
+++ b/paddle/fluid/operators/load_combine_op_xpu.cc
@@ -15,11 +15,11 @@ limitations under the License. */
 #include "paddle/fluid/operators/load_combine_op.h"
 
 namespace ops = paddle::operators;
+using XPUCtx = paddle::platform::XPUDeviceContext;
 
-REGISTER_OP_XPU_KERNEL(
-    load_combine,
-    ops::LoadCombineOpKernel<paddle::platform::XPUDeviceContext, float>,
-    ops::LoadCombineOpKernel<paddle::platform::XPUDeviceContext, double>,
-    ops::LoadCombineOpKernel<paddle::platform::XPUDeviceContext, int>,
-    ops::LoadCombineOpKernel<paddle::platform::XPUDeviceContext, int8_t>,
-    ops::LoadCombineOpKernel<paddle::platform::XPUDeviceContext, int64_t>);
+REGISTER_OP_XPU_KERNEL(load_combine,
+                       ops::LoadCombineOpKernel<float, XPUCtx>,
+                       ops::LoadCombineOpKernel<double, XPUCtx>,
+                       ops::LoadCombineOpKernel<int, XPUCtx>,
+                       ops::LoadCombineOpKernel<int8_t, XPUCtx>,
+                       ops::LoadCombineOpKernel<int64_t, XPUCtx>);
diff --git a/paddle/fluid/operators/lod_reset_op.cc b/paddle/fluid/operators/lod_reset_op.cc
index e4d4ee2ce9319021d2f9886307347d9d5b4a0b4c..aa5c6bc249cda11c0004ca0a63faf7cd90c394f6 100644
--- a/paddle/fluid/operators/lod_reset_op.cc
+++ b/paddle/fluid/operators/lod_reset_op.cc
@@ -235,6 +235,7 @@ DECLARE_NO_NEED_BUFFER_VARS_INFERER(LoDResetGradNoNeedBufferVarInferer, "X");
 }  // namespace paddle
 
 namespace ops = paddle::operators;
+namespace plat = paddle::platform;
 REGISTER_OPERATOR(lod_reset,
                   ops::LoDResetOp,
                   ops::LoDResetOpMaker,
@@ -247,30 +248,32 @@ REGISTER_OPERATOR(lod_reset_grad,
                   ops::LoDResetGradNoNeedBufferVarInferer,
                   ops::LoDResetGradInplaceInferer);
 
-REGISTER_OP_CPU_KERNEL(
-    lod_reset,
-    ops::LoDResetKernel<paddle::platform::CPUPlace, paddle::platform::float16>,
-    ops::LoDResetKernel<paddle::platform::CPUPlace, float>,
-    ops::LoDResetKernel<paddle::platform::CPUPlace, double>,
-    ops::LoDResetKernel<paddle::platform::CPUPlace, int>,
-    ops::LoDResetKernel<paddle::platform::CPUPlace, int64_t>);
+PD_REGISTER_STRUCT_KERNEL(lod_reset,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::LoDResetKernel,
+                          plat::float16,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
 
 #ifdef PADDLE_WITH_XPU
-REGISTER_OP_XPU_KERNEL(
-    lod_reset,
-    ops::LoDResetKernel<paddle::platform::XPUDeviceContext,
-                        paddle::platform::float16>,
-    ops::LoDResetKernel<paddle::platform::XPUDeviceContext, float>,
-    ops::LoDResetKernel<paddle::platform::XPUDeviceContext, double>,
-    ops::LoDResetKernel<paddle::platform::XPUDeviceContext, int>,
-    ops::LoDResetKernel<paddle::platform::XPUDeviceContext, int64_t>);
+using XPUCtx = paddle::platform::XPUDeviceContext;
+REGISTER_OP_XPU_KERNEL(lod_reset,
+                       ops::LoDResetKernel<paddle::platform::float16, XPUCtx>,
+                       ops::LoDResetKernel<float, XPUCtx>,
+                       ops::LoDResetKernel<double, XPUCtx>,
+                       ops::LoDResetKernel<int, XPUCtx>,
+                       ops::LoDResetKernel<int64_t, XPUCtx>);
 #endif
 
-REGISTER_OP_CPU_KERNEL(
-    lod_reset_grad,
-    ops::LoDResetGradKernel<paddle::platform::CPUPlace,
-                            paddle::platform::float16>,
-    ops::LoDResetGradKernel<paddle::platform::CPUPlace, float>,
-    ops::LoDResetGradKernel<paddle::platform::CPUPlace, double>,
-    ops::LoDResetGradKernel<paddle::platform::CPUPlace, int>,
-    ops::LoDResetGradKernel<paddle::platform::CPUPlace, int64_t>);
+PD_REGISTER_STRUCT_KERNEL(lod_reset_grad,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::LoDResetGradKernel,
+                          plat::float16,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/lod_reset_op.cu b/paddle/fluid/operators/lod_reset_op.cu
index 25aad4c4afce1d4ce93cb5f77b0f907ad936790c..ea1347ee7cd642911a5412b6715ae7b89adc73f1 100644
--- a/paddle/fluid/operators/lod_reset_op.cu
+++ b/paddle/fluid/operators/lod_reset_op.cu
@@ -16,13 +16,19 @@ limitations under the License. */
 
 namespace ops = paddle::operators;
 
-REGISTER_OP_CUDA_KERNEL(lod_reset,
-                        ops::LoDResetKernel<phi::GPUContext, float>,
-                        ops::LoDResetKernel<phi::GPUContext, double>,
-                        ops::LoDResetKernel<phi::GPUContext, int>,
-                        ops::LoDResetKernel<phi::GPUContext, int64_t>);
-REGISTER_OP_CUDA_KERNEL(lod_reset_grad,
-                        ops::LoDResetGradKernel<phi::GPUContext, float>,
-                        ops::LoDResetGradKernel<phi::GPUContext, double>,
-                        ops::LoDResetGradKernel<phi::GPUContext, int>,
-                        ops::LoDResetGradKernel<phi::GPUContext, int64_t>);
+PD_REGISTER_STRUCT_KERNEL(lod_reset,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::LoDResetKernel,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
+PD_REGISTER_STRUCT_KERNEL(lod_reset_grad,
+                          GPU,
+                          ALL_LAYOUT,
+                          ops::LoDResetGradKernel,
+                          float,
+                          double,
+                          int,
+                          int64_t) {}
diff --git a/paddle/fluid/operators/lod_reset_op.h b/paddle/fluid/operators/lod_reset_op.h
index 3b31933174d002a36501d83337b4f1a83d7a0d2f..a4af52472ca9cd8e8b45099140a82983db43a78a 100644
--- a/paddle/fluid/operators/lod_reset_op.h
+++ b/paddle/fluid/operators/lod_reset_op.h
@@ -35,7 +35,7 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class LoDResetKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const {
@@ -123,7 +123,7 @@ class LoDResetKernel : public framework::OpKernel<T> {
   }
 };
 
-template <typename DeviceContext, typename T>
+template <typename T, typename DeviceContext>
 class LoDResetGradKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const {
diff --git a/paddle/fluid/operators/lookup_table_dequant_op.cc b/paddle/fluid/operators/lookup_table_dequant_op.cc
index 09636f600a6611c261eed16dd9c23eade868bf62..f9258f9e0185cd205299bfbc5a09321162367279 100644
--- a/paddle/fluid/operators/lookup_table_dequant_op.cc
+++ b/paddle/fluid/operators/lookup_table_dequant_op.cc
@@ -133,5 +133,9 @@ REGISTER_OPERATOR(
     ops::LookupTableDequantOpMaker,
     paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
     paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>);
-REGISTER_OP_CPU_KERNEL(lookup_table_dequant,
-                       ops::LookupTableDequantKernel<float>);
+
+PD_REGISTER_STRUCT_KERNEL(lookup_table_dequant,
+                          CPU,
+                          ALL_LAYOUT,
+                          ops::LookupTableDequantKernel,
+                          float) {}
diff --git a/paddle/fluid/operators/lookup_table_dequant_op.h b/paddle/fluid/operators/lookup_table_dequant_op.h
index 1c8001e371764cde288a84de418a31a7b827d93f..2f5a3d0fd7a16a92d581da46ca9200b51994465b 100644
--- a/paddle/fluid/operators/lookup_table_dequant_op.h
+++ b/paddle/fluid/operators/lookup_table_dequant_op.h
@@ -46,7 +46,7 @@ void dequant(const unsigned char *in,
 
 constexpr int64_t kNoPadding = -1;
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class LookupTableDequantKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext &context) const override {
diff --git a/paddle/fluid/operators/marker_op.cc b/paddle/fluid/operators/marker_op.cc
index 0cd3ccd686d2c38362b110516555167c087ff37b..91a2131ec8e2cbf916555085e0470488559a8ec3 100644
--- a/paddle/fluid/operators/marker_op.cc
+++ b/paddle/fluid/operators/marker_op.cc
@@ -54,7 +54,7 @@ class MarkerOpMaker : public framework::OpProtoAndCheckerMaker {
   }
 };
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class MarkerOpCPUKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
@@ -75,4 +75,5 @@ class MarkerOpCPUKernel : public framework::OpKernel<T> {
 namespace ops = paddle::operators;
 
 REGISTER_OP_WITHOUT_GRADIENT(marker, ops::MarkerOp, ops::MarkerOpMaker);
-REGISTER_OP_CPU_KERNEL(marker, ops::MarkerOpCPUKernel<float>);
+PD_REGISTER_STRUCT_KERNEL(
+    marker, CPU, ALL_LAYOUT, ops::MarkerOpCPUKernel, float) {}
diff --git a/paddle/fluid/operators/marker_op.cu b/paddle/fluid/operators/marker_op.cu
index 7c6fe79ab7ff128c6b13acf6d0361f90ce24169b..1feb6a2b2616f8cfe4c77bccd2463036991c3649 100644
--- a/paddle/fluid/operators/marker_op.cu
+++ b/paddle/fluid/operators/marker_op.cu
@@ -29,7 +29,7 @@ __global__ void SimpleMarkerKernel(T* in, T* out, int ndim) {
   }
 }
 
-template <typename T>
+template <typename T, typename DeviceContext>
 class MarkerOpCUDAKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
@@ -61,4 +61,5 @@ class MarkerOpCUDAKernel : public framework::OpKernel<T> {
 namespace ops = paddle::operators;
 namespace plat = paddle::platform;
 
-REGISTER_OP_CUDA_KERNEL(marker, ops::MarkerOpCUDAKernel<float>);
+PD_REGISTER_STRUCT_KERNEL(
+    marker, GPU, ALL_LAYOUT, ops::MarkerOpCUDAKernel, float) {}
diff --git a/paddle/fluid/operators/save_load_combine_op_test.cc b/paddle/fluid/operators/save_load_combine_op_test.cc
index 14ff038376b43d900e10e034b367e7aee41b6ce9..896fbe0983681c97a9ee4dd020fba2813674a5db 100644
--- a/paddle/fluid/operators/save_load_combine_op_test.cc
+++ b/paddle/fluid/operators/save_load_combine_op_test.cc
@@ -23,8 +23,9 @@ limitations under the License. */
 #include "paddle/phi/core/kernel_registry.h"
 
 USE_OP_ITSELF(save_combine);
+USE_OP_ITSELF(load_combine);
 PD_DECLARE_KERNEL(save_combine_tensor, CPU, ALL_LAYOUT);
-USE_CPU_ONLY_OP(load_combine);
+PD_DECLARE_KERNEL(load_combine, CPU, ALL_LAYOUT);
 
 template <typename T, typename U>
 T* CreateForSaveCombineOp(int x,