Replace platform::GpuLaunchConfig in fluid/operators (#64236)

co63oc · web-flow · commit 34ebff8d47e1 · 2024-05-13T16:54:29.000+08:00
* Fix

* Fix
diff --git a/paddle/fluid/operators/bilateral_slice_op.cu b/paddle/fluid/operators/bilateral_slice_op.cu
@@ -13,7 +13,7 @@
 #include <string>
 
 #include "paddle/fluid/operators/bilateral_slice_op.h"
-#include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
 #include "paddle/phi/backends/gpu/gpu_primitives.h"
 
 namespace paddle {
@@ -169,8 +169,9 @@ class BilateralSliceOpCUDAKernel : public framework::OpKernel<T> {
 
     int total_count = batch_size * h * w * output_dims[1];
 
-    platform::GpuLaunchConfig config =
-        platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), total_count);
+    phi::backends::gpu::GpuLaunchConfig config =
+        phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                                 total_count);
 
     BilateralSliceCudaForwardKernel<T>
         <<<config.block_per_grid,
@@ -502,8 +503,9 @@ class BilateralSliceGradOpCUDAKernel : public framework::OpKernel<T> {
     grid_sizes.gw = gw;
     grid_sizes.input_chans = input_chans;
 
-    platform::GpuLaunchConfig config =
-        platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), grid_count);
+    phi::backends::gpu::GpuLaunchConfig config =
+        phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                                 grid_count);
 
     BilateralSliceCudaGridGradKernel<T>
         <<<config.block_per_grid,
@@ -518,8 +520,8 @@ class BilateralSliceGradOpCUDAKernel : public framework::OpKernel<T> {
                                                  grid_count,
                                                  output_chans);
 
-    config =
-        platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), guide_count);
+    config = phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                                      guide_count);
 
     BilateralSliceCudaGuideGradKernel<T>
         <<<config.block_per_grid,
@@ -535,8 +537,8 @@ class BilateralSliceGradOpCUDAKernel : public framework::OpKernel<T> {
                                                  guide_count,
                                                  output_chans);
 
-    config =
-        platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), input_count);
+    config = phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                                      input_count);
 
     BilateralSliceCudaInputGradKernel<T>
         <<<config.block_per_grid,
