[Fluid] Move random routing to phi (#55773)

paddle/fluid/operators/random_routing_op.cc

@@ -97,10 +97,3 @@ REGISTER_OPERATOR(
     paddle::framework::EmptyGradOpMaker<paddle::framework::OpDesc>,
     paddle::framework::EmptyGradOpMaker<paddle::imperative::OpBase>,
     ops::RandomRoutingInplaceInferer)
-
-PD_REGISTER_STRUCT_KERNEL(random_routing,
-                          CPU,
-                          ALL_LAYOUT,
-                          ops::RandomRoutingOpCPUKernel,
-                          float,
-                          double) {}

paddle/fluid/operators/random_routing_op.h

@@ -27,10 +27,7 @@ namespace operators {
 template <typename T, typename DeviceContext>
 class RandomRoutingOpCPUKernel : public framework::OpKernel<T> {
  public:
-  void Compute(const framework::ExecutionContext& ctx) const override {
-    PADDLE_THROW(platform::errors::Unavailable(
-        "Do not support expert count op for cpu kernel now."));
-  }
+  void Compute(const framework::ExecutionContext& ctx) const override {}
 };
 
 }  // namespace operators

paddle/phi/kernels/cpu/random_routing_kernel.cc

+// Copyright (c) 2023 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "paddle/phi/kernels/random_routing_kernel.h"
+#include "paddle/phi/core/errors.h"
+#include "paddle/phi/core/kernel_registry.h"
+
+namespace phi {
+namespace fusion {
+
+template <typename T, typename Context>
+void RandomRoutingKernel(const Context& dev_ctx,
+                         const DenseTensor& prob,
+                         const DenseTensor& topk_value,
+                         const DenseTensor& topk_idx,
+                         DenseTensor* out) {
+  PADDLE_THROW(phi::errors::Unavailable(
+      "Do not support expert count op for cpu kernel now."));
+}
+
+}  // namespace fusion
+}  // namespace phi
+
+PD_REGISTER_KERNEL(random_routing,
+                   CPU,
+                   ALL_LAYOUT,
+                   phi::fusion::RandomRoutingKernel,
+                   float,
+                   double) {}

paddle/fluid/operators/random_routing_op.cu → paddle/phi/kernels/gpu/random_routing_kernel.cu

-// Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved.
+// Copyright (c) 2023 PaddlePaddle Authors. All Rights Reserved.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -12,13 +12,14 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "paddle/fluid/operators/random_routing_op.h"
-#include "paddle/fluid/framework/op_registry.h"
-#include "paddle/fluid/platform/float16.h"
+#include "paddle/phi/kernels/random_routing_kernel.h"
+#include "paddle/phi/backends/gpu/gpu_helper.h"
 #include "paddle/phi/backends/gpu/gpu_primitives.h"
+#include "paddle/phi/core/dense_tensor.h"
+#include "paddle/phi/core/kernel_registry.h"
+#include "paddle/phi/core/tensor_utils.h"
 
-namespace paddle {
-namespace operators {
+namespace phi {
 
 #define CEIL(_x_, _y_) (((_x_)-1) / (_y_) + 1)
 #define PERTHREAD_EXPERTS 256
@@ -47,45 +48,35 @@ __global__ void random_routing_kernel(int64_t* data,
   }
 }
 
-template <typename T, typename DeviceContext>
-class RandomRoutingOpCUDAKernel : public framework::OpKernel<T> {
- public:
-  void Compute(const framework::ExecutionContext& context) const override {
-    auto topk_idx = context.Input<phi::DenseTensor>("TopK_Idx");
-    auto topk_value = context.Input<phi::DenseTensor>("TopK_Value");
-    auto prob = context.Input<phi::DenseTensor>("Prob");
-    auto out = context.Output<phi::DenseTensor>("Out");
+template <typename T, typename Context>
+void RandomRoutingKernel(const Context& dev_ctx,
+                         const DenseTensor& prob,
+                         const DenseTensor& topk_value,
+                         const DenseTensor& topk_idx,
+                         DenseTensor* out) {
+  phi::Copy(dev_ctx, topk_idx, dev_ctx.GetPlace(), false, out);
 
-    auto place = context.GetPlace();
-    const auto& dev_ctx = context.template device_context<phi::GPUContext>();
-    framework::TensorCopy(*topk_idx, place, out);
+  size_t N = topk_idx.dims()[0];
+  size_t D = topk_idx.dims()[1];
 
-    size_t N = topk_idx->dims()[0];
-    size_t D = topk_idx->dims()[1];
+  int64_t num_idx = topk_idx.numel();
 
-    int64_t num_idx = topk_idx->numel();
+  auto prob_data = prob.data<T>();
+  auto topk_value_data = topk_value.data<T>();
+  auto topk_idx_data = topk_idx.data<int64_t>();
+  auto out_data = out->data<int64_t>();
 
-    auto prob_data = prob->data<T>();
-    auto topk_value_data = topk_value->data<T>();
-    auto topk_idx_data = topk_idx->data<int64_t>();
-    auto out_data = out->data<int64_t>();
-
-    random_routing_kernel<T>
-        <<<GET_BLOCKS(num_idx), CUDA_NUM_THREADS, 0, dev_ctx.stream()>>>(
-            out_data, num_idx, N, D, prob_data, topk_idx_data, topk_value_data);
-  }
-};
-
-}  // namespace operators
-}  // namespace paddle
+  random_routing_kernel<T>
+      <<<GET_BLOCKS(num_idx), CUDA_NUM_THREADS, 0, dev_ctx.stream()>>>(
+          out_data, num_idx, N, D, prob_data, topk_idx_data, topk_value_data);
+}
 
-namespace ops = paddle::operators;
-namespace plat = paddle::platform;
+}  // namespace phi
 
-PD_REGISTER_STRUCT_KERNEL(random_routing,
-                          GPU,
-                          ALL_LAYOUT,
-                          ops::RandomRoutingOpCUDAKernel,
-                          float,
-                          double,
-                          plat::float16) {}
+PD_REGISTER_KERNEL(random_routing,
+                   GPU,
+                   ALL_LAYOUT,
+                   phi::RandomRoutingKernel,
+                   float,
+                   double,
+                   phi::dtype::float16) {}

paddle/phi/kernels/random_routing_kernel.h

+// Copyright (c) 2023 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#pragma once
+
+#include "paddle/phi/common/scalar.h"
+#include "paddle/phi/core/dense_tensor.h"
+
+namespace phi {
+
+template <typename T, typename Context>
+void RandomRoutingKernel(const Context& dev_ctx,
+                         const DenseTensor& prob,
+                         const DenseTensor& topk_value,
+                         const DenseTensor& topk_idx,
+                         DenseTensor* out);
+
+}  // namespace phi

paddle/phi/ops/compat/random_routing_sig.cc

+/* Copyright (c) 2023 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/phi/core/compat/op_utils.h"
+
+namespace phi {
+
+KernelSignature RandomRoutingOpArgumentMapping(
+    const ArgumentMappingContext& ctx) {
+  return KernelSignature(
+      "random_routing", {"Prob", "TopK_Value", "TopK_Idx"}, {}, {"Out"});
+}
+
+}  // namespace phi
+
+PD_REGISTER_ARG_MAPPING_FN(random_routing, phi::RandomRoutingOpArgumentMapping);