   int i = blockIdx.x;
   AccT x_mean = static_cast<AccT>(0);
   AccT x_var = static_cast<AccT>(0);
+  x += i * size;
   const int input_offset = ((uint64_t)x) % ALIGN_BYTES / sizeof(T);
-  x += i * size;
   phi::Array<const T*, 1> ins;
   ins[0] = x;
   ThreadReduce<T, AccT, VecSize, 1>(ins, size, input_offset, &x_mean, &x_var);
-
-  x_mean = kps::details::BlockXReduce<AccT, kps::AddFunctor<AccT>>(
-      x_mean, kps::AddFunctor<AccT>());
-  x_var = kps::details::BlockXReduce<AccT, kps::AddFunctor<AccT>>(
-      x_var, kps::AddFunctor<AccT>());
-  __syncthreads();
-  if (threadIdx.x == 0) {
-    mean[i] = static_cast<T>(x_mean / size);
-    var[i] = static_cast<T>(x_var / size);
-  }
+  ReduceMeanAndVar<AccT>(mean, var, x_mean, x_var, size);
 }
 
 template <typename T, int flags>

    ds_sum += dy[index] * x[index];
    db_sum += dy[index];
  }
  CudaAtomicAddWithWarp(&ds[nc], ds_sum);
  CudaAtomicAddWithWarp(&db[nc], db_sum);
}

template <typename T>

            imsize, C, group_size, groups, p1_data, p2_data, p3_data, x_data,
            dy_data, d_x_data);
      }
    } else {
      if (d_scale) {
        set_zero(dev_ctx, d_scale, static_cast<T>(0));
