Fix

Author: co63oc

paddle/phi/kernels/funcs/binomial_functor.h

@@ -40,11 +40,11 @@ inline T stirling_approx_tail(int64_t k) {
 }
 
 template <typename T, typename Context>
-inline int64_t btrs(const Context& ctx, const T n, const T p) {
+inline int64_t btrs(const Context& dev_ctx, const T n, const T p) {
  int64_t k;
  T U, V, us;
  std::uniform_real_distribution<T> dist(0.0, 1.0);
- auto gen_ptr = ctx.GetGenerator();
+ auto gen_ptr = dev_ctx.GetGenerator();
  auto engine = gen_ptr->GetCPUEngine();
 
  const T stddev = std::sqrt(n * p * (1 - p));
@@ -87,13 +87,15 @@ inline int64_t btrs(const Context& ctx, const T n, const T p) {
 }
 
 template <typename T, typename Context>
-inline int64_t binomial_inversion(const Context& ctx, const T n, const T p) {
+inline int64_t binomial_inversion(const Context& dev_ctx,
+ const T n,
+ const T p) {
  T unif;
  T geom_sum = 0.0;
  int64_t num_geom = 0;
  T logprob = std::log1p(-p);
  std::uniform_real_distribution<T> dist(0.0, 1.0);
- auto gen_ptr = ctx.GetGenerator();
+ auto gen_ptr = dev_ctx.GetGenerator();
  auto engine = gen_ptr->GetCPUEngine();
 
  while (1) {
@@ -109,23 +111,23 @@ inline int64_t binomial_inversion(const Context& ctx, const T n, const T p) {
 }
 
 template <typename T, typename Context>
-inline int64_t BinomialFunctor(const Context& ctx, const T n, const T p) {
+inline int64_t BinomialFunctor(const Context& dev_ctx, const T n, const T p) {
  if (n <= 0.0 || p <= 0.0) {
  return 0;
  } else if (p >= 1.0) {
  return static_cast<int64_t>(n);
  } else if (p <= 0.5) {
  if (n * p >= 10.0) {
- return btrs<T>(ctx, n, p);
+ return btrs<T>(dev_ctx, n, p);
  } else {
- return binomial_inversion<T>(ctx, n, p);
+ return binomial_inversion<T>(dev_ctx, n, p);
  }
  } else {
  T qprob = 1.0 - p;
  if (n * qprob >= 10.0) {
- return static_cast<int64_t>(n) - btrs<T>(ctx, n, qprob);
+ return static_cast<int64_t>(n) - btrs<T>(dev_ctx, n, qprob);
  } else {
- return static_cast<int64_t>(n) - binomial_inversion<T>(ctx, n, qprob);
+ return static_cast<int64_t>(n) - binomial_inversion<T>(dev_ctx, n, qprob);
  }
  }
 }

paddle/phi/kernels/funcs/broadcast_function.h

@@ -436,15 +436,15 @@ __global__ void VectorizedBroadcastKernel(
 
 template <typename OutT, typename Functor, int Arity, int NumOuts, int VecSize>
 void LaunchBroadcastKernel(
- const KPDevice &ctx,
+ const KPDevice &dev_ctx,
  const BroadcastTypeClassifier<OutT, Functor, Arity, NumOuts> &classifier,
  Functor func) {
 #ifdef PADDLE_WITH_XPU_KP
  int numel = classifier.numel;
  const int threads = 64;
  const int blocks = 8;
  int read_lens = configs[0].buf_len;
- auto stream = ctx.x_context()->xpu_stream;
+ auto stream = dev_ctx.x_context()->xpu_stream;
  int main_offset = (numel / (read_lens * threads)) * read_lens * threads;
  int tail_tid = numel % (read_lens * threads);
 
@@ -461,8 +461,8 @@ void LaunchBroadcastKernel(
 #else
  const int &numel = classifier.numel;
  auto gpu_config =
- phi::backends::gpu::GetGpuLaunchConfig1D(ctx, numel, VecSize);
- auto stream = ctx.stream();
+ phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, numel, VecSize);
+ auto stream = dev_ctx.stream();
  auto threads = gpu_config.GetBlockSize();
  auto blocks = gpu_config.block_per_grid;
  int main_offset = (numel / (VecSize * threads)) * VecSize * threads;
@@ -513,20 +513,20 @@ void LaunchBroadcastKernel(
 
 template <typename OutT, typename Functor, int Arity, int NumOuts = 1>
 typename std::enable_if<!NeedVectorized<OutT>::value, void>::type
-BroadcastKernelForDifferentVecSize(const KPDevice &ctx,
+BroadcastKernelForDifferentVecSize(const KPDevice &dev_ctx,
  const std::vector<const DenseTensor *> &ins,
  std::vector<DenseTensor *> *outs,
  int axis,
  Functor func) {
  auto classifier =
  BroadcastTypeClassifier<OutT, Functor, Arity, NumOuts>(ins, outs, axis);
  LaunchBroadcastKernel<OutT, Functor, Arity, NumOuts, VecSizeS>(
- ctx, classifier, func);
+ dev_ctx, classifier, func);
 }
 
 template <typename OutT, typename Functor, int Arity, int NumOuts = 1>
 typename std::enable_if<NeedVectorized<OutT>::value, void>::type
-BroadcastKernelForDifferentVecSize(const KPDevice &ctx,
+BroadcastKernelForDifferentVecSize(const KPDevice &dev_ctx,
  const std::vector<const DenseTensor *> &ins,
  std::vector<DenseTensor *> *outs,
  int axis,
@@ -545,17 +545,17 @@ BroadcastKernelForDifferentVecSize(const KPDevice &ctx,
  switch (vec_size) {
  case VecSizeL: {
  LaunchBroadcastKernel<OutT, Functor, Arity, NumOuts, VecSizeL>(
- ctx, classifier, func);
+ dev_ctx, classifier, func);
  break;
  }
  case VecSizeM: {
  LaunchBroadcastKernel<OutT, Functor, Arity, NumOuts, VecSizeM>(
- ctx, classifier, func);
+ dev_ctx, classifier, func);
  break;
  }
  case VecSizeS: {
  LaunchBroadcastKernel<OutT, Functor, Arity, NumOuts, VecSizeS>(
- ctx, classifier, func);
+ dev_ctx, classifier, func);
  break;
  }
  default: {
@@ -591,7 +591,7 @@ static void SliceTensor(DenseTensor *x,
 }
 
 template <typename OutT, typename Functor, int kArity, int NumOuts = 1>
-void BroadcastKernelSplit(const KPDevice &ctx,
+void BroadcastKernelSplit(const KPDevice &dev_ctx,
  const std::vector<const DenseTensor *> &ins,
  std::vector<DenseTensor *> *outs,
  int axis,
@@ -728,12 +728,12 @@ void BroadcastKernelSplit(const KPDevice &ctx,
  }
 
  BroadcastKernelForDifferentVecSize<OutT, Functor, kArity, NumOuts>(
- ctx, new_ins, &new_outs, axis, func);
+ dev_ctx, new_ins, &new_outs, axis, func);
  }
 }
 
 template <typename OutT, typename Functor, int kArity, int NumOuts = 1>
-void BroadcastKernelApply(const KPDevice &ctx,
+void BroadcastKernelApply(const KPDevice &dev_ctx,
  const std::vector<const DenseTensor *> &ins,
  std::vector<DenseTensor *> *outs,
  int axis,
@@ -748,16 +748,16 @@ void BroadcastKernelApply(const KPDevice &ctx,
  }
  if (use_int64_index_kernel) { // use_int64_index_kernel
  BroadcastKernelSplit<OutT, Functor, kArity, NumOuts>(
- ctx, ins, outs, axis, func, compute_size);
+ dev_ctx, ins, outs, axis, func, compute_size);
  return;
  }
 #endif
  BroadcastKernelForDifferentVecSize<OutT, Functor, kArity, NumOuts>(
- ctx, ins, outs, axis, func);
+ dev_ctx, ins, outs, axis, func);
 }
 
 template <typename OutT, typename Functor, int NumOuts = 1>
-void BroadcastKernel(const KPDevice &ctx,
+void BroadcastKernel(const KPDevice &dev_ctx,
  const std::vector<const DenseTensor *> &ins,
  std::vector<DenseTensor *> *outs,
  Functor func,
@@ -805,7 +805,7 @@ void BroadcastKernel(const KPDevice &ctx,
  "%d-th output tensor`s shape is not.",
  i));
  }
- ctx.template Alloc<OutT>((*outs)[i]);
+ dev_ctx.template Alloc<OutT>((*outs)[i]);
  }
  if ((*outs)[0]->numel() == 0) {
  return;
@@ -823,7 +823,7 @@ void BroadcastKernel(const KPDevice &ctx,
  }
  axis = axis == -1 ? max_rank - min_rank : axis;
  BroadcastKernelApply<OutT, Functor, kArity, NumOuts>(
- ctx, ins, outs, axis, func);
+ dev_ctx, ins, outs, axis, func);
 }
 
 template <typename Functor, typename T, typename OutType = T>