[slice]add IndexPutWithSortKernel in index_elementwise_get_grad slice…

paddle/fluid/pybind/slice_utils.h

@@ -782,6 +782,7 @@ static paddle::Tensor getValueForBoolTensor(const paddle::Tensor& tensor,
     }
 
     AdvancedIndex ad = AdvancedIndex(tensor, indices_int64);
+    const bool is_combined = false;
     const bool accumulate = false;
 
     return index_elementwise_get_ad_func(tensor,
@@ -791,7 +792,8 @@ static paddle::Tensor getValueForBoolTensor(const paddle::Tensor& tensor,
                                          ad.indexed_sizes,
                                          ad.indexed_strides,
                                          slice_offset,
-                                         accumulate);
+                                         accumulate,
+                                         is_combined);
   } else {
     if (bool_index.shape().size() == 1)
       return gather_ad_func(tensor, bool_2_idx);
@@ -1238,23 +1240,17 @@ static void ApplyGetitem(const int index_size,
                     &transed_index_int64);
 
       AdvancedIndex ad = AdvancedIndex(*transed_tensor, transed_index_int64);
-      if (index_size == 1) {
-        paddle::Tensor flattened_tensor =
-            flatten_ad_func((*transed_index)[0], 0, -1);
-        *out = gather_ad_func(*transed_tensor, flattened_tensor);
-        *out = reshape_ad_func(*out, ad.src_sizes);
-      } else {
-        const bool accumulate = true;
-        *out = index_elementwise_get_ad_func(*self_tensor,
-                                             ad.indices,
-                                             ad.src_sizes,
-                                             ad.src_strides,
-                                             ad.indexed_sizes,
-                                             ad.indexed_strides,
-                                             slice_offset,
-                                             accumulate);
-      }
-
+      const bool is_combined = (index_size == 1) ? false : true;
+      const bool accumulate = true;
+      *out = index_elementwise_get_ad_func(*self_tensor,
+                                           ad.indices,
+                                           ad.src_sizes,
+                                           ad.src_strides,
+                                           ad.indexed_sizes,
+                                           ad.indexed_strides,
+                                           slice_offset,
+                                           accumulate,
+                                           is_combined);
       return;
     } else {
       paddle::Tensor transed_advanced_index_tensor;

paddle/phi/infermeta/backward.cc

@@ -2168,6 +2168,7 @@ void IndexElementwiseGetGradInferMeta(
     const std::vector<int64_t>& index_strides,
     const int64_t slice_offset,
     const bool accumulate,
+    const bool is_combined,
     MetaTensor* x_grad) {
   if (x_grad) {
     x_grad->share_meta(x);

paddle/phi/infermeta/backward.h

@@ -788,5 +788,6 @@ void IndexElementwiseGetGradInferMeta(
     const std::vector<int64_t>& index_strides,
     const int64_t slice_offset,
     const bool accumulate,
+    const bool is_combined,
     MetaTensor* x_grad);
 }  // namespace phi

paddle/phi/infermeta/binary.cc

@@ -2599,6 +2599,7 @@ void IndexElementwiseGetInferMeta(const MetaTensor& x,
                                   const std::vector<int64_t>& index_stride,
                                   const int64_t slice_offset,
                                   const bool accumulate,
+                                  const bool is_combined,
                                   MetaTensor* out) {
   out->set_dims(common::make_ddim(input_dims));
   out->set_dtype(x.dtype());

paddle/phi/infermeta/binary.h

@@ -493,6 +493,7 @@ void IndexElementwiseGetInferMeta(const MetaTensor& x,
                                   const std::vector<int64_t>& index_stride,
                                   const int64_t slice_offset,
                                   const bool accumulate,
+                                  const bool is_combined,
                                   MetaTensor* out);
 
 void KronInferMeta(const MetaTensor& x, const MetaTensor& y, MetaTensor* out);

paddle/phi/kernels/cpu/index_elementwise_get_grad_kernel.cc

@@ -131,6 +131,7 @@ void IndexElementwiseGetGradKernel(const Context& dev_ctx,
                                    const std::vector<int64_t>& index_strides,
                                    const int64_t slice_offset,
                                    const bool accumulate,
+                                   const bool is_combined,
                                    DenseTensor* x_grad) {
   dev_ctx.template Alloc<T>(x_grad);
   auto dxt = phi::EigenVector<T>::Flatten(*x_grad);

paddle/phi/kernels/cpu/index_elementwise_get_kernel.cc

@@ -100,6 +100,7 @@ void IndexElementwiseGetKernel(const Context& dev_ctx,
                                const std::vector<int64_t>& index_stride,
                                const int64_t slice_offset,
                                const bool accumulate,
+                               const bool is_combined,
                                DenseTensor* out) {
   const auto& index_type = index[0]->dtype();
   PADDLE_ENFORCE_EQ(index_type == phi::DataType::INT64,

paddle/phi/kernels/funcs/radix_sort.cu

@@ -0,0 +1,117 @@
+// Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "paddle/phi/kernels/funcs/radix_sort.h"
+#include "paddle/phi/common/memory_utils.h"
+
+namespace phi {
+namespace funcs {
+
+namespace {
+
+template <typename T>
+struct CudaType {
+  using type = T;
+};
+
+template <>
+struct CudaType<int64_t> {
+  using type = long long;  // NOLINT
+};
+
+#define PADDLE_CUB_WRAPPER(func, ...)                                     \
+  do {                                                                    \
+    size_t temp_storage_bytes = 0;                                        \
+    func(nullptr, temp_storage_bytes, __VA_ARGS__);                       \
+    auto temp_storage =                                                   \
+        phi::memory_utils::Alloc(dev_ctx.GetPlace(), temp_storage_bytes); \
+    func(temp_storage->ptr(), temp_storage_bytes, __VA_ARGS__);           \
+  } while (0)
+
+}  // namespace
+
+template <typename key_t, int value_size>
+void RadixSortPairsImpl(const phi::GPUContext& dev_ctx,
+                        const key_t* keys_in,
+                        key_t* keys_out,
+                        const OpaqueTypeRadix<value_size>* values_in,
+                        OpaqueTypeRadix<value_size>* values_out,
+                        int64_t n,
+                        bool descending,
+                        int64_t begin_bit,
+                        int64_t end_bit) {
+  PADDLE_ENFORCE_LE(
+      n,
+      std::numeric_limits<int>::max(),
+      phi::errors::InvalidArgument(
+          "CUB sort does not support sorting more than INT_MAX elements"));
+
+  using key_t_ = typename CudaType<key_t>::type;
+
+  phi::Allocator::AllocationPtr keys_out_owner;
+  if (keys_out == nullptr) {
+    keys_out_owner =
+        phi::memory_utils::Alloc(dev_ctx.GetPlace(), n * sizeof(key_t));
+    keys_out = reinterpret_cast<key_t*>(keys_out_owner->ptr());
+  }
+
+  const key_t_* keys_in_ = reinterpret_cast<const key_t_*>(keys_in);
+  key_t_* keys_out_ = reinterpret_cast<key_t_*>(keys_out);
+
+  if (descending) {
+    PADDLE_CUB_WRAPPER(cub::DeviceRadixSort::SortPairsDescending,
+                       keys_in_,
+                       keys_out_,
+                       values_in,
+                       values_out,
+                       static_cast<int>(n),
+                       begin_bit,
+                       end_bit,
+                       dev_ctx.stream());
+  } else {
+    PADDLE_CUB_WRAPPER(cub::DeviceRadixSort::SortPairs,
+                       keys_in_,
+                       keys_out_,
+                       values_in,
+                       values_out,
+                       static_cast<int>(n),
+                       begin_bit,
+                       end_bit,
+                       dev_ctx.stream());
+  }
+}
+
+#define INSTANTIATE_SORT_PAIRS(key_t, value_size)      \
+  template void RadixSortPairsImpl<key_t, value_size>( \
+      const phi::GPUContext&,                          \
+      const key_t*,                                    \
+      key_t*,                                          \
+      const OpaqueTypeRadix<value_size>*,              \
+      OpaqueTypeRadix<value_size>*,                    \
+      int64_t,                                         \
+      bool,                                            \
+      int64_t,                                         \
+      int64_t);
+
+INSTANTIATE_SORT_PAIRS(int32_t, 1)
+INSTANTIATE_SORT_PAIRS(int32_t, 2)
+INSTANTIATE_SORT_PAIRS(int32_t, 4)
+INSTANTIATE_SORT_PAIRS(int64_t, 1)
+INSTANTIATE_SORT_PAIRS(int64_t, 2)
+INSTANTIATE_SORT_PAIRS(int64_t, 4)
+INSTANTIATE_SORT_PAIRS(int32_t, 8)
+INSTANTIATE_SORT_PAIRS(int64_t, 8)
+
+}  // namespace funcs
+}  // namespace phi

paddle/phi/kernels/funcs/radix_sort.h

@@ -0,0 +1,80 @@
+// Copyright (c) 2025 PaddlePaddle Authors. All Rights Reserved.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#pragma once
+#include <cub/cub.cuh>
+#include "paddle/phi/backends/gpu/gpu_context.h"
+#include "paddle/phi/core/dense_tensor.h"
+
+namespace phi {
+namespace funcs {
+
+template <int kValueSize>
+struct OpaqueTypeRadix {
+  uint8_t data[kValueSize];
+  __device__ __host__ OpaqueTypeRadix() = default;
+};
+
+template <typename key_t, int kValueSize>
+void RadixSortPairsImpl(const phi::GPUContext& dev_ctx,
+                        const key_t* keys_in,
+                        key_t* keys_out,
+                        const OpaqueTypeRadix<kValueSize>* values_in,
+                        OpaqueTypeRadix<kValueSize>* values_out,
+                        int64_t n,
+                        bool descending = false,
+                        int64_t begin_bit = 0,
+                        int64_t end_bit = sizeof(key_t) * 8);
+
+template <typename key_t, typename value_t>
+void RadixSortPairs(const phi::GPUContext& dev_ctx,
+                    const key_t* keys_in,
+                    key_t* keys_out,
+                    const value_t* values_in,
+                    value_t* values_out,
+                    int64_t n,
+                    bool descending = false,
+                    int64_t begin_bit = 0,
+                    int64_t end_bit = sizeof(key_t) * 8) {
+  PADDLE_ENFORCE_EQ(
+      std::is_trivially_copyable<value_t>::value,
+      true,
+      phi::errors::InvalidArgument(
+          "RadixSortPairs value type must be trivially copyable"));
+
+  using opaque_t = OpaqueTypeRadix<sizeof(value_t)>;
+  PADDLE_ENFORCE_EQ(
+      sizeof(value_t) <= 8 && (sizeof(value_t) & (sizeof(value_t) - 1)) == 0,
+      true,
+      phi::errors::InvalidArgument(
+          "Unsupported value_t size (must be 1, 2, 4, or 8 bytes)"));
+  PADDLE_ENFORCE_EQ(
+      sizeof(value_t),
+      alignof(value_t),
+      phi::errors::InvalidArgument("Expected value_t to be size-aligned"));
+
+  RadixSortPairsImpl<key_t, sizeof(value_t)>(
+      dev_ctx,
+      keys_in,
+      keys_out,
+      reinterpret_cast<const opaque_t*>(values_in),
+      reinterpret_cast<opaque_t*>(values_out),
+      n,
+      descending,
+      begin_bit,
+      end_bit);
+}
+
+}  // namespace funcs
+}  // namespace phi