From 61486bf286d729c201130f3723a154d64a5c5087 Mon Sep 17 00:00:00 2001
From: Chen Weihang <chenweihang@baidu.com>
Date: Fri, 2 Dec 2022 10:29:04 +0800
Subject: [PATCH] polish fusion kernel naming (#48609)

---
 paddle/phi/kernels/fusion/README.md                    |  4 +++-
 .../kernels/fusion/fused_softmax_mask_grad_kernel.h    |  8 ++++----
 paddle/phi/kernels/fusion/fused_softmax_mask_kernel.h  |  8 ++++----
 .../fusion/gpu/fused_softmax_mask_grad_kernel.cu       | 10 +++++-----
 .../kernels/fusion/gpu/fused_softmax_mask_kernel.cu    | 10 +++++-----
 5 files changed, 21 insertions(+), 19 deletions(-)

diff --git a/paddle/phi/kernels/fusion/README.md b/paddle/phi/kernels/fusion/README.md
index 2080a37dd0f..1e9e2bb7e43 100644
--- a/paddle/phi/kernels/fusion/README.md
+++ b/paddle/phi/kernels/fusion/README.md
@@ -10,4 +10,6 @@
   - Fusion Kernel is generally used to accelerate the combined operation on a certain device. If all devices need to be implemented, the cost is relatively high.
   - We don't recommend implementing a pseudo kernel that just throws exception, if not required, it can be not implemented.
 
-3. Fusion Kernel needs to be in the `phi/fusion` namespace
+3. Fusion Kernel needs to be in the `phi/fusion` namespace.
+
+4. The file naming of the Fusion Kernel needs to follow the format of `fused_[fusion operation name]_kernel.h/cc/cu`, the kernel function naming of the Fusion Kernel needs to follow the format of `Fused[fusion operation name]Kernel`, and the kernel registration naming of the Fusion Kernel needs to follow the format of `fused_[fusion operation name]`.
diff --git a/paddle/phi/kernels/fusion/fused_softmax_mask_grad_kernel.h b/paddle/phi/kernels/fusion/fused_softmax_mask_grad_kernel.h
index 391c614801f..8f4486aa490 100644
--- a/paddle/phi/kernels/fusion/fused_softmax_mask_grad_kernel.h
+++ b/paddle/phi/kernels/fusion/fused_softmax_mask_grad_kernel.h
@@ -19,9 +19,9 @@
 namespace phi {
 
 template <typename T, typename Context>
-void SoftmaxMaskFuseGradKernel(const Context& dev_ctx,
-                               const DenseTensor& out,
-                               const DenseTensor& out_grad,
-                               DenseTensor* x_grad);
+void FusedSoftmaxMaskGradKernel(const Context& dev_ctx,
+                                const DenseTensor& out,
+                                const DenseTensor& out_grad,
+                                DenseTensor* x_grad);
 
 }  // namespace phi
diff --git a/paddle/phi/kernels/fusion/fused_softmax_mask_kernel.h b/paddle/phi/kernels/fusion/fused_softmax_mask_kernel.h
index dd08373f428..1263e8c5d64 100644
--- a/paddle/phi/kernels/fusion/fused_softmax_mask_kernel.h
+++ b/paddle/phi/kernels/fusion/fused_softmax_mask_kernel.h
@@ -19,9 +19,9 @@
 namespace phi {
 
 template <typename T, typename Context>
-void SoftmaxMaskFuseKernel(const Context& dev_ctx,
-                           const DenseTensor& x,
-                           const DenseTensor& mask,
-                           DenseTensor* out);
+void FusedSoftmaxMaskKernel(const Context& dev_ctx,
+                            const DenseTensor& x,
+                            const DenseTensor& mask,
+                            DenseTensor* out);
 
 }  // namespace phi
diff --git a/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_grad_kernel.cu b/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_grad_kernel.cu
index ab731f8f239..d55802fdb96 100644
--- a/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_grad_kernel.cu
+++ b/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_grad_kernel.cu
@@ -118,10 +118,10 @@ __global__ void SoftmaxMaskFuseGradGPUKernel(const T* grad_input,
 }
 
 template <typename T, typename Context>
-void SoftmaxMaskFuseGradKernel(const Context& dev_ctx,
-                               const DenseTensor& out,
-                               const DenseTensor& out_grad,
-                               DenseTensor* x_grad) {
+void FusedSoftmaxMaskGradKernel(const Context& dev_ctx,
+                                const DenseTensor& out,
+                                const DenseTensor& out_grad,
+                                DenseTensor* x_grad) {
   auto* grad_x_data = dev_ctx.template Alloc<T>(x_grad);
   auto* grad_y_data = out_grad.data<T>();
   auto* softmax_rst_data = out.data<T>();
@@ -196,6 +196,6 @@ void SoftmaxMaskFuseGradKernel(const Context& dev_ctx,
 PD_REGISTER_KERNEL(fused_softmax_mask_grad,
                    GPU,
                    ALL_LAYOUT,
-                   phi::fusion::SoftmaxMaskFuseGradKernel,
+                   phi::fusion::FusedSoftmaxMaskGradKernel,
                    float,
                    phi::dtype::float16) {}
diff --git a/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_kernel.cu b/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_kernel.cu
index e86b4841e92..0902b9448ec 100644
--- a/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_kernel.cu
+++ b/paddle/phi/kernels/fusion/gpu/fused_softmax_mask_kernel.cu
@@ -146,10 +146,10 @@ __global__ void SoftmaxMaskFuseGPUKernel(const T* x_data,
 // T only supports fp16
 // leave as template only for future update
 template <typename T, typename Context>
-void SoftmaxMaskFuseKernel(const Context& dev_ctx,
-                           const DenseTensor& x,
-                           const DenseTensor& mask,
-                           DenseTensor* out) {
+void FusedSoftmaxMaskKernel(const Context& dev_ctx,
+                            const DenseTensor& x,
+                            const DenseTensor& mask,
+                            DenseTensor* out) {
   auto* x_data = x.data<T>();
   auto* mask_data = mask.data<T>();
   auto* y_data = dev_ctx.template Alloc<T>(out);
@@ -275,6 +275,6 @@ void SoftmaxMaskFuseKernel(const Context& dev_ctx,
 PD_REGISTER_KERNEL(fused_softmax_mask,
                    GPU,
                    ALL_LAYOUT,
-                   phi::fusion::SoftmaxMaskFuseKernel,
+                   phi::fusion::FusedSoftmaxMaskKernel,
                    float,
                    phi::dtype::float16) {}
-- 
GitLab