From d2b0d63fc412f0a72d492930571c359ac7a928f2 Mon Sep 17 00:00:00 2001 From: zhangyuqin1998 <75946871+zhangyuqin1998@users.noreply.github.com> Date: Mon, 17 Apr 2023 11:43:50 +0800 Subject: [PATCH] rename_SliceKernel (#52863) --- paddle/phi/kernels/accuracy_kernel.h | 14 ++++---- paddle/phi/kernels/cpu/accuracy_kernel.cc | 16 +++++----- paddle/phi/kernels/cpu/slice_kernel.cc | 2 +- paddle/phi/kernels/gpu/accuracy_kernel.cu | 16 +++++----- paddle/phi/kernels/gpu/qr_kernel.cu | 12 +++---- paddle/phi/kernels/gpu/slice_kernel.cu.cc | 2 +- paddle/phi/kernels/impl/qr_grad_kernel_impl.h | 12 +++---- paddle/phi/kernels/impl/slice_kernel_impl.h | 16 +++++----- .../phi/kernels/impl/svd_grad_kernel_impl.h | 26 +++++---------- paddle/phi/kernels/onednn/slice_kernel.cc | 18 +++++------ paddle/phi/kernels/slice_kernel.h | 32 +++++++++---------- paddle/phi/kernels/xpu/slice_kernel.cc | 18 +++++------ 12 files changed, 85 insertions(+), 99 deletions(-) diff --git a/paddle/phi/kernels/accuracy_kernel.h b/paddle/phi/kernels/accuracy_kernel.h index 8f2dbb96f86..7ce31052b82 100644 --- a/paddle/phi/kernels/accuracy_kernel.h +++ b/paddle/phi/kernels/accuracy_kernel.h @@ -20,11 +20,11 @@ namespace phi { template -void AccuracyRawKernel(const Context& dev_ctx, - const DenseTensor& out, - const DenseTensor& indices, - const DenseTensor& label, - DenseTensor* accuracy, - DenseTensor* correct, - DenseTensor* total); +void AccuracyKernel(const Context& dev_ctx, + const DenseTensor& out, + const DenseTensor& indices, + const DenseTensor& label, + DenseTensor* accuracy, + DenseTensor* correct, + DenseTensor* total); } // namespace phi diff --git a/paddle/phi/kernels/cpu/accuracy_kernel.cc b/paddle/phi/kernels/cpu/accuracy_kernel.cc index d426e98fbc5..68dd475cb91 100644 --- a/paddle/phi/kernels/cpu/accuracy_kernel.cc +++ b/paddle/phi/kernels/cpu/accuracy_kernel.cc @@ -21,13 +21,13 @@ namespace phi { template -void AccuracyRawKernel(const Context& dev_ctx, - const DenseTensor& inference, - const DenseTensor& indices, - const DenseTensor& label, - DenseTensor* accuracy, - DenseTensor* correct, - DenseTensor* total) { +void AccuracyKernel(const Context& dev_ctx, + const DenseTensor& inference, + const DenseTensor& indices, + const DenseTensor& label, + DenseTensor* accuracy, + DenseTensor* correct, + DenseTensor* total) { int* correct_data = dev_ctx.template Alloc(correct); int* total_data = dev_ctx.template Alloc(total); float* accuracy_data = dev_ctx.template Alloc(accuracy); @@ -93,7 +93,7 @@ void AccuracyRawKernel(const Context& dev_ctx, // TODO(add supported dtype.) PD_REGISTER_KERNEL( - accuracy, CPU, ALL_LAYOUT, phi::AccuracyRawKernel, float, double) { + accuracy, CPU, ALL_LAYOUT, phi::AccuracyKernel, float, double) { kernel->InputAt(1).SetDataType(phi::DataType::INT64); kernel->InputAt(2).SetDataType(phi::DataType::INT64); kernel->OutputAt(0).SetDataType(phi::DataType::FLOAT32); diff --git a/paddle/phi/kernels/cpu/slice_kernel.cc b/paddle/phi/kernels/cpu/slice_kernel.cc index 89f00e71f3a..a6eec654775 100644 --- a/paddle/phi/kernels/cpu/slice_kernel.cc +++ b/paddle/phi/kernels/cpu/slice_kernel.cc @@ -21,7 +21,7 @@ PD_REGISTER_KERNEL(slice, CPU, ALL_LAYOUT, - phi::SliceRawKernel, + phi::SliceKernel, bool, uint8_t, int, diff --git a/paddle/phi/kernels/gpu/accuracy_kernel.cu b/paddle/phi/kernels/gpu/accuracy_kernel.cu index c12afe94b0b..61f20c1a57c 100644 --- a/paddle/phi/kernels/gpu/accuracy_kernel.cu +++ b/paddle/phi/kernels/gpu/accuracy_kernel.cu @@ -73,13 +73,13 @@ __global__ void AccuracyCudaKernel(const int N, } template -void AccuracyRawKernel(const Context& dev_ctx, - const DenseTensor& inference, - const DenseTensor& indices, - const DenseTensor& label, - DenseTensor* accuracy, - DenseTensor* correct, - DenseTensor* total) { +void AccuracyKernel(const Context& dev_ctx, + const DenseTensor& inference, + const DenseTensor& indices, + const DenseTensor& label, + DenseTensor* accuracy, + DenseTensor* correct, + DenseTensor* total) { // FIXME(typhoonzero): only support indices currently // if add support for output values, how to detect the data type? const int64_t* indices_data = indices.data(); @@ -137,7 +137,7 @@ void AccuracyRawKernel(const Context& dev_ctx, PD_REGISTER_KERNEL(accuracy, GPU, ALL_LAYOUT, - phi::AccuracyRawKernel, + phi::AccuracyKernel, phi::dtype::float16, phi::dtype::bfloat16, float, diff --git a/paddle/phi/kernels/gpu/qr_kernel.cu b/paddle/phi/kernels/gpu/qr_kernel.cu index 845a18077c3..14f602cc95b 100644 --- a/paddle/phi/kernels/gpu/qr_kernel.cu +++ b/paddle/phi/kernels/gpu/qr_kernel.cu @@ -101,8 +101,8 @@ void QrKernel(const Context& ctx, if (reduced_mode) { auto trans_qr = TransposeLast2Dim(ctx, qr); - auto sliced_qr = SliceKernel( - ctx, trans_qr, {trans_qr.dims().size() - 2}, {0}, {min_mn}, {1}, {}); + auto sliced_qr = Slice( + ctx, trans_qr, {trans_qr.dims().size() - 2}, {0}, {min_mn}); auto tmp_r = TrilTriu(ctx, sliced_qr, 0, false); // Transpose 'tmp_r' to retore the original row-major order phi::Copy(ctx, tmp_r, r->place(), false, r); @@ -128,8 +128,8 @@ void QrKernel(const Context& ctx, qr_stride, tau_stride); auto trans_q = TransposeLast2Dim(ctx, qr); - auto sliced_q = SliceKernel( - ctx, trans_q, {trans_q.dims().size() - 1}, {0}, {min_mn}, {1}, {}); + auto sliced_q = Slice( + ctx, trans_q, {trans_q.dims().size() - 1}, {0}, {min_mn}); phi::Copy(ctx, sliced_q, q->place(), false, q); } else { if (m > n) { @@ -170,8 +170,8 @@ void QrKernel(const Context& ctx, qr_stride, tau_stride); auto trans_q = TransposeLast2Dim(ctx, qr); - auto sliced_q = SliceKernel( - ctx, trans_q, {trans_q.dims().size() - 1}, {0}, {m}, {1}, {}); + auto sliced_q = Slice( + ctx, trans_q, {trans_q.dims().size() - 1}, {0}, {m}); phi::Copy(ctx, sliced_q, q->place(), false, q); } } diff --git a/paddle/phi/kernels/gpu/slice_kernel.cu.cc b/paddle/phi/kernels/gpu/slice_kernel.cu.cc index 492dc82998b..5b011b32169 100644 --- a/paddle/phi/kernels/gpu/slice_kernel.cu.cc +++ b/paddle/phi/kernels/gpu/slice_kernel.cu.cc @@ -21,7 +21,7 @@ PD_REGISTER_KERNEL(slice, GPU, ALL_LAYOUT, - phi::SliceRawKernel, + phi::SliceKernel, bool, uint8_t, int, diff --git a/paddle/phi/kernels/impl/qr_grad_kernel_impl.h b/paddle/phi/kernels/impl/qr_grad_kernel_impl.h index 5c04d9bb90c..ae0609766eb 100644 --- a/paddle/phi/kernels/impl/qr_grad_kernel_impl.h +++ b/paddle/phi/kernels/impl/qr_grad_kernel_impl.h @@ -149,17 +149,13 @@ void QrGradKernel(const Context& ctx, // Calculate dX and dY individually and concatenate them to get dA ctx.template Alloc>(&dA); - auto Y = SliceKernel( - ctx, A, {A.dims().size() - 1}, {m}, {n}, {1}, {}); - auto U = SliceKernel( - ctx, R, {R.dims().size() - 1}, {0}, {m}, {1}, {}); + auto Y = Slice(ctx, A, {A.dims().size() - 1}, {m}, {n}); + auto U = Slice(ctx, R, {R.dims().size() - 1}, {0}, {m}); DenseTensor dY, dX, dV, dR_tmp, dQ_prime; if (dR.initialized()) { - dV = SliceKernel( - ctx, dR, {dR.dims().size() - 1}, {m}, {n}, {1}, {}); - dR_tmp = SliceKernel( - ctx, dR, {dR.dims().size() - 1}, {0}, {m}, {1}, {}); + dV = Slice(ctx, dR, {dR.dims().size() - 1}, {m}, {n}); + dR_tmp = Slice(ctx, dR, {dR.dims().size() - 1}, {0}, {m}); // Y * dV^H dQ_prime = Matmul(ctx, Y, TransposeLast2Dim(ctx, dV)); diff --git a/paddle/phi/kernels/impl/slice_kernel_impl.h b/paddle/phi/kernels/impl/slice_kernel_impl.h index 8a8925e3e0e..3f81873884f 100644 --- a/paddle/phi/kernels/impl/slice_kernel_impl.h +++ b/paddle/phi/kernels/impl/slice_kernel_impl.h @@ -100,14 +100,14 @@ void SliceCompute(const Context& ctx, } template -void SliceRawKernel(const Context& ctx, - const DenseTensor& input, - const std::vector& axes, - const IntArray& starts_arr, - const IntArray& ends_arr, - const std::vector& infer_flags, - const std::vector& decrease_axis, - DenseTensor* out) { +void SliceKernel(const Context& ctx, + const DenseTensor& input, + const std::vector& axes, + const IntArray& starts_arr, + const IntArray& ends_arr, + const std::vector& infer_flags, + const std::vector& decrease_axis, + DenseTensor* out) { int rank = input.dims().size(); auto& starts = starts_arr.GetData(); diff --git a/paddle/phi/kernels/impl/svd_grad_kernel_impl.h b/paddle/phi/kernels/impl/svd_grad_kernel_impl.h index ee7cab21789..f96073a21bb 100644 --- a/paddle/phi/kernels/impl/svd_grad_kernel_impl.h +++ b/paddle/phi/kernels/impl/svd_grad_kernel_impl.h @@ -83,27 +83,17 @@ void SvdGradKernel(const Context& dev_ctx, DenseTensor U, VH, dU, dV, dVH; if (full_matrices) { // if full_matrices is set, slice the U and VT to k columns - U = SliceKernel( - dev_ctx, u, {u.dims().size() - 1}, {0}, {k}, {1}, {}); - VH = SliceKernel( - dev_ctx, vh, {vh.dims().size() - 2}, {0}, {k}, {1}, {}); + U = Slice(dev_ctx, u, {u.dims().size() - 1}, {0}, {k}); + // If m < n for input matrices A, we partition A = [X|Y] and R = [U|V] + + VH = Slice(dev_ctx, vh, {vh.dims().size() - 2}, {0}, {k}); if (u_grad.get_ptr() != nullptr) { - dU = SliceKernel(dev_ctx, - *(u_grad.get_ptr()), - {u.dims().size() - 1}, - {0}, - {k}, - {1}, - {}); + dU = Slice( + dev_ctx, *(u_grad.get_ptr()), {u.dims().size() - 1}, {0}, {k}); } if (vh_grad.get_ptr() != nullptr) { - dVH = SliceKernel(dev_ctx, - *(vh_grad.get_ptr()), - {vh.dims().size() - 2}, - {0}, - {k}, - {1}, - {}); + dVH = Slice( + dev_ctx, *(vh_grad.get_ptr()), {vh.dims().size() - 2}, {0}, {k}); } } else { U = u; diff --git a/paddle/phi/kernels/onednn/slice_kernel.cc b/paddle/phi/kernels/onednn/slice_kernel.cc index 252c464532d..6c927018264 100644 --- a/paddle/phi/kernels/onednn/slice_kernel.cc +++ b/paddle/phi/kernels/onednn/slice_kernel.cc @@ -20,14 +20,14 @@ namespace phi { template -void SliceRawKernel(const Context& dev_ctx, - const DenseTensor& x, - const std::vector& axes, - const IntArray& starts, - const IntArray& ends, - const std::vector& infer_flags, - const std::vector& decrease_axis, - DenseTensor* out) { +void SliceKernel(const Context& dev_ctx, + const DenseTensor& x, + const std::vector& axes, + const IntArray& starts, + const IntArray& ends, + const std::vector& infer_flags, + const std::vector& decrease_axis, + DenseTensor* out) { const auto& onednn_engine = dev_ctx.GetEngine(); auto x_vec_dims = vectorize(x.dims()); @@ -102,7 +102,7 @@ void SliceRawKernel(const Context& dev_ctx, PD_REGISTER_KERNEL(slice, OneDNN, ONEDNN, - phi::SliceRawKernel, + phi::SliceKernel, float, int8_t, uint8_t, diff --git a/paddle/phi/kernels/slice_kernel.h b/paddle/phi/kernels/slice_kernel.h index 160fde880d7..91de2a3250e 100644 --- a/paddle/phi/kernels/slice_kernel.h +++ b/paddle/phi/kernels/slice_kernel.h @@ -22,14 +22,14 @@ namespace phi { template -void SliceRawKernel(const Context& ctx, - const DenseTensor& input, - const std::vector& axes, - const IntArray& starts, - const IntArray& ends, - const std::vector& infer_flags, - const std::vector& decrease_axis, - DenseTensor* out); +void SliceKernel(const Context& ctx, + const DenseTensor& input, + const std::vector& axes, + const IntArray& starts, + const IntArray& ends, + const std::vector& infer_flags, + const std::vector& decrease_axis, + DenseTensor* out); template void SliceArrayKernel(const Context& dev_ctx, @@ -45,18 +45,18 @@ void SliceArrayDenseKernel(const Context& dev_ctx, DenseTensor* out); template -DenseTensor SliceKernel(const Context& ctx, - const DenseTensor& input, - const std::vector& axes, - const IntArray& starts, - const IntArray& ends, - const std::vector& infer_flags, - const std::vector& decrease_axis) { +DenseTensor Slice(const Context& ctx, + const DenseTensor& input, + const std::vector& axes, + const IntArray& starts, + const IntArray& ends) { DenseTensor dense_out; MetaTensor meta_out(&dense_out); + std::vector infer_flags = {1}; + std::vector decrease_axis = {}; SliceRawInferMeta( input, axes, starts, ends, infer_flags, decrease_axis, &meta_out); - SliceRawKernel( + SliceKernel( ctx, input, axes, starts, ends, infer_flags, decrease_axis, &dense_out); return dense_out; } diff --git a/paddle/phi/kernels/xpu/slice_kernel.cc b/paddle/phi/kernels/xpu/slice_kernel.cc index b30c6908357..cf38c4ecf85 100644 --- a/paddle/phi/kernels/xpu/slice_kernel.cc +++ b/paddle/phi/kernels/xpu/slice_kernel.cc @@ -21,14 +21,14 @@ namespace phi { template -void SliceRawKernel(const Context& ctx, - const DenseTensor& input, - const std::vector& axes, - const IntArray& starts_t, - const IntArray& ends_t, - const std::vector& infer_flags, - const std::vector& decrease_axis, - DenseTensor* out) { +void SliceKernel(const Context& ctx, + const DenseTensor& input, + const std::vector& axes, + const IntArray& starts_t, + const IntArray& ends_t, + const std::vector& infer_flags, + const std::vector& decrease_axis, + DenseTensor* out) { using XPUType = typename XPUTypeTrait::Type; // Step 1: Get the accurate attribute value of starts and ends @@ -110,7 +110,7 @@ void SliceRawKernel(const Context& ctx, PD_REGISTER_KERNEL(slice, XPU, ALL_LAYOUT, - phi::SliceRawKernel, + phi::SliceKernel, float, int, phi::dtype::float16, -- GitLab