PaddlePaddle
diff --git a/‎paddle/phi/kernels/fusion/gpu/cast_with_ptr.h
Copy file name to clipboard
+7-7Lines changed: 7 additions & 7 deletions b/‎paddle/phi/kernels/fusion/gpu/cast_with_ptr.h
Copy file name to clipboard
+7-7Lines changed: 7 additions & 7 deletions
diff --git a/‎paddle/phi/kernels/fusion/gpu/cudnn_bn_stats_finalize.cu.h
Copy file name to clipboard
+16-16Lines changed: 16 additions & 16 deletions b/‎paddle/phi/kernels/fusion/gpu/cudnn_bn_stats_finalize.cu.h
Copy file name to clipboard
+16-16Lines changed: 16 additions & 16 deletions
@@ -27,15 +27,15 @@ struct CastFunctor {
   HOSTDEVICE OutT operator()(InT x) const { return static_cast<OutT>(x); }
 };
 template <typename InT, typename OutT, int VecSize>
-static void VecCastKernel(const phi::GPUContext &ctx,
+static void VecCastKernel(const phi::GPUContext &dev_ctx,
                           const InT *x,
                           OutT *y,
                           size_t n) {
-  auto config = phi::backends::gpu::GetGpuLaunchConfig1D(ctx, n, VecSize);
+  auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, n, VecSize);
   auto block = config.GetGridSize();
   auto thread = config.GetBlockSize();
   auto main_offset = n / (VecSize * thread) * VecSize * thread;
-  auto stream = ctx.stream();
+  auto stream = dev_ctx.stream();
   using FunctorT = CastFunctor<InT, OutT>;
   Array<const _ptr_ char *__restrict__, 1> in_arr;
   in_arr[0] = reinterpret_cast<const _ptr_ char *>(x);
@@ -47,7 +47,7 @@ static void VecCastKernel(const phi::GPUContext &ctx,
 }
 
 template <typename InT, typename OutT>
-static void LaunchCastKernel(const phi::GPUContext &ctx,
+static void LaunchCastKernel(const phi::GPUContext &dev_ctx,
                              const InT *x,
                              OutT *y,
                              size_t n) {
@@ -59,11 +59,11 @@ static void LaunchCastKernel(const phi::GPUContext &ctx,
   int vec_size = std::min(phi::GetVectorizedSize(x), phi::GetVectorizedSize(y));
   switch (vec_size) {
     case 4:
-      return VecCastKernel<InT, OutT, 4>(ctx, x, y, n);
+      return VecCastKernel<InT, OutT, 4>(dev_ctx, x, y, n);
     case 2:
-      return VecCastKernel<InT, OutT, 2>(ctx, x, y, n);
+      return VecCastKernel<InT, OutT, 2>(dev_ctx, x, y, n);
     case 1:
-      return VecCastKernel<InT, OutT, 1>(ctx, x, y, n);
+      return VecCastKernel<InT, OutT, 1>(dev_ctx, x, y, n);
     default:
       PADDLE_THROW(
           errors::InvalidArgument("The vectorized size must be 1, 2 or 4."));
 
@@ -61,15 +61,15 @@ struct BNStatsFinalizeArgs {
 template <typename T>
 class CudnnBNStatsFinalize {
  public:
-  CudnnBNStatsFinalize(const phi::GPUContext &ctx,
+  CudnnBNStatsFinalize(const phi::GPUContext &dev_ctx,
                        const std::vector<int> &param_shape)
       : train_op_(CUDNN_FUSED_BN_FINALIZE_STATISTICS_TRAINING),
         inference_op_(CUDNN_FUSED_BN_FINALIZE_STATISTICS_INFERENCE) {
     args_.Set(param_shape);
   }
   ~CudnnBNStatsFinalize() {}
 
-  void Forward(const phi::GPUContext &ctx,
+  void Forward(const phi::GPUContext &dev_ctx,
                const phi::DenseTensor &sum,
                const phi::DenseTensor &sum_of_squares,
                const phi::DenseTensor &scale,
@@ -85,9 +85,9 @@ class CudnnBNStatsFinalize {
                int64_t ele_count,
                bool is_train) {
     if (is_train) {
-      TrainInit(ctx);
+      TrainInit(dev_ctx);
     } else {
-      InferenceInit(ctx);
+      InferenceInit(dev_ctx);
     }
     auto &op = is_train ? train_op_ : inference_op_;
 
@@ -97,18 +97,18 @@ class CudnnBNStatsFinalize {
         const_cast<float *>(sum_of_squares.data<float>());
     float *scale_ptr = const_cast<float *>(scale.data<float>());
     float *bias_ptr = const_cast<float *>(bias.data<float>());
-    float *saved_mean_ptr = ctx.template Alloc<float>(
+    float *saved_mean_ptr = dev_ctx.template Alloc<float>(
         saved_mean, saved_mean->numel() * sizeof(float));
-    float *saved_invstd_ptr = ctx.template Alloc<float>(
+    float *saved_invstd_ptr = dev_ctx.template Alloc<float>(
         saved_invstd, saved_invstd->numel() * sizeof(float));
-    float *running_mean_ptr = ctx.template Alloc<float>(
+    float *running_mean_ptr = dev_ctx.template Alloc<float>(
         running_mean, running_mean->numel() * sizeof(float));
-    float *running_var_ptr = ctx.template Alloc<float>(
+    float *running_var_ptr = dev_ctx.template Alloc<float>(
         running_var, running_var->numel() * sizeof(float));
-    T *equiv_scale_ptr =
-        ctx.template Alloc<T>(equiv_scale, equiv_scale->numel() * sizeof(T));
+    T *equiv_scale_ptr = dev_ctx.template Alloc<T>(
+        equiv_scale, equiv_scale->numel() * sizeof(T));
     T *equiv_bias_ptr =
-        ctx.template Alloc<T>(equiv_bias, equiv_bias->numel() * sizeof(T));
+        dev_ctx.template Alloc<T>(equiv_bias, equiv_bias->numel() * sizeof(T));
     op.SetOpVariantParamAttrPtr(CUDNN_PTR_BN_SCALE, scale_ptr);
     op.SetOpVariantParamAttrPtr(CUDNN_PTR_BN_BIAS, bias_ptr);
     op.SetOpVariantParamAttrPtr(CUDNN_PTR_BN_RUNNING_MEAN, running_mean_ptr);
@@ -130,12 +130,12 @@ class CudnnBNStatsFinalize {
                                   &avg_factor);
     }
     // fused op execute
-    auto handle = ctx.cudnn_handle();
+    auto handle = dev_ctx.cudnn_handle();
     op.Execute(handle);
   }
 
  private:
-  void TrainInit(const phi::GPUContext &ctx) {
+  void TrainInit(const phi::GPUContext &dev_ctx) {
     // Set constant_param for train op
     train_op_.SetOpConstParamAttr({CUDNN_PARAM_YSUM_PLACEHOLDER,
                                    CUDNN_PARAM_YSQSUM_PLACEHOLDER,
@@ -156,7 +156,7 @@ class CudnnBNStatsFinalize {
                                   args_.out_desc.desc());
 
     // Get workspace
-    auto handle = ctx.cudnn_handle();
+    auto handle = dev_ctx.cudnn_handle();
     train_op_.SetOpConstParamAttr(CUDNN_PARAM_BN_MODE,
                                   CUDNN_BATCHNORM_SPATIAL_PERSISTENT);
     // Check workspace size, also creates plan.
@@ -172,7 +172,7 @@ class CudnnBNStatsFinalize {
                                        &workspace_size_bytes);
   }
 
-  void InferenceInit(const phi::GPUContext &ctx) {
+  void InferenceInit(const phi::GPUContext &dev_ctx) {
     // Set constant_param for inference op
     inference_op_.SetOpConstParamAttr({CUDNN_PARAM_BN_SCALE_PLACEHOLDER,
                                        CUDNN_PARAM_BN_BIAS_PLACEHOLDER,
@@ -188,7 +188,7 @@ class CudnnBNStatsFinalize {
                                       args_.out_desc.desc());
 
     // Get workspace
-    auto handle = ctx.cudnn_handle();
+    auto handle = dev_ctx.cudnn_handle();
     inference_op_.SetOpConstParamAttr(CUDNN_PARAM_BN_MODE,
                                       CUDNN_BATCHNORM_SPATIAL_PERSISTENT);
     // Check workspace size, also creates plan.