From 3c863dad9146be24ccec93816c3cb0752d40d9ca Mon Sep 17 00:00:00 2001 From: Istvan Kiss Date: Thu, 12 Dec 2024 13:59:08 +0100 Subject: [PATCH] SWDEV-502543 - Update doxygen to surface funtions and Coop Group API Change-Id: Id4df63b8ae64a1113f85d89aa250ac9f7cc8b9bb --- .../amd_detail/amd_hip_cooperative_groups.h | 593 ++++++++++++++---- .../hip/amd_detail/amd_surface_functions.h | 145 ++++- 2 files changed, 599 insertions(+), 139 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 418f96949d..1901b49aa4 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -21,9 +21,9 @@ THE SOFTWARE. */ /** - * @file amd_detail/hip_cooperative_groups.h + * @file amd_detail/amd_hip_cooperative_groups.h * - * @brief Device side implementation of `Cooperative Group` feature. + * \brief Device side implementation of `Cooperative Group` feature. * * Defines new types and device API wrappers related to `Cooperative Group` * feature, which the programmer can directly use in his kernel(s) in order to @@ -39,26 +39,26 @@ THE SOFTWARE. namespace cooperative_groups { -/** @brief The base type of all cooperative group types +/** \brief The base type of all cooperative group types. * * \details Holds the key properties of a constructed cooperative group types - * object, like the group type, its size, etc + * object, like the group type, its size, etc. * - * @note Cooperative groups feature is implemented on Linux, under developement - * on Windows. + * \note Cooperative groups feature is implemented on Linux, under development + * on Microsoft Windows. */ class thread_group { protected: - uint32_t _type; // thread_group type - uint32_t _size; // total number of threads in the tread_group - uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types, - // LSB represents lane 0, and MSB represents lane 63 + uint32_t _type; //! Type of the thread_group. + uint32_t _size; //! Total number of threads in the tread_group. + uint64_t _mask; //! Lanemask for coalesced and tiled partitioned group types, + //! LSB represents lane 0, and MSB represents lane 63 - // Construct a thread group, and set thread group type and other essential - // thread group properties. This generic thread group is directly constructed - // only when the group is supposed to contain only the calling the thread - // (throurh the API - `this_thread()`), and in all other cases, this thread - // group object is a sub-object of some other derived thread group object + //! Construct a thread group, and set thread group type and other essential + //! thread group properties. This generic thread group is directly constructed + //! only when the group is supposed to contain only the calling the thread + //! (through the API - `this_thread()`), and in all other cases, this thread + //! group object is a sub-object of some other derived thread group object. __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size = static_cast(0), uint64_t mask = static_cast(0)) { _type = type; @@ -85,21 +85,32 @@ class thread_group { friend class thread_block; public: - // Total number of threads in the thread group, and this serves the purpose - // for all derived cooperative group types since their `size` is directly - // saved during the construction + //! Total number of threads in the thread_group, and this serves the purpose + //! for all derived cooperative group types because their `size` is directly + //! saved during the construction. __CG_QUALIFIER__ uint32_t size() const { return _size; } + //! Returns the type of the group. __CG_QUALIFIER__ unsigned int cg_type() const { return _type; } - // Rank of the calling thread within [0, size()) + //! Rank of the calling thread within [0, \link size() size() \endlink). __CG_QUALIFIER__ uint32_t thread_rank() const; - // Is this cooperative group type valid? + //! Returns true if the group has not violated any API constraints. __CG_QUALIFIER__ bool is_valid() const; - // synchronize the threads in the thread group + + /** \brief Synchronizes the threads in the group. + * + * \details Causes all threads in the group to wait at this synchronization point, + * and for all shared and global memory accesses by the threads to complete, + * before running synchronization. This guarantees the visibility of accessed data + * for all threads in the group. + * + * \note There are potential read-after-write (RAW), write-after-read (WAR), or + * write-after-write (WAW) hazards, when threads in the group access the + * same addresses in shared or global memory. The data hazards can + * be avoided with synchronization of the group. + */ __CG_QUALIFIER__ void sync() const; }; /** - *------------------------------------------------------------------------------------------------- - *------------------------------------------------------------------------------------------------- * @defgroup CooperativeG Cooperative Groups * @ingroup API * @{ @@ -108,100 +119,116 @@ class thread_group { * The cooperative groups provides flexible thread parallel programming algorithms, threads * cooperate and share data to perform collective computations. * - * @note Cooperative groups feature is implemented on Linux, under developement - * on Windows. + * \note Cooperative groups feature is implemented on Linux, under development + * on Microsoft Windows. * */ -/** \brief The multi-grid cooperative group type + +/** \brief The multi-grid cooperative group type. * - * \details Represents an inter-device cooperative group type where the - * participating threads within the group spans across multple - * devices, running the (same) kernel on these devices - * @note The multi-grid cooperative group type is implemented on Linux, under developement - * on Windows. + * \details Represents an inter-device cooperative group type, where the + * participating threads within the group span across multiple + * devices, running the (same) kernel on these devices. + * \note The multi-grid cooperative group type is implemented on Linux, + * under development on Microsoft Windows. */ class multi_grid_group : public thread_group { - // Only these friend functions are allowed to construct an object of this class - // and access its resources + //! Only these friend functions are allowed to construct an object of this class + //! and access its resources. friend __CG_QUALIFIER__ multi_grid_group this_multi_grid(); protected: - // Construct mutli-grid thread group (through the API this_multi_grid()) + //! Construct mutli-grid thread group (through the API this_multi_grid()) explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size) : thread_group(internal::cg_multi_grid, size) {} public: - // Number of invocations participating in this multi-grid group. In other - // words, the number of GPUs + + //! Number of invocations participating in this multi-grid group. In other + //! words, the number of GPUs. __CG_QUALIFIER__ uint32_t num_grids() { return internal::multi_grid::num_grids(); } - // Rank of this invocation. In other words, an ID number within the range - // [0, num_grids()) of the GPU, this kernel is running on + + //! Rank of this invocation. In other words, an ID number within the range + //! [0, num_grids()) of the GPU that kernel is running on. __CG_QUALIFIER__ uint32_t grid_rank() { return internal::multi_grid::grid_rank(); } + //! @copydoc thread_group::thread_rank __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); } + //! @copydoc thread_group::is_valid __CG_QUALIFIER__ bool is_valid() const { return internal::multi_grid::is_valid(); } + //! @copydoc thread_group::sync __CG_QUALIFIER__ void sync() const { internal::multi_grid::sync(); } }; -/** @brief User exposed API interface to construct multi-grid cooperative - * group type object - `multi_grid_group` +/** \addtogroup CooperativeGConstruct Construct functions of Cooperative groups + * \ingroup CooperativeG + * @{ */ + +/** \brief User-exposed API interface to construct grid cooperative group type + * object - `multi_grid_group`. * - * \details User is not allowed to directly construct an object of type - * `multi_grid_group`. Instead, he should construct it through this - * API function - * @note This multi-grid cooperative API type is implemented on Linux, under developement - * on Windows. + * \details User is not allowed to construct an object of type + * `multi_grid_group` directly. Instead, they should construct it + * through this API function. + * \note This multi-grid cooperative API type is implemented on Linux, under + * development on Microsoft Windows. */ __CG_QUALIFIER__ multi_grid_group this_multi_grid() { return multi_grid_group(internal::multi_grid::size()); } +// Doxygen end group CooperativeGConstruct +/** @} */ -/** @brief The grid cooperative group type +/** \brief The grid cooperative group type. * - * \details Represents an inter-workgroup cooperative group type where the + * \details Represents an inter-workgroup cooperative group type, where the * participating threads within the group spans across multiple - * workgroups running the (same) kernel on the same device - * @note This is implemented on Linux, under developement - * on Windows. + * workgroups running the (same) kernel on the same device. + * \note This is implemented on Linux and is under development + * on Microsoft Windows. */ class grid_group : public thread_group { - // Only these friend functions are allowed to construct an object of this class - // and access its resources + //! Only these friend functions are allowed to construct an object of this class + //! and access its resources. friend __CG_QUALIFIER__ grid_group this_grid(); protected: - // Construct grid thread group (through the API this_grid()) + //! Construct grid thread group (through the API this_grid()) explicit __CG_QUALIFIER__ grid_group(uint32_t size) : thread_group(internal::cg_grid, size) {} public: + //! @copydoc thread_group::thread_rank __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); } + //! @copydoc thread_group::is_valid __CG_QUALIFIER__ bool is_valid() const { return internal::grid::is_valid(); } + //! @copydoc thread_group::sync __CG_QUALIFIER__ void sync() const { internal::grid::sync(); } __CG_QUALIFIER__ dim3 group_dim() const { return internal::workgroup::block_dim(); } }; -/** @brief User exposed API interface to construct grid cooperative group type - * object - `grid_group` +/** \ingroup CooperativeGConstruct + * \brief User-exposed API interface to construct grid cooperative group type + * object - `grid_group`. * - * \details User is not allowed to directly construct an object of type - * `multi_grid_group`. Instead, he should construct it through this - * API function - * @note This function is implemented on Linux, under developement - * on Windows. + * \details User is not allowed to construct an object of type `grid_group` + * directly. Instead, they should construct it through this + * API function. + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ __CG_QUALIFIER__ grid_group this_grid() { return grid_group(internal::grid::size()); } -/** @brief The workgroup (thread-block in CUDA terminology) cooperative group - * type +/** \brief The workgroup (thread-block in CUDA terminology) cooperative group + * type. * - * \details Represents an intra-workgroup cooperative group type where the - * participating threads within the group are exactly the same threads - * which are participated in the currently executing `workgroup` - * @note This is implemented on Linux, under developement - * on Windows. + * \details Represents an intra-workgroup cooperative group type, where the + * participating threads within the group are the same threads that + * participated in the currently executing `workgroup`. + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ class thread_block : public thread_group { - // Only these friend functions are allowed to construct an object of thi - // class and access its resources + //! Only these friend functions are allowed to construct an object of thi + //! class and access its resources friend __CG_QUALIFIER__ thread_block this_thread_block(); friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size); @@ -234,25 +261,31 @@ class thread_block : public thread_group { } public: - // 3-dimensional block index within the grid + //! Returns 3-dimensional block index within the grid. __CG_STATIC_QUALIFIER__ dim3 group_index() { return internal::workgroup::group_index(); } - // 3-dimensional thread index within the block + //! Returns 3-dimensional thread index within the block. __CG_STATIC_QUALIFIER__ dim3 thread_index() { return internal::workgroup::thread_index(); } + //! @copydoc thread_group::thread_rank __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return internal::workgroup::thread_rank(); } + //! @copydoc thread_group::size __CG_STATIC_QUALIFIER__ uint32_t size() { return internal::workgroup::size(); } + //! @copydoc thread_group::is_valid __CG_STATIC_QUALIFIER__ bool is_valid() { return internal::workgroup::is_valid(); } + //! @copydoc thread_group::sync __CG_STATIC_QUALIFIER__ void sync() { internal::workgroup::sync(); } + //! Returns the group dimensions. __CG_QUALIFIER__ dim3 group_dim() { return internal::workgroup::block_dim(); } }; -/** \brief User exposed API interface to construct workgroup cooperative +/** \ingroup CooperativeGConstruct + * \brief User-exposed API interface to construct workgroup cooperative * group type object - `thread_block`. * - * \details User is not allowed to directly construct an object of type - * `thread_block`. Instead, he should construct it through this API + * \details User is not allowed to construct an object of type `thread_block` + * directly. Instead, they should construct it through this API * function. - * @note This function is implemented on Linux, under developement - * on Windows. + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ __CG_QUALIFIER__ thread_block this_thread_block() { return thread_block(internal::workgroup::size()); @@ -262,10 +295,9 @@ __CG_QUALIFIER__ thread_block this_thread_block() { * * \details Represents one tiled thread group in a wavefront. * This group type also supports sub-wave level intrinsics. - * @note This is implemented on Linux, under developement - * on Windows. + * \note This is implemented on Linux and is under development + * on Microsoft Windows. */ - class tiled_group : public thread_group { private: friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, @@ -297,12 +329,15 @@ class tiled_group : public thread_group { } public: + //! @copydoc thread_group::size __CG_QUALIFIER__ unsigned int size() const { return (coalesced_info.tiled_info.size); } + //! @copydoc thread_group::thread_rank __CG_QUALIFIER__ unsigned int thread_rank() const { return (internal::workgroup::thread_rank() & (coalesced_info.tiled_info.size - 1)); } + //! @copydoc thread_group::sync __CG_QUALIFIER__ void sync() const { internal::tiled_group::sync(); } @@ -312,10 +347,10 @@ template class thread_block_tile; /** \brief The coalesced_group cooperative group type * - * \details Represents a active thread group in a wavefront. + * \details Represents an active thread group in a wavefront. * This group type also supports sub-wave level intrinsics. - * @note This is implemented on Linux, under developement - * on Windows. + * \note This is implemented on Linux and is under development + * on Microsoft Windows. */ class coalesced_group : public thread_group { private: @@ -387,26 +422,44 @@ class coalesced_group : public thread_group { } public: + //! @copydoc thread_group::size __CG_QUALIFIER__ unsigned int size() const { return coalesced_info.size; } + //! @copydoc thread_group::thread_rank __CG_QUALIFIER__ unsigned int thread_rank() const { return internal::coalesced_group::masked_bit_count(coalesced_info.member_mask); - } + } + //! @copydoc thread_group::sync __CG_QUALIFIER__ void sync() const { internal::coalesced_group::sync(); } + //! Returns the linear rank of the group within the set of tiles partitioned + //! from a parent group (bounded by meta_group_size). __CG_QUALIFIER__ unsigned int meta_group_rank() const { return coalesced_info.tiled_info.meta_group_rank; } + //! Returns the number of groups created when the parent group was partitioned. __CG_QUALIFIER__ unsigned int meta_group_size() const { return coalesced_info.tiled_info.meta_group_size; } + /** \brief Shuffle operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle operation is a direct copy of ``var`` from ``srcRank`` + * thread ID of group. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. Only the srcRank thread ID of + * group is copied to other threads. + * \param srcRank [in] The source thread ID of the group for copy. + */ template __CG_QUALIFIER__ T shfl(T var, int srcRank) const { static_assert(is_valid_type::value, "Neither an integer or float type."); @@ -420,6 +473,20 @@ class coalesced_group : public thread_group { return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); } + /** \brief Shuffle down operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle down operation is copy of ``var`` from thread with + * thread ID of group relative higher with ``lane_delta`` to caller + * thread ID. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. + * \param lane_delta [in] The lane_delta is the relative thread ID difference + * between caller thread ID and source of copy thread + * ID. sourceID = (threadID + lane_delta) % size() + */ template __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const { static_assert(is_valid_type::value, "Neither an integer or float type."); @@ -447,6 +514,20 @@ class coalesced_group : public thread_group { return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); } + /** \brief Shuffle up operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle up operation is copy of ``var`` from thread with + * thread ID of group relative lower with ``lane_delta`` to caller + * thread ID. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. + * \param lane_delta [in] The lane_delta is the relative thread ID difference + * between caller thread ID and source of copy thread + * ID. sourceID = (threadID - lane_delta) % size() + */ template __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const { static_assert(is_valid_type::value, "Neither an integer or float type."); @@ -474,26 +555,65 @@ class coalesced_group : public thread_group { return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE); } #ifdef HIP_ENABLE_WARP_SYNC_BUILTINS + + /** \brief Ballot function on group level. + * + * \details Returns a bit mask with the Nth bit set to one if the specified + * predicate evaluates as true on the Nth thread. + * + * \param pred [in] The predicate to evaluate on group threads. + */ __CG_QUALIFIER__ unsigned long long ballot(int pred) const { return internal::helper::adjust_mask( coalesced_info.member_mask, __ballot_sync(coalesced_info.member_mask, pred)); } + /** \brief Any function on group level. + * + * \details Returns non-zero if a predicate evaluates true for any threads. + * + * \param pred [in] The predicate to evaluate on group threads. + */ __CG_QUALIFIER__ int any(int pred) const { return __any_sync(static_cast(coalesced_info.member_mask), pred); } + /** \brief All function on group level. + * + * \details Returns non-zero if a predicate evaluates true for all threads. + * + * \param pred [in] The predicate to evaluate on group threads. + */ __CG_QUALIFIER__ int all(int pred) const { return __all_sync(static_cast(coalesced_info.member_mask), pred); } + /** \brief Match any function on group level. + * + * \details Returns a bit mask containing a 1-bit for every participating + * thread if that thread has the same value in ``value`` as the + * caller thread. + * + * \param value [in] The value to examine on the current thread in group. + */ template __CG_QUALIFIER__ unsigned long long match_any(T value) const { return internal::helper::adjust_mask( coalesced_info.member_mask, __match_any_sync(static_cast(coalesced_info.member_mask), value)); } + /** \brief Match all function on group level. + * + * \details Returns a bit mask containing a 1-bit for every participating + * thread if they all have the same value in ``value`` as the caller + * thread. The predicate ``pred`` is set to true if all + * participating threads have the same value in ``value``. + * + * \param value [in] The value to examine on the current thread in group. + * \param pred [out] The predicate is set to true if all participating + * threads in the thread group have the same value. + */ template __CG_QUALIFIER__ unsigned long long match_all(T value, int& pred) const { return internal::helper::adjust_mask( coalesced_info.member_mask, @@ -503,21 +623,24 @@ class coalesced_group : public thread_group { #endif }; -/** \brief User exposed API to create coalesced groups. +/** \ingroup CooperativeGConstruct + * \brief User-exposed API to create coalesced groups. * - * \details A collective operation that groups all active lanes into a new thread group. - * @note This function is implemented on Linux, under developement - * on Windows. + * \details A collective operation that groups all active lanes into a new + * thread group. + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ - __CG_QUALIFIER__ coalesced_group coalesced_threads() { return cooperative_groups::coalesced_group(__builtin_amdgcn_read_exec()); } +#ifndef DOXYGEN_SHOULD_SKIP_THIS + /** - * Implemenation of all publicly exposed base class APIs - * @note This function is implemented on Linux, under developement - * on Windows. + * Implementation of all publicly exposed base class APIs + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { switch (this->_type) { @@ -542,10 +665,11 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { } } } + /** - * Implemenation of all publicly exposed thread group API - * @note This function is implemented on Linux, under developement - * on Windows. + * Implementation of all publicly exposed thread group API + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ __CG_QUALIFIER__ bool thread_group::is_valid() const { switch (this->_type) { @@ -570,10 +694,11 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const { } } } + /** - * Implemenation of all publicly exposed thread group sync API - * @note This function is implemented on Linux, under developement - * on Windows. + * Implementation of all publicly exposed thread group sync API + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ __CG_QUALIFIER__ void thread_group::sync() const { switch (this->_type) { @@ -603,58 +728,90 @@ __CG_QUALIFIER__ void thread_group::sync() const { } } -/** - * Implemenation of publicly exposed `wrapper` API on top of basic cooperative - * group type APIs - * @note This function is implemented on Linux, under developement - * on Windows. +#endif + +/** \addtogroup CooperativeGAPI User-exposed API of Cooperative groups + * \ingroup CooperativeG + * @{ */ + +/** \brief Returns the size of the group. + * + * \details Total number of threads in the thread group, and this serves the + * purpose for all derived cooperative group types because their + * `size` is directly saved during the construction. + * + * \tparam CGTy The cooperative group class template parameter. + * \param g [in] The cooperative group for size returns. + * + * \note Implementation of publicly exposed `wrapper` API on top of basic + * cooperative group type APIs. This function is implemented on Linux + * and is under development on Microsoft Windows. */ template __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); } -/** - * Implemenation of publicly exposed `wrapper` API on top of basic cooperative - * group type APIs - * @note This function is implemented on Linux, under developement - * on Windows. + +/** \brief Returns the rank of thread of the group. + * + * \details Rank of the calling thread within [0, \link size() size() \endlink). + * + * \tparam CGTy The cooperative group class template parameter. + * \param g [in] The cooperative group for rank returns. + * + * \note Implementation of publicly exposed `wrapper` API on top of basic + * cooperative group type APIs. This function is implemented on Linux + * and is under development on Microsoft Windows. */ template __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) { return g.thread_rank(); } -/** - * Implemenation of publicly exposed `wrapper` API on top of basic cooperative - * group type APIs - * @note This function is implemented on Linux, under developement - * on Windows. + +/** \brief Returns true if the group has not violated any API constraints. + * + * \tparam CGTy The cooperative group class template parameter. + * \param g [in] The cooperative group for validity check. + * + * \note Implementation of publicly exposed `wrapper` API on top of basic + * cooperative group type APIs. This function is implemented on Linux + * and is under development on Microsoft Windows. */ template __CG_QUALIFIER__ bool is_valid(CGTy const& g) { return g.is_valid(); } -/** - * Implemenation of publicly exposed `wrapper` API on top of basic cooperative - * group type APIs - * @note This function is implemented on Linux, under developement - * on Windows. + +/** \brief Synchronizes the threads in the group. + * + * \tparam CGTy The cooperative group class template parameter. + * \param g [in] The cooperative group for synchronization. + * + * \note Implementation of publicly exposed `wrapper` API on top of basic + * cooperative group type APIs. This function is implemented on Linux + * and is under development on Microsoft Windows. */ template __CG_QUALIFIER__ void sync(CGTy const& g) { g.sync(); } + +// Doxygen end group CooperativeGAPI +/** @} */ + /** * template class tile_base - * @note This class is implemented on Linux, under developement - * on Windows. + * \note This function is implemented on Linux and is under development + * on Microsoft Windows. */ template class tile_base { protected: _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize; public: - // Rank of the thread within this tile + //! Rank of the thread within this tile _CG_STATIC_CONST_DECL_ unsigned int thread_rank() { return (internal::workgroup::thread_rank() & (numThreads - 1)); } - // Number of threads within this tile + //! Number of threads within this tile __CG_STATIC_QUALIFIER__ unsigned int size() { return numThreads; } }; + /** * template class thread_block_tile_base - * @note This class is implemented on Linux, under developement - * on Windows. + * \note This class is implemented on Linux, under development + * on Microsoft Windows. */ template class thread_block_tile_base : public tile_base { static_assert(is_valid_tile_size::value, @@ -673,6 +830,7 @@ template class thread_block_tile_base : public tile_base class thread_block_tile_base : public tile_base class parent_group_info { public: - // Returns the linear rank of the group within the set of tiles partitioned - // from a parent group (bounded by meta_group_size) + //! Returns the linear rank of the group within the set of tiles partitioned + //! from a parent group (bounded by meta_group_size) __CG_STATIC_QUALIFIER__ unsigned int meta_group_rank() { return ParentCGTy::thread_rank() / tileSize; } - // Returns the number of groups created when the parent group was partitioned. + //! Returns the number of groups created when the parent group was partitioned. __CG_STATIC_QUALIFIER__ unsigned int meta_group_size() { return (ParentCGTy::size() + tileSize - 1) / tileSize; } @@ -739,8 +897,8 @@ public: /** \brief Group type - thread_block_tile * * \details Represents one tile of thread group. - * @note This type is implemented on Linux, under developement - * on Windows. + * \note This type is implemented on Linux, under development + * on Microsoft Windows. */ template class thread_block_tile_type : public thread_block_tile_base, @@ -783,14 +941,17 @@ class thread_block_tile_type : public thread_block_tile_base class thread_block_tile : public impl::thread_block_tile_internal { protected: @@ -862,8 +1032,132 @@ class thread_block_tile : public impl::thread_block_tile_internal() const { return thread_block_tile(*this); } -}; +#ifdef DOXYGEN_SHOULD_INCLUDE_THIS + + //! @copydoc thread_group::thread_rank + __CG_QUALIFIER__ unsigned int thread_rank() const; + + //! @copydoc thread_group::sync + __CG_QUALIFIER__ void sync(); + + //! Returns the linear rank of the group within the set of tiles partitioned + //! from a parent group (bounded by meta_group_size) + __CG_QUALIFIER__ unsigned int meta_group_rank() const; + + //! Returns the number of groups created when the parent group was partitioned. + __CG_QUALIFIER__ unsigned int meta_group_size() const; + + /** \brief Shuffle operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle operation is a direct copy of ``var`` from ``srcRank`` + * thread ID of group. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. Only the srcRank thread ID of + * group is copied to other threads. + * \param srcRank [in] The source thread ID of the group for copy. + */ + template __CG_QUALIFIER__ T shfl(T var, int srcRank) const; + + /** \brief Shuffle down operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle down operation is copy of ``var`` from thread with + * thread ID of group relative higher with ``lane_delta`` to caller + * thread ID. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. + * \param lane_delta [in] The lane_delta is the relative thread ID difference + * between caller thread ID and source of copy thread + * ID. sourceID = (threadID + lane_delta) % size() + */ + template __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const; + + /** \brief Shuffle up operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle up operation is copy of ``var`` from thread with + * thread ID of group relative lower with ``lane_delta`` to caller + * thread ID. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. + * \param lane_delta [in] The lane_delta is the relative thread ID difference + * between caller thread ID and source of copy thread + * ID. sourceID = (threadID - lane_delta) % size() + */ + template __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const; + + /** \brief Shuffle xor operation on group level. + * + * \details Exchanging variables between threads without use of shared memory. + * Shuffle xor operation is copy of var from thread with thread ID + * of group based on laneMask XOR of the caller thread ID. + * + * \tparam T The type can be a 32-bit integer or single-precision + * floating point. + * \param var [in] The source variable to copy. + * \param laneMask [in] The laneMask is the mask for XOR operation. + * sourceID = threadID ^ laneMask + */ + template __CG_QUALIFIER__ T shfl_xor(T var, unsigned int laneMask) const; + + /** \brief Ballot function on group level. + * + * \details Returns a bit mask with the Nth bit set to one if the Nth thread + * predicate evaluates true. + * + * \param pred [in] The predicate to evaluate on group threads. + */ + __CG_QUALIFIER__ unsigned long long ballot(int pred) const; + + /** \brief Any function on group level. + * + * \details Returns non-zero if a predicate evaluates true for any threads. + * + * \param pred [in] The predicate to evaluate on group threads. + */ + __CG_QUALIFIER__ int any(int pred) const; + + /** \brief All function on group level. + * + * \details Returns non-zero if a predicate evaluates true for all threads. + * + * \param pred [in] The predicate to evaluate on group threads. + */ + __CG_QUALIFIER__ int all(int pred) const; + + /** \brief Match any function on group level. + * + * \details Returns a bit mask containing a 1-bit for every participating + * thread if that thread has the same value in ``value`` as the + * caller thread. + * + * \param value [in] The value to examine on the current thread in group. + */ + template __CG_QUALIFIER__ unsigned long long match_any(T value) const; + + /** \brief Match all function on group level. + * + * \details Returns a bit mask containing a 1-bit for every participating + * thread if they all have the same value in ``value`` as the caller + * thread. The predicate ``pred`` is set to true if all + * participating threads have the same value in ``value``. + * + * \param value [in] The value to examine on the current thread in group. + * \param pred [out] The predicate is set to true if all participating + * threads in the thread group have the same value. + */ + template __CG_QUALIFIER__ unsigned long long match_all(T value, int& pred) const; + +#endif +}; template class thread_block_tile : public impl::thread_block_tile_internal { @@ -889,10 +1183,17 @@ struct tiled_partition_internal : public thread_block_tile __CG_QUALIFIER__ thread_block_tile tiled_partition(const ParentCGTy& g) { @@ -902,9 +1203,14 @@ __CG_QUALIFIER__ thread_block_tile tiled_partition(const Paren } #ifdef HIP_ENABLE_WARP_SYNC_BUILTINS -/** \brief Binary partition + +/** \ingroup CooperativeGConstruct + * \brief Binary partition. * - * \details This splits the input thread group into two partitions determined by predicate + * \details This splits the input thread group into two partitions determined by predicate. + * + * \param cgrp [in] The coalesced group for split. + * \param pred [in] The predicate used during the group split up. */ __CG_QUALIFIER__ coalesced_group binary_partition(const coalesced_group& cgrp, bool pred) { auto mask = __ballot_sync(cgrp.coalesced_info.member_mask, pred); @@ -916,6 +1222,17 @@ __CG_QUALIFIER__ coalesced_group binary_partition(const coalesced_group& cgrp, b } } +/** \ingroup CooperativeGConstruct + * \brief Binary partition. + * + * \details This splits the input thread group into two partitions determined by predicate. + * + * \tparam size The size of the input thread block tile group. + * \tparam parent The cooperative group class template parameter of the input group. + * + * \param tgrp [in] The thread block tile group for split. + * \param pred [in] The predicate used during the group split up. + */ template __CG_QUALIFIER__ coalesced_group binary_partition(const thread_block_tile& tgrp, bool pred) { diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h index 556988a8e7..a4919f0533 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -39,7 +39,12 @@ THE SOFTWARE. #endif #define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \ - unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj; + unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj; + +/** + * @defgroup SurfaceAPI Surface API + * @{ + */ // CUDA is using byte address, need map to pixel address for HIP static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) { @@ -96,6 +101,14 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order]; } +/** \brief Reads the value at coordinate x from the one-dimensional surface. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The coordinate where the value will be read out. + * \param boundaryMode [in] The boundary mode is currently ignored. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -107,6 +120,13 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the one-dimensional surface at coordinate x. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The coordinate where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -117,6 +137,15 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su __ockl_image_store_1D(i, x, tmp); } + +/** \brief Reads the value from the two-dimensional surface at coordinate x, y. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the value will be read out. + * \param y [in] The y coordinate where the value will be read out. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -127,6 +156,15 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the two-dimensional surface at coordinate + * x, y. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param y [in] The y coordinate where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -137,6 +175,16 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su __ockl_image_store_2D(i, int2(x, y).data, tmp); } +/** \brief Reads the value from the three-dimensional surface at coordinate + * x, y, z. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the value will be read out. + * \param y [in] The y coordinate where the value will be read out. + * \param z [in] The z coordinate where the value will be read out. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -147,6 +195,16 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the three-dimensional surface at coordinate + * x, y, z. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param y [in] The y coordinate where the data will be written. + * \param z [in] The z coordinate where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -157,6 +215,15 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp); } +/** \brief Reads the value from the one-dimensional layered surface at + * coordinate x and layer index. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The coordinate where the value will be read out. + * \param layer [in] The layer index where the value will be read out. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -167,6 +234,15 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the one-dimensional layered surface at + * coordinate x and layer index. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param layer [in] The layer index where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -177,6 +253,16 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje __ockl_image_store_lod_1D(i, x, layer, tmp); } +/** \brief Reads the value from the two-dimensional layered surface at + * coordinate x, y and layer index. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the value will be read out. + * \param y [in] The y coordinate where the value will be read out. + * \param layer [in] The layer index where the value will be read out. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -187,6 +273,16 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the two-dimensional layered surface at + * coordinate x, y and layer index. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param y [in] The y coordinate where the data will be written. + * \param layer [in] The layer index where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -197,6 +293,16 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp); } +/** \brief Reads the value from the cubemap surface at coordinate x, y and + * face index. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the value will be read out. + * \param y [in] The y coordinate where the value will be read out. + * \param face [in] The face index where the value will be read out. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -207,6 +313,16 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the cubemap surface at coordinate x, y and + * face index. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value is written to surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param y [in] The y coordinate where the data will be written. + * \param face [in] The face index where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -217,6 +333,17 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject __ockl_image_store_CM(i, int2(x, y).data, face, tmp); } +/** \brief Reads the value from the layered cubemap surface at coordinate x, y + * and face, layer index. + * + * \tparam T The data type of the surface. + * \param data [out] The T type result is stored in this pointer. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the value will be read out. + * \param y [in] The y coordinate where the value will be read out. + * \param face [in] The face index where the value will be read out. + * \param layer [in] The layer index where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -228,6 +355,17 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac *data = __hipMapFrom(tmp); } +/** \brief Writes the value data to the layered cubemap surface at coordinate + * x, y and face, layer index. + * + * \tparam T The data type of the surface. + * \param data [in] The T type value to write to the surface. + * \param surfObj [in] The surface descriptor. + * \param x [in] The x coordinate where the data will be written. + * \param y [in] The y coordinate where the data will be written. + * \param face [in] The face index where the data will be written. + * \param layer [in] The layer index where the data will be written. + */ template < typename T, typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> @@ -239,6 +377,11 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfa __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp); } +// Doxygen end group SurfaceAPI +/** +* @} +*/ + #endif #endif