Merge branch 'update_rocprim_docs' into 'master'

Update rocPRIM docs

See merge request amd/rocPRIM!176

Author: Jakub Szuppe

CMakeLists.txt

@@ -26,7 +26,7 @@ cmake_minimum_required(VERSION 3.5.1 FATAL_ERROR)
 set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "Install path prefix, prepended onto install directories")
 
 # rocPRIM project
-project(rocprim VERSION 0.3.2.0 LANGUAGES CXX)
+project(rocprim VERSION 1.0.0.0 LANGUAGES CXX)
 
 # CMake modules
 list(APPEND CMAKE_MODULE_PATH

README.md

@@ -1,7 +1,3 @@
-# Warning (pre-production state)
-
-rocPRIM is in its pre-production state and should be used for development purposes only.
-
 # rocPRIM
 
 The rocPRIM is a header-only library providing HIP and HC parallel primitives for developing

rocprim/include/rocprim/device/config_types.hpp

@@ -39,13 +39,13 @@ struct default_config { };
 /// \brief Configuration of particular kernels launched by device-level operation
 ///
 /// \tparam BlockSize - number of threads in a block.
-/// \tparam ItemsPerThread - number of items in processed by each thread.
+/// \tparam ItemsPerThread - number of items processed by each thread.
 template<unsigned int BlockSize, unsigned int ItemsPerThread>
 struct kernel_config
 {
-    /// value of BlockSize - number of threads in a block.
+    /// \brief Number of threads in a block.
     static constexpr unsigned int block_size = BlockSize;
-    /// value of ItemsPerThread - number of items in processed by each thread.
+    /// \brief Number of items processed by each thread.
     static constexpr unsigned int items_per_thread = ItemsPerThread;
 };
 

rocprim/include/rocprim/device/device_radix_sort_config.hpp

@@ -54,12 +54,14 @@ template<
 >
 struct radix_sort_config
 {
-#ifndef DOXYGEN_SHOULD_SKIP_THIS
+    /// \brief Number of bits in long iterations.
     static constexpr unsigned int long_radix_bits = LongRadixBits;
+    /// \brief Number of bits in short iterations.
     static constexpr unsigned int short_radix_bits = ShortRadixBits;
+    /// \brief Configuration of digits scan kernel.
     using scan = ScanConfig;
+    /// \brief Configuration of radix sort kernel.
     using sort = SortConfig;
-#endif
 };
 
 namespace detail

rocprim/include/rocprim/device/device_reduce_by_key_config.hpp

@@ -43,7 +43,9 @@ template<
 >
 struct reduce_by_key_config
 {
+    /// \brief Configuration of carry-outs scan kernel.
     using scan = ScanConfig;
+    /// \brief Configuration of the main reduce-by-key kernel.
     using reduce = ReduceConfig;
 };
 

rocprim/include/rocprim/device/device_reduce_config.hpp

@@ -38,7 +38,7 @@ BEGIN_ROCPRIM_NAMESPACE
 /// \brief Configuration of device-level reduce primitives.
 ///
 /// \tparam BlockSize - number of threads in a block.
-/// \tparam ItemsPerThread - number of items in processed by each thread.
+/// \tparam ItemsPerThread - number of items processed by each thread.
 /// \tparam BlockReduceMethod - algorithm for block reduce.
 template<
     unsigned int BlockSize,
@@ -47,9 +47,11 @@ template<
 >
 struct reduce_config
 {
+    /// \brief Number of threads in a block.
     static constexpr unsigned int block_size = BlockSize;
+    /// \brief Number of items processed by each thread.
     static constexpr unsigned int items_per_thread = ItemsPerThread;
-
+    /// \brief Algorithm for block reduce.
     static constexpr block_reduce_algorithm block_reduce_method = BlockReduceMethod;
 };
 

rocprim/include/rocprim/device/device_run_length_encode_config.hpp

@@ -45,7 +45,9 @@ template<
 >
 struct run_length_encode_config
 {
+    /// \brief Configuration of device-level reduce-by-key operation.
     using reduce_by_key = ReduceByKeyConfig;
+    /// \brief Configuration of device-level select operation.
     using select = SelectConfig;
 };
 

rocprim/include/rocprim/device/device_scan_config.hpp

@@ -40,7 +40,7 @@ BEGIN_ROCPRIM_NAMESPACE
 /// \brief Configuration of device-level scan primitives.
 ///
 /// \tparam BlockSize - number of threads in a block.
-/// \tparam ItemsPerThread - number of items in processed by each thread.
+/// \tparam ItemsPerThread - number of items processed by each thread.
 /// \tparam BlockLoadMethod - method for loading input values.
 /// \tparam StoreLoadMethod - method for storing values.
 /// \tparam BlockScanMethod - algorithm for block scan.
@@ -53,11 +53,15 @@ template<
 >
 struct scan_config
 {
+    /// \brief Number of threads in a block.
     static constexpr unsigned int block_size = BlockSize;
+    /// \brief Number of items processed by each thread.
     static constexpr unsigned int items_per_thread = ItemsPerThread;
-
+    /// \brief Method for loading input values.
     static constexpr block_load_method block_load_method = BlockLoadMethod;
+    /// \brief Method for storing values.
     static constexpr block_store_method block_store_method = BlockStoreMethod;
+    /// \brief Algorithm for block scan.
     static constexpr block_scan_algorithm block_scan_method = BlockScanMethod;
 };
 

rocprim/include/rocprim/device/device_segmented_radix_sort_config.hpp

@@ -52,11 +52,12 @@ template<
 >
 struct segmented_radix_sort_config
 {
-#ifndef DOXYGEN_SHOULD_SKIP_THIS
+    /// \brief Number of bits in long iterations.
     static constexpr unsigned int long_radix_bits = LongRadixBits;
+    /// \brief Number of bits in short iterations
     static constexpr unsigned int short_radix_bits = ShortRadixBits;
+    /// \brief Configuration of radix sort kernel.
     using sort = SortConfig;
-#endif
 };
 
 namespace detail

rocprim/include/rocprim/device/device_select_config.hpp

@@ -39,7 +39,7 @@ BEGIN_ROCPRIM_NAMESPACE
 /// \brief Configuration of device-level select operation.
 ///
 /// \tparam BlockSize - number of threads in a block.
-/// \tparam ItemsPerThread - number of items in processed by each thread.
+/// \tparam ItemsPerThread - number of items processed by each thread.
 /// \tparam ValueBlockLoadMethod - method for loading input values.
 /// \tparam FlagBlockLoadMethod - method for loading flag values.
 /// \tparam BlockScanMethod - algorithm for block scan.
@@ -52,12 +52,15 @@ template<
 >
 struct select_config
 {
+    /// \brief Number of threads in a block.
     static constexpr unsigned int block_size = BlockSize;
+    /// \brief Number of items processed by each thread.
     static constexpr unsigned int items_per_thread = ItemsPerThread;
-
+    /// \brief Method for loading input values.
     static constexpr block_load_method value_block_load_method = ValueBlockLoadMethod;
+    /// \brief Method for loading flag values.
     static constexpr block_load_method flag_block_load_method = FlagBlockLoadMethod;
-
+    /// \brief Algorithm for block scan.
     static constexpr block_scan_algorithm block_scan_method = BlockScanMethod;
 };
 

rocprim/include/rocprim/warp/warp_reduce.hpp

@@ -323,6 +323,25 @@ class warp_reduce
         base_type::reduce(input, output, valid_items, storage, reduce_op);
     }
 