diff --git a/paddle/fluid/operators/fused/fused_dropout_act_bias.h b/paddle/fluid/operators/fused/fused_dropout_act_bias.h
@@ -454,8 +454,8 @@ void LaunchDropoutActBiasGrad(Functor act_functor,
     }
   } else {
     const uint64_t n = rows * cols;
-    platform::GpuLaunchConfig config =
-        platform::GetGpuLaunchConfig1D(ctx, n / real_vec_size);
+    phi::backends::gpu::GpuLaunchConfig config =
+        phi::backends::gpu::GetGpuLaunchConfig1D(ctx, n / real_vec_size);
     if (n % VecSize == 0) {
       FusedDropoutActGrad<T, MaskType, VecSize, Functor>
           <<<config.block_per_grid, config.thread_per_block, 0, ctx.stream()>>>(
diff --git a/paddle/fluid/operators/fused/fused_dropout_common.h b/paddle/fluid/operators/fused/fused_dropout_common.h
@@ -20,9 +20,9 @@ limitations under the License. */
 
 #include "paddle/fluid/memory/memory.h"
 #include "paddle/fluid/operators/fused/quant_dequant_kernel.h"
-#include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
 #include "paddle/fluid/platform/device_context.h"
 #include "paddle/phi/backends/gpu/gpu_device_function.h"
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
 #include "paddle/phi/common/amp_type_traits.h"
 #include "paddle/phi/common/float16.h"
 #include "paddle/phi/kernels/funcs/aligned_vector.h"
@@ -40,7 +40,7 @@ namespace operators {
  * 1D blocks: blockDim.x = cols
  * 2D grids: gridDim.y = rows
  */
-inline platform::GpuLaunchConfig Get1DBlocksAnd2DGrids(
+inline phi::backends::gpu::GpuLaunchConfig Get1DBlocksAnd2DGrids(
     const phi::GPUContext &ctx,
     const uint32_t rows,
     const uint32_t cols,
@@ -60,7 +60,7 @@ inline platform::GpuLaunchConfig Get1DBlocksAnd2DGrids(
   const auto blocks_x =
       std::max(static_cast<uint32_t>(1), (tmp_cols + threads - 1) / threads);
   const auto blocks_y = std::max(static_cast<uint32_t>(1), rows);
-  platform::GpuLaunchConfig config;
+  phi::backends::gpu::GpuLaunchConfig config;
   config.block_per_grid.x = blocks_x;
   config.block_per_grid.y = blocks_y;
   config.thread_per_block.x = threads;
diff --git a/paddle/fluid/operators/fused/fused_residual_dropout_bias.h b/paddle/fluid/operators/fused/fused_residual_dropout_bias.h
@@ -484,8 +484,8 @@ void LaunchResidualDropoutBiasGrad(const T *dout,
                                 ctx.stream());                                \
       } else {                                                                \
         const uint64_t n = rows * cols;                                       \
-        platform::GpuLaunchConfig config =                                    \
-            platform::GetGpuLaunchConfig1D(ctx, n / real_vec_size);           \
+        phi::backends::gpu::GpuLaunchConfig config =                          \
+            phi::backends::gpu::GetGpuLaunchConfig1D(ctx, n / real_vec_size); \
         if (n % VecSize == 0) {                                               \
           FusedResidualDropoutGrad<T, MaskType, VecSize>                      \
               <<<config.block_per_grid,                                       \
diff --git a/paddle/fluid/operators/fused/fused_seqpool_cvm_op.cu b/paddle/fluid/operators/fused/fused_seqpool_cvm_op.cu
@@ -16,7 +16,7 @@
 
 #include "paddle/fluid/operators/fused/fused_seqpool_cvm_op.h"
 #include "paddle/fluid/platform/device/gpu/gpu_info.h"
-#include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
 #include "paddle/phi/core/mixed_vector.h"
 
 namespace paddle {
@@ -181,7 +181,8 @@ void FusedSeqpoolCVM(const framework::ExecutionContext
 #endif
 
   size_t N = static_cast<size_t>(batch_size * slot_num * embedding_size);
-  platform::GpuLaunchConfig config = platform::GetGpuLaunchConfig1D(dev_ctx, N);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, N);
   // first sum pool
   FusedSeqpoolKernelNormal<<<config.block_per_grid.x,
                              config.thread_per_block.x,
@@ -208,8 +209,8 @@ void FusedSeqpoolCVM(const framework::ExecutionContext
     // not need show click input
     N = static_cast<size_t>(batch_size * slot_num *
                             (embedding_size - cvm_offset));
-    platform::GpuLaunchConfig config =
-        platform::GetGpuLaunchConfig1D(dev_ctx, N);
+    phi::backends::gpu::GpuLaunchConfig config =
+        phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, N);
     FusedCVMKernelNoCVM<<<config.block_per_grid.x,
                           config.thread_per_block.x,
                           0,
@@ -390,7 +391,7 @@ void FusedSeqpoolCVMGrad(const framework::ExecutionContext &ctx,
 #endif
 
   size_t N = static_cast<size_t>(batch_size * slot_num * embedding_size);
-  auto config = platform::GetGpuLaunchConfig1D(dev_ctx, N);
+  auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, N);
   if (use_cvm) {
     // join grad
     FusedSeqpoolCVMGradKernelWithCVM<<<config.block_per_grid.x,
diff --git a/paddle/fluid/operators/interpolate_op.cu b/paddle/fluid/operators/interpolate_op.cu
@@ -13,7 +13,7 @@
 #include <string>
 
 #include "paddle/fluid/operators/interpolate_op.h"
-#include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
 #include "paddle/phi/backends/gpu/gpu_primitives.h"
 
 namespace paddle {
@@ -979,8 +979,9 @@ static void Interpolate1DCUDAFwd(const framework::ExecutionContext& ctx,
   int out_cw = c * out_w;
   int pixelNum = n * out_cw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("linear" == interp_method) {
     KeLinearInterpFw<T><<<config.block_per_grid,
@@ -1089,8 +1090,9 @@ static void Interpolate2DCUDAFwd(const framework::ExecutionContext& ctx,
 
   int pixelNum = n * out_chw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("nearest" == interp_method) {
     KeNearestNeighborInterpFw<T>
@@ -1255,8 +1257,9 @@ static void Interpolate3DCUDAFwd(const framework::ExecutionContext& ctx,
 
   int pixelNum = n * out_cdhw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("trilinear" == interp_method) {
     KeTrilinearInterpFw<T>
@@ -1353,8 +1356,9 @@ static void Interpolate1DCUDABwd(const framework::ExecutionContext& ctx,
   int out_cw = c * out_w;
   int pixelNum = n * out_cw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("linear" == interp_method) {
     KeLinearInterpBw<T>
@@ -1457,8 +1461,9 @@ static void Interpolate2DCUDABwd(const framework::ExecutionContext& ctx,
 
   int pixelNum = n * out_chw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("nearest" == interp_method) {
     KeNearestNeighborInterpBw<T>
@@ -1611,8 +1616,9 @@ static void Interpolate3DCUDABwd(const framework::ExecutionContext& ctx,
 
   int pixelNum = n * out_cdhw;
 
-  platform::GpuLaunchConfig config =
-      platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);
+  phi::backends::gpu::GpuLaunchConfig config =
+      phi::backends::gpu::GetGpuLaunchConfig1D(ctx.cuda_device_context(),
+                                               pixelNum);
 
   if ("trilinear" == interp_method) {
     KeTrilinearInterpBw<T>
diff --git a/paddle/fluid/operators/temporal_shift_op.cu b/paddle/fluid/operators/temporal_shift_op.cu
@@ -10,7 +10,7 @@
    limitations under the License. */
 
 #include "paddle/fluid/operators/temporal_shift_op.h"
-#include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"
 #include "paddle/phi/backends/gpu/gpu_primitives.h"
 
 namespace paddle {
