   }
 
   /**
-   * \brief Sorts items using a merge sorting method.
-   *
-   * \par
+   * @brief Sorts items using a merge sorting method.
+   *
+   * @par
    * StableSortPairs is stable: it preserves the relative ordering of equivalent
    * elements. That is, if x and y are elements such that x precedes y,
    * and if the two elements are equivalent (neither x < y nor y < x) then
    * a postcondition of stable_sort is that x still precedes y.
    *
-   * \par Snippet
-   * The code snippet below illustrates the sorting of a device vector of \p int keys
-   * with associated vector of \p int values.
-   * \par
-   * \code
-   * #include <cub/cub.cuh>   // or equivalently <cub/device/device_merge_sort.cuh>
+   * @par Snippet
+   * The code snippet below illustrates the sorting of a device vector of `int`
+   * keys with associated vector of `int` values.
+   * @par
+   * @code
+   * #include <cub/cub.cuh>
+   * // or equivalently <cub/device/device_merge_sort.cuh>
    *
    * // Declare, allocate, and initialize device-accessible pointers for sorting data
    * int  num_items;       // e.g., 7

   * // Determine temporary device storage requirements
   * void     *d_temp_storage = NULL;
   * size_t   temp_storage_bytes = 0;
   * cub::DeviceMergeSort::StableSortKeys(d_temp_storage, temp_storage_bytes, d_keys, num_items, custom_op);
   *
   * // Allocate temporary storage
   * cudaMalloc(&d_temp_storage, temp_storage_bytes);
   *
   * // Run sorting operation
   * cub::DeviceMergeSort::StableSortKeys(d_temp_storage, temp_storage_bytes, d_keys, num_items, custom_op);
   *
   * // d_keys      <-- [0, 3, 5, 6, 7, 8, 9]
   *
   * \endcode
   *
   * \tparam KeyIteratorT is a model of <a href="https://en.cppreference.com/w/cpp/iterator/random_access_iterator">Random Access Iterator</a>,
   *         \p KeyIteratorT is mutable, and \p KeyIteratorT's \c value_type is
   *         a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>,
   *         and the ordering relation on \p KeyIteratorT's \c value_type is a <em>strict weak ordering</em>, as defined in the
   *         <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a> requirements.
   * \tparam OffsetT is an integer type for global offsets.
   * \tparam CompareOpT functor type having member <tt>bool operator()(KeyT lhs, KeyT rhs)</tt>
   *         CompareOpT is a model of <a href="https://en.cppreference.com/w/cpp/concepts/strict_weak_order">Strict Weak Ordering</a>.
   */
  template <typename KeyIteratorT,
            typename OffsetT,
            typename CompareOpT>
  CUB_RUNTIME_FUNCTION static cudaError_t
  StableSortKeys(void *d_temp_storage,            ///< [in] Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
                 std::size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation
                 KeyIteratorT d_keys,             ///< [in,out] Pointer to the input sequence of unsorted input keys
                 OffsetT num_items,               ///< [in] Number of items to sort
                 CompareOpT compare_op,           ///< [in] Comparison function object which returns true if the first argument is ordered before the second
                 cudaStream_t stream = 0,         ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
                 bool debug_synchronous = false)  ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors.  Also causes launch configurations to be printed to the console.  Default is \p false.
  {
    return SortKeys<KeyIteratorT, OffsetT, CompareOpT>(d_temp_storage,
                                                       temp_storage_bytes,

   * // Determine temporary device storage requirements
   * void     *d_temp_storage = NULL;
   * size_t   temp_storage_bytes = 0;
   * cub::DeviceMergeSort::SortKeys(
   *   d_temp_storage, temp_storage_bytes,
   *   d_keys, num_items, custom_op);
   *
   * // Allocate temporary storage
   * cudaMalloc(&d_temp_storage, temp_storage_bytes);
   *
   * // Run sorting operation
   * cub::DeviceMergeSort::SortKeys(
   *   d_temp_storage, temp_storage_bytes,
   *   d_keys, num_items, custom_op);
   *
   * // d_keys      <-- [0, 3, 5, 6, 7, 8, 9]
   * @endcode
   *
   * @tparam KeyIteratorT
   *   is a model of [Random Access Iterator], `KeyIteratorT` is mutable, and
   *   `KeyIteratorT`'s `value_type` is a model of [LessThan Comparable], and
   *   the ordering relation on `KeyIteratorT`'s `value_type` is a
   *   *strict weak ordering*, as defined in the [LessThan Comparable]
   *   requirements
   *
   * @tparam OffsetT
   *   is an integer type for global offsets
   *
   * @tparam CompareOpT
   *   functor type having member `bool operator()(KeyT lhs, KeyT rhs)`
   *   CompareOpT is a model of [Strict Weak Ordering]
   *
   * @param[in] d_temp_storage
   *   Device-accessible allocation of temporary storage. When `nullptr`, the
   *   required allocation size is written to `temp_storage_bytes` and no work
   *   is done
   *
   * @param[in,out] temp_storage_bytes
   *   Reference to size in bytes of `d_temp_storage` allocation
   *
   * @param[in,out] d_keys
   *   Pointer to the input sequence of unsorted input keys
   *
   * @param[in] num_items
   *   Number of items to sort
   *
   * @param[in] compare_op
   *   Comparison function object which returns true if the first argument is
   *   ordered before the second
   *
   * @param[in] stream
   *   **[optional]** CUDA stream to launch kernels within. Default is
   *   stream<sub>0</sub>
   *
   * @param[in] debug_synchronous
   *   **[optional]** Whether or not to synchronize the stream after every
   *   kernel launch to check for errors. Also causes launch configurations to
   *   be printed to the console. Default is `false`
   *
   * [Random Access Iterator]: https://en.cppreference.com/w/cpp/iterator/random_access_iterator
   * [Strict Weak Ordering]: https://en.cppreference.com/w/cpp/concepts/strict_weak_order
   * [LessThan Comparable] https://en.cppreference.com/w/cpp/named_req/LessThanComparable
   */
  template <typename KeyIteratorT,
            typename OffsetT,
            typename CompareOpT>
  CUB_RUNTIME_FUNCTION static cudaError_t
  SortKeys(void *d_temp_storage,
           std::size_t &temp_storage_bytes,
           KeyIteratorT d_keys,
           OffsetT num_items,
           CompareOpT compare_op,
           cudaStream_t stream    = 0,
           bool debug_synchronous = false)
  {
    using DispatchMergeSortT = DispatchMergeSort<KeyIteratorT,
                                                 NullType *,