+    /// \brief Performs head-segmented reduction across threads in a logical warp.
+    ///
+    /// \tparam Flag - type of head flags. Must be contextually convertible to \p bool.
+    /// \tparam BinaryFunction - type of binary function used for reduce. Default type
+    /// is rocprim::plus<T>.
+    ///
+    /// \param [in] input - thread input value.
+    /// \param [out] output - reference to a thread output value. May be aliased with \p input.
+    /// \param [in] flag - thread head flag, \p true flags mark beginnings of segments.
+    /// \param [in] storage - reference to a temporary storage object of type storage_type.
+    /// \param [in] reduce_op - binary operation function object that will be used for reduce.
+    /// The signature of the function should be equivalent to the following:
+    /// <tt>T f(const T &a, const T &b);</tt>. The signature does not need to have
+    /// <tt>const &</tt>, but function object must not modify the objects passed to it.
+    ///
+    /// \par Storage reusage
+    /// Synchronization barrier should be placed before \p storage is reused
+    /// or repurposed: \p __syncthreads() in HIP, \p tile_barrier::wait() in HC, or
+    /// universal rocprim::syncthreads().
     template<class Flag, class BinaryFunction = ::rocprim::plus<T>>
     ROCPRIM_DEVICE inline
     void head_segmented_reduce(T input,
@@ -334,6 +353,25 @@ class warp_reduce
         base_type::head_segmented_reduce(input, output, flag, storage, reduce_op);
     }
 
+    /// \brief Performs tail-segmented reduction across threads in a logical warp.
+    ///
+    /// \tparam Flag - type of tail flags. Must be contextually convertible to \p bool.
+    /// \tparam BinaryFunction - type of binary function used for reduce. Default type
+    /// is rocprim::plus<T>.
+    ///
+    /// \param [in] input - thread input value.
+    /// \param [out] output - reference to a thread output value. May be aliased with \p input.
+    /// \param [in] flag - thread tail flag, \p true flags mark ends of segments.
+    /// \param [in] storage - reference to a temporary storage object of type storage_type.
+    /// \param [in] reduce_op - binary operation function object that will be used for reduce.
+    /// The signature of the function should be equivalent to the following:
+    /// <tt>T f(const T &a, const T &b);</tt>. The signature does not need to have
+    /// <tt>const &</tt>, but function object must not modify the objects passed to it.
+    ///
+    /// \par Storage reusage
+    /// Synchronization barrier should be placed before \p storage is reused
+    /// or repurposed: \p __syncthreads() in HIP, \p tile_barrier::wait() in HC, or
+    /// universal rocprim::syncthreads().
     template<class Flag, class BinaryFunction = ::rocprim::plus<T>>
     ROCPRIM_DEVICE inline
     void tail_segmented_reduce(T input,

rocprim/include/rocprim/warp/warp_scan.hpp

@@ -734,12 +734,14 @@ class warp_scan
         return base_type::broadcast(input, src_lane, storage);
     }
 
+#ifndef DOXYGEN_SHOULD_SKIP_THIS
 protected:
     ROCPRIM_DEVICE inline
     void to_exclusive(T inclusive_input, T& exclusive_output, storage_type& storage)
     {
         return base_type::to_exclusive(inclusive_input, exclusive_output, storage);
     }
+#endif
 };
 
 END_ROCPRIM_NAMESPACE