PaddlePaddle
diff --git a/‎paddle/phi/kernels/gpu/partial_sum_kernel.cu
Copy file name to clipboard
-1Lines changed: 0 additions & 1 deletion b/‎paddle/phi/kernels/gpu/partial_sum_kernel.cu
Copy file name to clipboard
-1Lines changed: 0 additions & 1 deletion
diff --git a/‎paddle/phi/kernels/gpu/poisson_kernel.cu
Copy file name to clipboard
+7-5Lines changed: 7 additions & 5 deletions b/‎paddle/phi/kernels/gpu/poisson_kernel.cu
Copy file name to clipboard
+7-5Lines changed: 7 additions & 5 deletions
diff --git a/‎paddle/phi/kernels/gpu/polygamma_grad_kernel.cu
Copy file name to clipboard
+3-3Lines changed: 3 additions & 3 deletions b/‎paddle/phi/kernels/gpu/polygamma_grad_kernel.cu
Copy file name to clipboard
+3-3Lines changed: 3 additions & 3 deletions
diff --git a/‎paddle/phi/kernels/gpu/polygamma_kernel.cu
Copy file name to clipboard
+3-3Lines changed: 3 additions & 3 deletions b/‎paddle/phi/kernels/gpu/polygamma_kernel.cu
Copy file name to clipboard
+3-3Lines changed: 3 additions & 3 deletions
diff --git a/‎paddle/phi/kernels/gpu/prior_box_kernel.cu
Copy file name to clipboard
+10-10Lines changed: 10 additions & 10 deletions b/‎paddle/phi/kernels/gpu/prior_box_kernel.cu
Copy file name to clipboard
+10-10Lines changed: 10 additions & 10 deletions
diff --git a/‎paddle/phi/kernels/gpu/psroi_pool_grad_kernel.cu
Copy file name to clipboard
+10-9Lines changed: 10 additions & 9 deletions b/‎paddle/phi/kernels/gpu/psroi_pool_grad_kernel.cu
Copy file name to clipboard
+10-9Lines changed: 10 additions & 9 deletions
diff --git a/‎paddle/phi/kernels/gpu/psroi_pool_kernel.cu
Copy file name to clipboard
+22-17Lines changed: 22 additions & 17 deletions b/‎paddle/phi/kernels/gpu/psroi_pool_kernel.cu
Copy file name to clipboard
+22-17Lines changed: 22 additions & 17 deletions
@@ -75,7 +75,6 @@ void PartialSumOpCUDAKernel(const Context &dev_ctx,
                             int start_index,
                             int length,
                             DenseTensor *out) {
-  auto ctx = dev_ctx;
   auto in_vars = x;
 
   PADDLE_ENFORCE_EQ(
 
@@ -44,18 +44,20 @@ __global__ void GetPoisson(
 }
 
 template <typename T, typename Context>
-void PoissonKernel(const Context& ctx, const DenseTensor& x, DenseTensor* out) {
+void PoissonKernel(const Context& dev_ctx,
+                   const DenseTensor& x,
+                   DenseTensor* out) {
   const T* x_data = x.data<T>();
-  T* out_data = ctx.template Alloc<T>(out);
+  T* out_data = dev_ctx.template Alloc<T>(out);
   const int size = x.numel();
   const int kMaxBlockDim = 256;
 
-  int block_size = std::min(kMaxBlockDim, ctx.GetMaxThreadsPerBlock());
+  int block_size = std::min(kMaxBlockDim, dev_ctx.GetMaxThreadsPerBlock());
   dim3 dim_block(block_size);
   dim3 dim_grid((size + block_size - 1) / block_size);
-  phi::backends::gpu::LimitGridDim(ctx, &dim_grid);
+  phi::backends::gpu::LimitGridDim(dev_ctx, &dim_grid);
 
-  auto gen_cuda = ctx.GetGenerator();
+  auto gen_cuda = dev_ctx.GetGenerator();
   auto seed_offset = gen_cuda->IncrementOffset(20);
   uint64_t seed = seed_offset.first;
   uint64_t offset = seed_offset.second;
 
@@ -22,16 +22,16 @@
 namespace phi {
 
 template <typename T, typename Context>
-void PolygammaGradKernel(const Context& ctx,
+void PolygammaGradKernel(const Context& dev_ctx,
                          const DenseTensor& x,
                          const DenseTensor& out_grad,
                          const int n,
                          DenseTensor* x_grad) {
-  ctx.template Alloc<T>(x_grad);
+  dev_ctx.template Alloc<T>(x_grad);
   std::vector<const DenseTensor*> ins = {&x, &out_grad};
   std::vector<DenseTensor*> outs = {x_grad};
   auto functor = CudaPolygammaGradFunctor<T>(n + 1);
-  phi::funcs::ElementwiseKernel<T>(ctx, ins, &outs, functor);
+  phi::funcs::ElementwiseKernel<T>(dev_ctx, ins, &outs, functor);
 }
 
 }  // namespace phi
 
@@ -22,16 +22,16 @@
 namespace phi {
 
 template <typename T, typename Context>
-void PolygammaKernel(const Context& ctx,
+void PolygammaKernel(const Context& dev_ctx,
                      const DenseTensor& x,
                      const int n,
                      DenseTensor* out) {
-  ctx.template Alloc<T>(out);
+  dev_ctx.template Alloc<T>(out);
   std::vector<const DenseTensor*> ins = {&x};
   std::vector<DenseTensor*> outs = {out};
 
   auto functor = CudaPolygammaFunctor<T>(n);
-  phi::funcs::ElementwiseKernel<T>(ctx, ins, &outs, functor);
+  phi::funcs::ElementwiseKernel<T>(dev_ctx, ins, &outs, functor);
 }
 
 }  // namespace phi
 
@@ -108,7 +108,7 @@ __global__ void SetVariance(T* out,
 }
 
 template <typename T, typename Context>
-void PriorBoxKernel(const Context& ctx,
+void PriorBoxKernel(const Context& dev_ctx,
                     const DenseTensor& input,
                     const DenseTensor& image,
                     const std::vector<float>& min_sizes,
@@ -125,9 +125,9 @@ void PriorBoxKernel(const Context& ctx,
                     DenseTensor* var) {
   if (input.numel() == 0 || image.numel() == 0) {
     phi::Full<T, Context>(
-        ctx, phi::IntArray(common::vectorize(out->dims())), 0, out);
+        dev_ctx, phi::IntArray(common::vectorize(out->dims())), 0, out);
     phi::Full<T, Context>(
-        ctx, phi::IntArray(common::vectorize(var->dims())), 0, var);
+        dev_ctx, phi::IntArray(common::vectorize(var->dims())), 0, var);
     return;
   }
 
@@ -163,21 +163,21 @@ void PriorBoxKernel(const Context& ctx,
   int block = 512;
   int grid = (box_num + block - 1) / block;
 
-  auto stream = ctx.stream();
+  auto stream = dev_ctx.stream();
 
-  ctx.template Alloc<T>(out);
-  ctx.template Alloc<T>(var);
+  dev_ctx.template Alloc<T>(out);
+  dev_ctx.template Alloc<T>(var);
 
   DenseTensor r;
-  phi::TensorFromVector(new_aspect_ratios, ctx, &r);
+  phi::TensorFromVector(new_aspect_ratios, dev_ctx, &r);
 
   DenseTensor min;
-  phi::TensorFromVector(min_sizes, ctx, &min);
+  phi::TensorFromVector(min_sizes, dev_ctx, &min);
 
   T* max_data = nullptr;
   DenseTensor max;
   if (max_sizes.size() > 0) {
-    phi::TensorFromVector(max_sizes, ctx, &max);
+    phi::TensorFromVector(max_sizes, dev_ctx, &max);
     max_data = max.data<T>();
   }
 
@@ -198,7 +198,7 @@ void PriorBoxKernel(const Context& ctx,
                                              min_max_aspect_ratios_order);
 
   DenseTensor v;
-  phi::TensorFromVector(variances, ctx, &v);
+  phi::TensorFromVector(variances, dev_ctx, &v);
   grid = (box_num * 4 + block - 1) / block;
   SetVariance<T><<<grid, block, 0, stream>>>(
       var->data<T>(), v.data<T>(), variances.size(), box_num * 4);
 
@@ -104,7 +104,7 @@ __global__ void GPUPSROIPoolBackward(const int nthreads,
 }
 
 template <typename T, typename Context>
-void PsroiPoolGradKernel(const Context& ctx,
+void PsroiPoolGradKernel(const Context& dev_ctx,
                          const DenseTensor& x,
                          const DenseTensor& rois,
                          const paddle::optional<DenseTensor>& rois_num,
@@ -123,14 +123,15 @@ void PsroiPoolGradKernel(const Context& ctx,
     // set roi batch id
     DenseTensor rois_batch_id_list;
     rois_batch_id_list.Resize({rois_num_t});
-    int* rois_batch_id_data = ctx.template HostAlloc<int>(&rois_batch_id_list);
+    int* rois_batch_id_data =
+        dev_ctx.template HostAlloc<int>(&rois_batch_id_list);
     int rois_batch_size;
     if (rois_num.get_ptr()) {
       rois_batch_size = rois_num->numel();
       std::vector<int> rois_num_list(rois_batch_size);
       memory_utils::Copy(CPUPlace(),
                          rois_num_list.data(),
-                         ctx.GetPlace(),
+                         dev_ctx.GetPlace(),
                          rois_num->data<int>(),
                          sizeof(int) * rois_batch_size,
                          0);
@@ -152,22 +153,22 @@ void PsroiPoolGradKernel(const Context& ctx,
     }
 
     DenseTensor rois_batch_id_list_gpu;
-    Copy(ctx,
+    Copy(dev_ctx,
          rois_batch_id_list,
-         ctx.GetPlace(),
+         dev_ctx.GetPlace(),
          false,
          &rois_batch_id_list_gpu);
 
-    ctx.template Alloc<T>(dx);
+    dev_ctx.template Alloc<T>(dx);
     funcs::SetConstant<Context, T> set_zero;
-    set_zero(ctx, dx, static_cast<T>(0));
+    set_zero(dev_ctx, dx, static_cast<T>(0));
 
     int dout_size = dout.numel();
     int blocks = NumBlocks(dout_size);
     int threads = kNumCUDAThreads;
 
     if (dout_size > 0) {
-      GPUPSROIPoolBackward<T><<<blocks, threads, 0, ctx.stream()>>>(
+      GPUPSROIPoolBackward<T><<<blocks, threads, 0, dev_ctx.stream()>>>(
           dout_size,
           rois.data<T>(),
           dout.data<T>(),
@@ -179,7 +180,7 @@ void PsroiPoolGradKernel(const Context& ctx,
           pooled_height,
           pooled_width,
           rois_batch_id_list_gpu.data<int>(),
-          ctx.template Alloc<T>(dx));
+          dev_ctx.template Alloc<T>(dx));
     }
   }
 }
 
@@ -105,7 +105,7 @@ __global__ void GPUPSROIPoolForward(const int nthreads,
 }
 
 template <typename T, typename Context>
-void PsroiPoolKernel(const Context& ctx,
+void PsroiPoolKernel(const Context& dev_ctx,
                      const DenseTensor& x,
                      const DenseTensor& rois,
                      const paddle::optional<DenseTensor>& rois_num,
@@ -136,7 +136,8 @@ void PsroiPoolKernel(const Context& ctx,
   int rois_batch_size;
   DenseTensor rois_batch_id_list;
   rois_batch_id_list.Resize({rois_num_t});
-  int* rois_batch_id_data = ctx.template HostAlloc<int>(&rois_batch_id_list);
+  int* rois_batch_id_data =
+      dev_ctx.template HostAlloc<int>(&rois_batch_id_list);
 
   if (rois_num.get_ptr()) {
     rois_batch_size = rois_num->numel();
@@ -152,7 +153,7 @@ void PsroiPoolKernel(const Context& ctx,
     std::vector<int> rois_num_list(rois_batch_size);
     memory_utils::Copy(CPUPlace(),
                        rois_num_list.data(),
-                       ctx.GetPlace(),
+                       dev_ctx.GetPlace(),
                        rois_num_data,
                        sizeof(int) * rois_batch_size,
                        0);
@@ -201,26 +202,30 @@ void PsroiPoolKernel(const Context& ctx,
     }
   }
   DenseTensor rois_batch_id_list_gpu;
-  Copy(ctx, rois_batch_id_list, ctx.GetPlace(), false, &rois_batch_id_list_gpu);
+  Copy(dev_ctx,
+       rois_batch_id_list,
+       dev_ctx.GetPlace(),
+       false,
+       &rois_batch_id_list_gpu);
 
   int output_size = out->numel();
   int blocks = NumBlocks(output_size);
   int threads = kNumCUDAThreads;
 
   // call cuda kernel function
-  GPUPSROIPoolForward<T>
-      <<<blocks, threads, 0, ctx.stream()>>>(output_size,
-                                             x.data<T>(),
-                                             rois.data<T>(),
-                                             spatial_scale,
-                                             input_channels,
-                                             height,
-                                             width,
-                                             output_channels,
-                                             pooled_height,
-                                             pooled_width,
-                                             rois_batch_id_list_gpu.data<int>(),
-                                             ctx.template Alloc<T>(out));
+  GPUPSROIPoolForward<T><<<blocks, threads, 0, dev_ctx.stream()>>>(
+      output_size,
+      x.data<T>(),
+      rois.data<T>(),
+      spatial_scale,
+      input_channels,
+      height,
+      width,
+      output_channels,
+      pooled_height,
+      pooled_width,
+      rois_batch_id_list_gpu.data<int>(),
+      dev_ctx.template Alloc<T>(out));
 }
 
 }  // namespace phi