Do not invoke map_back kernel if iter_nelems is 1

The map_back operation is a no-op then.

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/sorting/merge_sort.hpp

@@ -814,6 +814,10 @@ sycl::event stable_argsort_axis1_contig_impl(
         exec_q, iter_nelems, sort_nelems, res_tp, index_comp, sorted_block_size,
         {base_sort_ev});
 
+    // no need to map back if iter_nelems == 1
+    if (iter_nelems == 1u)
+        return merges_ev;
+
     using MapBackKernelName = index_map_to_rows_krn<argTy, IndexTy>;
     using dpctl::tensor::kernels::sort_utils_detail::map_back_impl;
 

dpctl/tensor/libtensor/include/kernels/sorting/radix_sort.hpp

@@ -1867,11 +1867,17 @@ radix_argsort_axis1_contig_impl(sycl::queue &exec_q,
     using MapBackKernelName = radix_argsort_index_write_out_krn<argTy, IndexTy>;
     using dpctl::tensor::kernels::sort_utils_detail::map_back_impl;
 
-    sycl::event map_back_ev = map_back_impl<MapBackKernelName, IndexTy>(
-        exec_q, total_nelems, res_tp, res_tp, sort_nelems, {radix_sort_ev});
+    sycl::event dep = radix_sort_ev;
+
+    // no need to perform map_back ( id % sort_nelems)
+    //   if total_nelems == sort_nelems
+    if (iter_nelems > 1u) {
+        dep = map_back_impl<MapBackKernelName, IndexTy>(
+            exec_q, total_nelems, res_tp, res_tp, sort_nelems, {dep});
+    }
 
     sycl::event cleanup_ev = dpctl::tensor::alloc_utils::async_smart_free(
-        exec_q, {map_back_ev}, workspace_owner);
+        exec_q, {dep}, workspace_owner);
 
     return cleanup_ev;
 }