Port device docs to rst

NVIDIA · Nov 29, 2023 · f27bc10 · f27bc10
1 parent f8fe8a9
commit f27bc10
Show file tree

Hide file tree

Showing 60 changed files with 9,022 additions and 9,511 deletions.
diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh
@@ -755,7 +755,7 @@ enum BlockLoadAlgorithm
 //!      using CUDA's built-in vectorized loads as a coalescing optimization.
 //!   #. :cpp:enumerator:`cub::BLOCK_LOAD_TRANSPOSE`:
 //!      A :ref:`striped arrangement <flexible-data-arrangement>` of data is read directly from memory and is then
-//!      locally transposed into a `blocked arrangement <flexible-data-arrangement>`.
+//!      locally transposed into a :ref:`blocked arrangement <flexible-data-arrangement>`.
 //!   #. :cpp:enumerator:`cub::BLOCK_LOAD_WARP_TRANSPOSE`:
 //!      A :ref:`warp-striped arrangement <flexible-data-arrangement>` of data is read directly from memory and is then
 //!      locally transposed into a :ref:`blocked arrangement <flexible-data-arrangement>`.

diff --git a/cub/cub/block/block_merge_sort.cuh b/cub/cub/block/block_merge_sort.cuh
@@ -682,7 +682,6 @@ private:
 /**
  * @brief The BlockMergeSort class provides methods for sorting items
  *        partitioned across a CUDA thread block using a merge sorting method.
- * @ingroup BlockModule
  *
  * @tparam KeyT
  *   KeyT type

diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh
@@ -276,7 +276,7 @@ private:
         BlockScan;
 
 
-    /// Shared memory storage layout type for BlockRadixRank
+    #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
     struct __align__(16) _TempStorage
     {
         union Aliasable
@@ -289,6 +289,7 @@ private:
         // Storage for scanning local ranks
         typename BlockScan::TempStorage block_scan;
     };
+    #endif // !DOXYGEN_SHOULD_SKIP_THIS
 
     /// Shared storage reference
     _TempStorage &temp_storage;
@@ -634,7 +635,7 @@ private:
         BlockScanT;
 
 
-    /// Shared memory storage layout type for BlockRadixRank
+    #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
     struct __align__(16) _TempStorage
     {
         typename BlockScanT::TempStorage            block_scan;
@@ -646,6 +647,7 @@ private:
 
         } aliasable;
     };
+    #endif // !DOXYGEN_SHOULD_SKIP_THIS
 
     /// Shared storage reference
     _TempStorage &temp_storage;
@@ -657,7 +659,7 @@ private:
 
 public:
 
-    /// @smemstorage{BlockScan}
+    /// @smemstorage{BlockRadixRankMatch}
     struct TempStorage : Uninitialized<_TempStorage> {};
 
 
@@ -957,9 +959,6 @@ struct BlockRadixRankMatchEarlyCounts
     // types
     typedef cub::BlockScan<int, BLOCK_THREADS, INNER_SCAN_ALGORITHM> BlockScan;
 
-
-
-    // temporary storage
     struct TempStorage
     {
         union

diff --git a/cub/cub/block/block_radix_sort.cuh b/cub/cub/block/block_radix_sort.cuh
@@ -182,8 +182,6 @@ CUB_NAMESPACE_BEGIN
 //! This example can be easily adapted to the storage required by BlockRadixSort.
 //! @endrst
 //!
-//! @ingroup BlockModule
-//!
 //! @tparam KeyT
 //!   KeyT type
 //!
@@ -2231,8 +2229,4 @@ public:
 
 };
 
-/**
- * \example example_block_radix_sort.cu
- */
-
 CUB_NAMESPACE_END
diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh
@@ -105,7 +105,7 @@ enum BlockReduceAlgorithm
     //!    single warp rake across segments of shared partial reductions.
     //! #. A warp-synchronous Kogge-Stone style reduction within the raking warp.
     //! 
-    //! @par Performance Considerations
+    //! Performance Considerations
     //! ++++++++++++++++++++++++++
     //!
     //! - This variant performs more communication than BLOCK_REDUCE_RAKING

diff --git a/cub/cub/block/block_run_length_decode.cuh b/cub/cub/block/block_run_length_decode.cuh
@@ -60,9 +60,9 @@ CUB_NAMESPACE_BEGIN
 //! the specified window will be returned.
 //! 
 //! .. note::
-//!   Trailing runs of length 0 are supported (i.e., they may only appear at the end of the run_lengths array). 
-//!   A run of length zero may not be followed by a run length that is not zero.
-//! 
+//!
+//!    Trailing runs of length 0 are supported (i.e., they may only appear at the end of the run_lengths array). 
+//!    A run of length zero may not be followed by a run length that is not zero.
 //!
 //! .. code-block:: c++
 //!

diff --git a/cub/cub/block/block_shuffle.cuh b/cub/cub/block/block_shuffle.cuh
@@ -185,7 +185,7 @@ public:
     //! 
     //! - @smemreuse
     //!
-    //! @rst
+    //! @endrst
     //! 
     //! @param[in] input
     //!   The calling thread's input item
@@ -311,7 +311,7 @@ public:
 
     //! @rst 
     //! The thread block rotates its :ref:`blocked arrangement <flexible-data-arrangement>` of input items,
-    //! shifting it down by one item. All threads receive ``input[0]` provided by *thread*\ :sub:`0`.
+    //! shifting it down by one item. All threads receive ``input[0]`` provided by *thread*\ :sub:`0`.
     //! 
     //! - @blocked
     //! - @granularity