include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge_sort…

….h - trace Kernel's execution time info
uxlfoundation · Jan 17, 2025 · cefb462 · cefb462
1 parent ad1d290
commit cefb462
Showing 1 changed file with 34 additions and 0 deletions.
diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge_sort.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_merge_sort.h
@@ -42,6 +42,12 @@
 const __SYCL_CONSTANT_AS char fmt_create_diags_storage[] = "Create base diagonals storage: __n_sorted = %8d, storage size = [%8d]\n";
 const __SYCL_CONSTANT_AS char fmt_eval_base_diag[] = "Evaluate base diag: __n_sorted = %8d, __data_in_temp = %s, sp[%8d] = (%8d, %8d)\n";
 const __SYCL_CONSTANT_AS char fmt_lookup_sp[] = "Lookup sp: __linear_id = %8d, sp_left[%8d] = (%8d, %8d), sp_right[%8d] = (%8d, %8d)\n";
+
+#define MERGE_SORT_DISPLAY_STATISTIC 1
+#define MERGE_SORT_EXCLUDE_NEW_IMPL  0
+
+#if MERGE_SORT_DISPLAY_STATISTIC
+#include <chrono>
 #endif
 
 namespace oneapi
@@ -625,12 +631,19 @@ struct __merge_sort_global_submitter<_IndexT, __internal::__optional_kernel_name
 
         for (std::int64_t __i = 0; __i < __n_iter; ++__i)
         {
+#if MERGE_SORT_DISPLAY_STATISTIC
+            const auto __start_time = std::chrono::high_resolution_clock::now();
+#endif
+
+#if !MERGE_SORT_EXCLUDE_NEW_IMPL
             if (2 * __n_sorted < __get_starting_size_limit_for_large_submitter<__value_type>())
+#endif
             {
                 // Process parallel merge
                 __event_chain = run_parallel_merge(__event_chain, __n_sorted, __data_in_temp, __exec, __rng, __temp_buf,
                                                    __comp, __nd_range_params);
             }
+#if !MERGE_SORT_EXCLUDE_NEW_IMPL
             else
             {
                 if (nullptr == __p_base_diagonals_sp_global_storage)
@@ -660,6 +673,15 @@ struct __merge_sort_global_submitter<_IndexT, __internal::__optional_kernel_name
                                                                   __rng, __temp_buf, __comp, __nd_range_params_this,
                                                                   *__p_base_diagonals_sp_global_storage);
             }
+#endif
+
+#if MERGE_SORT_DISPLAY_STATISTIC
+            __event_chain.wait_and_throw();
+            const auto __stop_time = std::chrono::high_resolution_clock::now();
+
+            const std::chrono::duration<double> diff = __stop_time - __start_time;
+            std::cout << "iteration: " << __i << " __n_sorted = " << __n_sorted << ", time = " << diff.count() << " (s) " << std::endl;
+#endif
 
             __n_sorted *= 2;
             __data_in_temp = !__data_in_temp;
@@ -732,9 +754,21 @@ __merge_sort(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp, _LeafSo
 
     sycl::queue __q = __exec.queue();
 
+#if MERGE_SORT_DISPLAY_STATISTIC
+    const auto __start_time = std::chrono::high_resolution_clock::now();
+#endif
+
     // 1. Perform sorting of the leaves of the merge sort tree
     sycl::event __event_leaf_sort = __merge_sort_leaf_submitter<_LeafSortKernel>()(__q, __rng, __comp, __leaf_sorter);
 
+#if MERGE_SORT_DISPLAY_STATISTIC
+    __event_leaf_sort.wait_and_throw();
+    const auto __stop_time = std::chrono::high_resolution_clock::now();
+
+    const std::chrono::duration<double> diff = __stop_time - __start_time;
+    std::cout << "leaf sorter: time = " << diff.count() << " (s) " << std::endl;
+#endif
+
     // 2. Merge sorting
     oneapi::dpl::__par_backend_hetero::__buffer<_ExecutionPolicy, _Tp> __temp(__exec, __rng.size());
     auto __temp_buf = __temp.get_buffer();