diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups.h index a5ae5f75b3..cd3e4bf085 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups.h @@ -39,15 +39,15 @@ namespace cooperative_groups { /** \brief The base type of all cooperative group types * - * \details Holds the key properties of a constructed cooperative group type + * \details Holds the key properties of a constructed cooperative group types * object, like the group type, its size, etc */ 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; // 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 // Construct a thread group, and set thread group type and other essential // thread group properties. This generic thread group is directly constructed @@ -61,13 +61,21 @@ class thread_group { _mask = mask; } + struct _tiled_info { + bool is_tiled; + unsigned int size; + } tiled_info; + + friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, + unsigned int tile_size); + 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 - __CG_QUALIFIER__ uint32_t size() const { - return _size; - } + __CG_QUALIFIER__ uint32_t size() const { return _size; } + __CG_QUALIFIER__ unsigned int cg_type() const { return _type; } // Rank of the calling thread within [0, size()) __CG_QUALIFIER__ uint32_t thread_rank() const; // Is this cooperative group type valid? @@ -90,28 +98,18 @@ class multi_grid_group : public thread_group { protected: // 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) { } + : thread_group(internal::cg_multi_grid, size) {} public: // 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(); - } + __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 - __CG_QUALIFIER__ uint32_t grid_rank() { - return internal::multi_grid::grid_rank(); - } - __CG_QUALIFIER__ uint32_t thread_rank() const { - return internal::multi_grid::thread_rank(); - } - __CG_QUALIFIER__ bool is_valid() const { - return internal::multi_grid::is_valid(); - } - __CG_QUALIFIER__ void sync() const { - internal::multi_grid::sync(); - } + __CG_QUALIFIER__ uint32_t grid_rank() { return internal::multi_grid::grid_rank(); } + __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); } + __CG_QUALIFIER__ bool is_valid() const { return internal::multi_grid::is_valid(); } + __CG_QUALIFIER__ void sync() const { internal::multi_grid::sync(); } }; /** \brief User exposed API interface to construct multi-grid cooperative @@ -121,8 +119,7 @@ class multi_grid_group : public thread_group { * `multi_grid_group`. Instead, he should construct it through this * API function */ -__CG_QUALIFIER__ multi_grid_group -this_multi_grid() { +__CG_QUALIFIER__ multi_grid_group this_multi_grid() { return multi_grid_group(internal::multi_grid::size()); } @@ -139,19 +136,12 @@ class grid_group : public thread_group { protected: // Construct grid thread group (through the API this_grid()) - explicit __CG_QUALIFIER__ grid_group(uint32_t size) - : thread_group(internal::cg_grid, size) { } + explicit __CG_QUALIFIER__ grid_group(uint32_t size) : thread_group(internal::cg_grid, size) {} public: - __CG_QUALIFIER__ uint32_t thread_rank() const { - return internal::grid::thread_rank(); - } - __CG_QUALIFIER__ bool is_valid() const { - return internal::grid::is_valid(); - } - __CG_QUALIFIER__ void sync() const { - internal::grid::sync(); - } + __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); } + __CG_QUALIFIER__ bool is_valid() const { return internal::grid::is_valid(); } + __CG_QUALIFIER__ void sync() const { internal::grid::sync(); } }; /** \brief User exposed API interface to construct grid cooperative group type @@ -161,60 +151,112 @@ class grid_group : public thread_group { * `multi_grid_group`. Instead, he should construct it through this * API function */ -__CG_QUALIFIER__ grid_group -this_grid() { - return grid_group(internal::grid::size()); -} +__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 exctly the same threads + * participating threads within the group are exactly the same threads * which are participated in the currently executing `workgroup` */ class thread_block : public thread_group { - // Only these friend functions are allowed to construct an object of this + // 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); + friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent, + unsigned int tile_size); protected: // Construct a workgroup thread group (through the API this_thread_block()) explicit __CG_QUALIFIER__ thread_block(uint32_t size) - : thread_group(internal::cg_workgroup, size) { } + : thread_group(internal::cg_workgroup, size) {} + + __CG_QUALIFIER__ thread_group new_tiled_group(unsigned int tile_size) const { + const bool pow2 = ((tile_size & (tile_size - 1)) == 0); + // Invalid tile size, assert + if (!tile_size || (tile_size > WAVEFRONT_SIZE) || !pow2) { + assert(false && "invalid tile size"); + } + + thread_group tiledGroup = thread_group(internal::cg_tiled_group, tile_size); + tiledGroup.tiled_info.size = tile_size; + tiledGroup.tiled_info.is_tiled = true; + return tiledGroup; + } public: // 3-dimensional block index within the grid - __CG_QUALIFIER__ dim3 group_index() { - return internal::workgroup::group_index(); - } + __CG_QUALIFIER__ dim3 group_index() { return internal::workgroup::group_index(); } // 3-dimensional thread index within the block - __CG_QUALIFIER__ dim3 thread_index() { - return internal::workgroup::thread_index(); - } - __CG_QUALIFIER__ uint32_t thread_rank() const { - return internal::workgroup::thread_rank(); - } - __CG_QUALIFIER__ bool is_valid() const { - return internal::workgroup::is_valid(); - } - __CG_QUALIFIER__ void sync() const { - internal::workgroup::sync(); - } + __CG_QUALIFIER__ dim3 thread_index() { return internal::workgroup::thread_index(); } + __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::workgroup::thread_rank(); } + __CG_QUALIFIER__ bool is_valid() const { return internal::workgroup::is_valid(); } + __CG_QUALIFIER__ void sync() const { internal::workgroup::sync(); } }; -/** \brief User exposed API interface to construct workgroup cooperative - * group type object - `thread_block` +/** \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 - * function + * function. */ -__CG_QUALIFIER__ thread_block -this_thread_block() { +__CG_QUALIFIER__ thread_block this_thread_block() { return thread_block(internal::workgroup::size()); } +/** \brief The tiled_group cooperative group type + * + * \details Represents one tiled thread group in a wavefront. + * This group type also supports sub-wave level intrinsics. + */ + +class tiled_group : public thread_group { + private: + friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, + unsigned int tile_size); + friend __CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent, + unsigned int tile_size); + + __CG_QUALIFIER__ tiled_group new_tiled_group(unsigned int tile_size) const { + const bool pow2 = ((tile_size & (tile_size - 1)) == 0); + + if (!tile_size || (tile_size > WAVEFRONT_SIZE) || !pow2) { + assert(false && "invalid tile size"); + } + + if (size() <= tile_size) { + return (*this); + } + + tiled_group tiledGroup = tiled_group(tile_size); + tiledGroup.tiled_info.is_tiled = true; + return tiledGroup; + } + + protected: + explicit __CG_QUALIFIER__ tiled_group(unsigned int tileSize) + : thread_group(internal::cg_tiled_group, tileSize) { + tiled_info.size = tileSize; + tiled_info.is_tiled = true; + } + + public: + __CG_QUALIFIER__ unsigned int size() const { return (tiled_info.size); } + + __CG_QUALIFIER__ unsigned int thread_rank() const { + return (internal::workgroup::thread_rank() & (tiled_info.size - 1)); + } + + __CG_QUALIFIER__ void sync() const { + // enforce memory ordering for memory instructions. + __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent"); + } +}; + /** * Implemenation of all publicly exposed base class APIs */ @@ -229,6 +271,9 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { case internal::cg_workgroup: { return (static_cast(this)->thread_rank()); } + case internal::cg_tiled_group: { + return (static_cast(this)->thread_rank()); + } default: { assert(false && "invalid cooperative group type"); return -1; @@ -247,6 +292,9 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const { case internal::cg_workgroup: { return (static_cast(this)->is_valid()); } + case internal::cg_tiled_group: { + return (static_cast(this)->is_valid()); + } default: { assert(false && "invalid cooperative group type"); return false; @@ -268,6 +316,10 @@ __CG_QUALIFIER__ void thread_group::sync() const { static_cast(this)->sync(); break; } + case internal::cg_tiled_group: { + static_cast(this)->sync(); + break; + } default: { assert(false && "invalid cooperative group type"); } @@ -278,27 +330,181 @@ __CG_QUALIFIER__ void thread_group::sync() const { * Implemenation of publicly exposed `wrapper` APIs on top of basic cooperative * group type APIs */ -template -__CG_QUALIFIER__ uint32_t group_size(CGTy const &g) { - return g.size(); -} +template __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); } -template -__CG_QUALIFIER__ uint32_t thread_rank(CGTy const &g) { +template __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) { return g.thread_rank(); } -template -__CG_QUALIFIER__ bool is_valid(CGTy const &g) { - return g.is_valid(); +template __CG_QUALIFIER__ bool is_valid(CGTy const& g) { return g.is_valid(); } + +template __CG_QUALIFIER__ void sync(CGTy const& g) { g.sync(); } + +template class tile_base { + protected: + _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize; + + public: + // 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 + __CG_STATIC_QUALIFIER__ unsigned int size() { return numThreads; } +}; + +template class thread_block_tile_base : public tile_base { + static_assert(is_valid_tile_size::value, + "Tile size is either not a power of 2 or greater than the wavefront size"); + using tile_base::numThreads; + + public: + __CG_STATIC_QUALIFIER__ void sync() { + // enforce ordering for memory instructions + __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent"); + } + + template __CG_QUALIFIER__ T shfl(T var, int srcRank) const { + static_assert(is_valid_type::value, "Neither an integer or float type."); + return (__shfl(var, srcRank, numThreads)); + } + + 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."); + return (__shfl_down(var, lane_delta, numThreads)); + } + + 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."); + return (__shfl_up(var, lane_delta, numThreads)); + } + + template __CG_QUALIFIER__ T shfl_xor(T var, unsigned int laneMask) const { + static_assert(is_valid_type::value, "Neither an integer or float type."); + return (__shfl_xor(var, laneMask, numThreads)); + } +}; + +/** \brief Group type - thread_block_tile + * + * \details Represents one tile of thread group. + */ + +template +class thread_block_tile_type : public thread_block_tile_base, public tiled_group { + _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize; + + friend class thread_block_tile_type; + + typedef thread_block_tile_base tbtBase; + + protected: + __CG_QUALIFIER__ thread_block_tile_type() : tiled_group(numThreads) { + tiled_info.size = numThreads; + tiled_info.is_tiled = true; + } + + public: + using tbtBase::size; + using tbtBase::sync; + using tbtBase::thread_rank; +}; + + +/** \brief User exposed API to partition groups. + * + * \details A collective operation that partitions the parent group into a one-dimensional, + * row-major, tiling of subgroups. + */ + +__CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size) { + if (parent.cg_type() == internal::cg_tiled_group) { + const tiled_group* cg = static_cast(&parent); + return cg->new_tiled_group(tile_size); + } else { + const thread_block* tb = static_cast(&parent); + return tb->new_tiled_group(tile_size); + } } -template -__CG_QUALIFIER__ void sync(CGTy const &g) { - g.sync(); +// Thread block type overload +__CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent, unsigned int tile_size) { + return (parent.new_tiled_group(tile_size)); } -} // namespace cooperative_groups +// Coalesced group type overload +__CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent, unsigned int tile_size) { + return (parent.new_tiled_group(tile_size)); +} -#endif // __cplusplus -#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H +template class thread_block_tile; + +namespace impl { +template class thread_block_tile_internal; + +template +class thread_block_tile_internal : public thread_block_tile_type { + protected: + template + __CG_QUALIFIER__ thread_block_tile_internal( + const thread_block_tile_internal& g) + : thread_block_tile_type() {} + + __CG_QUALIFIER__ thread_block_tile_internal(const thread_block& g) + : thread_block_tile_type() {} +}; +} // namespace impl + +template +class thread_block_tile : public impl::thread_block_tile_internal { + protected: + __CG_QUALIFIER__ thread_block_tile(const ParentCGTy& g) + : impl::thread_block_tile_internal(g) {} + + public: + __CG_QUALIFIER__ operator thread_block_tile() const { + return thread_block_tile(*this); + } +}; + + +template +class thread_block_tile : public impl::thread_block_tile_internal { + template friend class thread_block_tile; + + protected: + public: + template + __CG_QUALIFIER__ thread_block_tile(const thread_block_tile& g) + : impl::thread_block_tile_internal(g) {} +}; + +template class thread_block_tile; + +namespace impl { +template struct tiled_partition_internal; + +template +struct tiled_partition_internal : public thread_block_tile { + __CG_QUALIFIER__ tiled_partition_internal(const thread_block& g) + : thread_block_tile(g) {} +}; + +} // namespace impl + +/** \brief User exposed API to partition groups. + * + * \details This constructs a templated class derieved from thread_group. + * The template defines tile size of the new thread group at compile time. + */ +template +__CG_QUALIFIER__ thread_block_tile tiled_partition(const ParentCGTy& g) { + static_assert(is_valid_tile_size::value, + "Tiled partition with size > wavefront size. Currently not supported "); + return impl::tiled_partition_internal(g); +} +} // namespace cooperative_groups + +#endif // __cplusplus +#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index 6e71e4f1c3..90463485b6 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -47,12 +47,34 @@ THE SOFTWARE. #define __CG_STATIC_QUALIFIER__ __device__ static __forceinline__ #endif +#if !defined(_CG_STATIC_CONST_DECL_) +#define _CG_STATIC_CONST_DECL_ static constexpr +#endif + #if !defined(WAVEFRONT_SIZE) +#if __gfx1010__ || __gfx1011__ || __gfx1012__ || __gfx1030__ || __gfx1031__ +#define WAVEFRONT_SIZE 32 +#else #define WAVEFRONT_SIZE 64 #endif namespace cooperative_groups { +/* Global scope */ +template +using is_power_of_2 = std::integral_constant; + +template +using is_valid_wavefront = std::integral_constant; + +template +using is_valid_tile_size = + std::integral_constant::value && is_valid_wavefront::value>; + +template +using is_valid_type = + std::integral_constant::value || std::is_floating_point::value>; + namespace internal { /** \brief Enums representing different cooperative group types @@ -61,7 +83,8 @@ typedef enum { cg_invalid, cg_multi_grid, cg_grid, - cg_workgroup + cg_workgroup, + cg_tiled_group } group_type; /** @@ -69,31 +92,19 @@ typedef enum { */ namespace multi_grid { -__CG_STATIC_QUALIFIER__ uint32_t num_grids() { - return (uint32_t)__ockl_multi_grid_num_grids(); -} +__CG_STATIC_QUALIFIER__ uint32_t num_grids() { return (uint32_t)__ockl_multi_grid_num_grids(); } -__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { - return (uint32_t)__ockl_multi_grid_grid_rank(); -} +__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { return (uint32_t)__ockl_multi_grid_grid_rank(); } -__CG_STATIC_QUALIFIER__ uint32_t size() { - return (uint32_t)__ockl_multi_grid_size(); -} +__CG_STATIC_QUALIFIER__ uint32_t size() { return (uint32_t)__ockl_multi_grid_size(); } -__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - return (uint32_t)__ockl_multi_grid_thread_rank(); -} +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return (uint32_t)__ockl_multi_grid_thread_rank(); } -__CG_STATIC_QUALIFIER__ bool is_valid() { - return (bool)__ockl_multi_grid_is_valid(); -} +__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_multi_grid_is_valid(); } -__CG_STATIC_QUALIFIER__ void sync() { - __ockl_multi_grid_sync(); -} +__CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); } -} // namespace multi_grid +} // namespace multi_grid /** * Functionalities related to grid cooperative group type @@ -101,41 +112,32 @@ __CG_STATIC_QUALIFIER__ void sync() { namespace grid { __CG_STATIC_QUALIFIER__ uint32_t size() { - return (uint32_t)((hipBlockDim_z * hipGridDim_z) * - (hipBlockDim_y * hipGridDim_y) * + return (uint32_t)((hipBlockDim_z * hipGridDim_z) * (hipBlockDim_y * hipGridDim_y) * (hipBlockDim_x * hipGridDim_x)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute global id of the workgroup to which the current thread belongs to - uint32_t blkIdx = - (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) + - (hipBlockIdx_y * hipGridDim_x) + - (hipBlockIdx_x)); + uint32_t blkIdx = (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) + + (hipBlockIdx_y * hipGridDim_x) + (hipBlockIdx_x)); // Compute total number of threads being passed to reach current workgroup // within grid uint32_t num_threads_till_current_workgroup = - (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); + (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); // Compute thread local rank within current workgroup - uint32_t local_thread_rank = - (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + - (hipThreadIdx_y * hipBlockDim_x) + - (hipThreadIdx_x)); + uint32_t local_thread_rank = (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + + (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x)); return (num_threads_till_current_workgroup + local_thread_rank); } -__CG_STATIC_QUALIFIER__ bool is_valid() { - return (bool)__ockl_grid_is_valid(); -} +__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_grid_is_valid(); } -__CG_STATIC_QUALIFIER__ void sync() { - __ockl_grid_sync(); -} +__CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); } -} // namespace grid +} // namespace grid /** * Functionalities related to `workgroup` (thread_block in CUDA terminology) @@ -144,39 +146,35 @@ __CG_STATIC_QUALIFIER__ void sync() { namespace workgroup { __CG_STATIC_QUALIFIER__ dim3 group_index() { - return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y, - (uint32_t)hipBlockIdx_z)); + return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y, (uint32_t)hipBlockIdx_z)); } __CG_STATIC_QUALIFIER__ dim3 thread_index() { - return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y, - (uint32_t)hipThreadIdx_z)); + return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y, (uint32_t)hipThreadIdx_z)); } __CG_STATIC_QUALIFIER__ uint32_t size() { - return((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); + return ((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + - (hipThreadIdx_y * hipBlockDim_x) + - (hipThreadIdx_x))); + return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + + (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x))); } __CG_STATIC_QUALIFIER__ bool is_valid() { - //TODO(mahesha) any functionality need to be added here? I believe not + // TODO(mahesha) any functionality need to be added here? I believe not return true; } -__CG_STATIC_QUALIFIER__ void sync() { - __syncthreads(); -} +__CG_STATIC_QUALIFIER__ void sync() { __syncthreads(); } -} // namespace workgroup +} // namespace workgroup -} // namespace internal +} // namespace internal -} // namespace cooperative_groups +} // namespace cooperative_groups -#endif // __cplusplus -#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H +#endif // __cplusplus +#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H +#endif diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_tiled_partition.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_tiled_partition.cpp new file mode 100644 index 0000000000..d9397da211 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_tiled_partition.cpp @@ -0,0 +1,400 @@ +/* +Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// Test Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. The second part, partitions this parent group into child subgroups + a.k.a tiles using using tiled_partition() collective operation. This can be called + with a static tile size, passed in templated non-type variable-tiled_partition, + or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these + cases. + This test tests functionality of cg group partitioning, (static and dynamic) and its respective + API's size(), thread_rank(), and sync(). +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +/* Parallel reduce kernel. + * + * Step complexity: O(log n) + * Work complexity: O(n) + * + * Note: This kernel works only with power of 2 input arrays. + */ +__device__ int reduction_kernel(thread_group g, int* x, int val) { + int lane = g.thread_rank(); + int sz = g.size(); + + for (int i = g.size() / 2; i > 0; i /= 2) { + // use lds to store the temporary result + x[lane] = val; + // Ensure all the stores are completed. + g.sync(); + + if (lane < i) { + val += x[lane + i]; + } + // It must work on one tiled thread group at a time, + // and it must make sure all memory operations are + // completed before moving to the next stride. + // sync() here just does that. + g.sync(); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +template +__global__ void kernel_cg_group_partition_static(int* result, bool isGlobalMem, int* globalMem) { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + int* workspace = NULL; + + if (isGlobalMem) { + workspace = globalMem; + } else { + // Declare a shared memory + extern __shared__ int sharedMem[]; + workspace = sharedMem; + } + + int input, outputSum, expectedOutput; + + // we pass its own thread rank as inputs + input = threadBlockCGTy.thread_rank(); + + expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2; + + outputSum = reduction_kernel(threadBlockCGTy, workspace, input); + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d (expected %d)\n\n", + (int)threadBlockCGTy.size() - 1, outputSum, expectedOutput); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + thread_block_tile tiledPartition = tiled_partition(threadBlockCGTy); + + // This offset allows each group to have its own unique area in the workspace array + int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank(); + + outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread " + "rank: %d\n", + tiledPartition.size() - 1, outputSum, input); + result[input / (tileSz)] = outputSum; + } + return; +} + + +__global__ void kernel_cg_group_partition_dynamic(unsigned int tileSz, int* result, + bool isGlobalMem, int* globalMem) { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + int* workspace = NULL; + + if (isGlobalMem) { + workspace = globalMem; + } else { + // Declare a shared memory + extern __shared__ int sharedMem[]; + workspace = sharedMem; + } + + int input, outputSum, expectedOutput; + + // input to reduction, for each thread, is its' rank in the group + input = threadBlockCGTy.thread_rank(); + + expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2; + + outputSum = reduction_kernel(threadBlockCGTy, workspace, input); + + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d\n\n", + (int)threadBlockCGTy.size() - 1, outputSum); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + thread_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + + // This offset allows each group to have its own unique area in the workspace array + int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank(); + + outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread " + "rank: %d\n", + tiledPartition.size() - 1, outputSum, input); + + result[input / (tileSz)] = outputSum; + } + return; +} + +// Search if the sum exists in the expected results array +void verifyResults(int* hPtr, int* dPtr, int size) { + int i = 0, j = 0; + for (i = 0; i < size; i++) { + for (j = 0; j < size; j++) { + if (hPtr[i] == dPtr[j]) { + break; + } + } + if (j == size) { + failed(" Result verification failed!"); + } + } +} + + +template static void test_group_partition(bool useGlobalMem) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 64; + + int numTiles = (blockSize * threadsPerBlock) / tileSz; + + // Build an array of expected reduction sum output on the host + // based on the sum of their respective thread ranks for verification. + // eg: parent group has 64threads. + // child thread ranks: 0-15, 16-31, 32-47, 48-63 + // expected sum: 120, 376, 632, 888 + int* expectedSum = new int[numTiles]; + int temp = 0, sum = 0; + + for (int i = 1; i <= numTiles; i++) { + sum = temp; + temp = (((tileSz * i) - 1) * (tileSz * i)) / 2; + expectedSum[i-1] = temp - sum; + } + + int* dResult = NULL; + hipMalloc((void**)&dResult, numTiles * sizeof(int)); + + int* globalMem = NULL; + if (useGlobalMem) { + hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int)); + } + + int* hResult = NULL; + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + if (useGlobalMem) { + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, blockSize, threadsPerBlock, 0, 0, + dResult, useGlobalMem, globalMem); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } else { + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dResult, useGlobalMem, globalMem); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } + + hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost); + + verifyResults(expectedSum, hResult, numTiles); + + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + if (useGlobalMem) { + hipFree(globalMem); + } + delete[] expectedSum; + + printf("\n...PASSED.\n\n"); +} + +static void test_group_partition(unsigned int tileSz, bool useGlobalMem) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 64; + + int numTiles = (blockSize * threadsPerBlock) / tileSz; + // Build an array of expected reduction sum output on the host + // based on the sum of their respective thread ranks to use for verification + int* expectedSum = new int[numTiles]; + int temp = 0, sum = 0; + for (int i = 1; i <= numTiles; i++) { + sum = temp; + temp = (((tileSz * i) - 1) * (tileSz * i)) / 2; + expectedSum[i-1] = temp - sum; + } + + int* dResult = NULL; + hipMalloc(&dResult, sizeof(int) * numTiles); + + int* globalMem = NULL; + if (useGlobalMem) { + hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int)); + } + + int* hResult = NULL; + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + // Launch Kernel + if (useGlobalMem) { + hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock, 0, 0, tileSz, + dResult, useGlobalMem, globalMem); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } else { + hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } + + hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost); + + verifyResults(expectedSum, hResult, numTiles); + + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + if (useGlobalMem) { + hipFree(globalMem); + } + delete[] expectedSum; + + printf("\n...PASSED.\n\n"); +} + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + } + + bool useGlobalMem = true; + std::cout << "Testing static tiled_partition for different tile sizes" << std::endl; + std::cout << "\nUsing global memory for computation\n"; + /* Test static tile_partition */ + std::cout << "TEST 1:" << '\n' << std::endl; + test_group_partition<2>(useGlobalMem); + std::cout << "TEST 2:" << '\n' << std::endl; + test_group_partition<4>(useGlobalMem); + std::cout << "TEST 3:" << '\n' << std::endl; + test_group_partition<8>(useGlobalMem); + std::cout << "TEST 4:" << '\n' << std::endl; + test_group_partition<16>(useGlobalMem); + std::cout << "TEST 5:" << '\n' << std::endl; + test_group_partition<32>(useGlobalMem); + + useGlobalMem = false; + std::cout << "Testing static tiled_partition for different tile sizes" << std::endl; + std::cout << "\nUsing shared memory for computation\n"; + /* Test static tile_partition */ + std::cout << "TEST 1:" << '\n' << std::endl; + test_group_partition<2>(useGlobalMem); + std::cout << "TEST 2:" << '\n' << std::endl; + test_group_partition<4>(useGlobalMem); + std::cout << "TEST 3:" << '\n' << std::endl; + test_group_partition<8>(useGlobalMem); + std::cout << "TEST 4:" << '\n' << std::endl; + test_group_partition<16>(useGlobalMem); + std::cout << "TEST 5:" << '\n' << std::endl; + test_group_partition<32>(useGlobalMem); + + + std::cout << "Now testing dynamic tiled_partition for different tile sizes" << '\n' << std::endl; + + /* Test dynamic group partition*/ + useGlobalMem = true; + int testNo = 1; + std::vector tileSizes = {2, 4, 8, 16, 32}; + std::cout << "\nUsing global memory for computation\n"; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i, useGlobalMem); + testNo++; + } + + useGlobalMem = false; + testNo = 1; + std::cout << "\nUsing shared memory for computation\n"; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i, useGlobalMem); + testNo++; + } + + passed(); + return 0; +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tile_shfl_ops.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tile_shfl_ops.cpp new file mode 100644 index 0000000000..002401968a --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tile_shfl_ops.cpp @@ -0,0 +1,231 @@ +/* +Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. + This sample tests functionality of intrinsics provided by thread_block_tile type, + shfl_down and shfl_xor. +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +template +__device__ int reduction_kernel_shfl_down(thread_block_tile const& g, volatile int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_down(val, i); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +template +__device__ int reduction_kernel_shfl_xor(thread_block_tile const& g, int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_xor(val, i); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +template +__global__ void kernel_cg_group_partition_static(int* result, bool runShflDown) { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + int input, outputSum, expectedSum; + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + thread_block_tile tiledPartition = tiled_partition(threadBlockCGTy); + int threadRank = tiledPartition.thread_rank(); + + input = tiledPartition.thread_rank(); + + // (n-1)(n)/2 + expectedSum = ((tileSz - 1) * tileSz / 2); + + if (runShflDown) { + outputSum = reduction_kernel_shfl_down(tiledPartition, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected " + "%d)\n", + tiledPartition.size() - 1, outputSum, expectedSum); + result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum; + } + } else { + outputSum = reduction_kernel_shfl_xor(tiledPartition, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group using shfl_xor is %d (expected " + "%d)\n", + tiledPartition.size() - 1, outputSum, expectedSum); + result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum; + } + } + + return; +} + +void verifyResults(int* ptr, int expectedResult, int numTiles) { + for (int i = 0; i < numTiles; i++) { + if (ptr[i] != expectedResult) { + failed(" Results do not match! "); + } + } +} + +template static void test_group_partition(bool runShflDown) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 64; + + int numTiles = (blockSize * threadsPerBlock) / tileSz; + int expectedSum = ((tileSz - 1) * tileSz / 2); + int* expectedResult = new int[numTiles]; + + for (int i = 0; i < numTiles; i++) { + expectedResult[i] = expectedSum; + } + + int* dResult = NULL; + int* hResult = NULL; + + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + hipMalloc(&dResult, numTiles * sizeof(int)); + + if (runShflDown) { + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dResult, runShflDown); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } else { + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dResult, runShflDown); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } + + hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost); + + verifyResults(hResult, expectedSum, numTiles); + + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + delete[] expectedResult; + + printf("\n...PASSED.\n\n"); +} + + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + + bool runShflDown = true; + std::cout << "Testing static tiled_partition for different tile sizes using shfl_down" + << std::endl; + /* Test static tile_partition */ + std::cout << "TEST 1:" << '\n' << std::endl; + test_group_partition<2>(runShflDown); + std::cout << "TEST 2:" << '\n' << std::endl; + test_group_partition<4>(runShflDown); + std::cout << "TEST 3:" << '\n' << std::endl; + test_group_partition<8>(runShflDown); + std::cout << "TEST 4:" << '\n' << std::endl; + test_group_partition<16>(runShflDown); + std::cout << "TEST 5:" << '\n' << std::endl; + test_group_partition<32>(runShflDown); + + runShflDown = false; + std::cout << "Testing static tiled_partition for different tile sizes using shfl_xor" + << std::endl; + /* Test static tile_partition */ + std::cout << "TEST 1:" << '\n' << std::endl; + test_group_partition<2>(runShflDown); + std::cout << "TEST 2:" << '\n' << std::endl; + test_group_partition<4>(runShflDown); + std::cout << "TEST 3:" << '\n' << std::endl; + test_group_partition<8>(runShflDown); + std::cout << "TEST 4:" << '\n' << std::endl; + test_group_partition<16>(runShflDown); + std::cout << "TEST 5:" << '\n' << std::endl; + test_group_partition<32>(runShflDown); + + passed(); +} diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tiled_shfl_up.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tiled_shfl_up.cpp new file mode 100644 index 0000000000..06e3b5da09 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tiled_shfl_up.cpp @@ -0,0 +1,179 @@ +/* +Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/* This test implements prefix sum(scan) kernel, first with each threads own rank + as input and comparing the sum with expected serial summation output on CPU. + + This sample tests functionality of intrinsics provided by thread_block_tile type, + shfl_up. +*/ +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +template +__device__ int prefix_sum_kernel(thread_block_tile const& g, volatile int val) { + int sz = g.size(); +#pragma unroll + for (int i = 1; i < sz; i <<= 1) { + int temp = g.shfl_up(val, i); + + if (g.thread_rank() >= i) { + val += temp; + } + } + return val; +} + +template __global__ void kernel_cg_group_partition_static(int* dPtr) { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + int input, outputSum; + + // we pass its own thread rank as inputs + input = threadBlockCGTy.thread_rank(); + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + thread_block_tile tiledPartition = tiled_partition(threadBlockCGTy); + + input = tiledPartition.thread_rank(); + + outputSum = prefix_sum_kernel(tiledPartition, input); + + // Update the result array with the corresponsing prefix sum + dPtr[threadBlockCGTy.thread_rank()] = outputSum; + return; +} + +void serialScan(int* ptr, int size) { + // Fill up the array + for (int i = 0; i < size; i++) { + ptr[i] = i; + } + + int acc = 0; + for (int i = 0; i < size; i++) { + acc = acc + ptr[i]; + ptr[i] = acc; + } +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +void verifyResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + failed(" Prefix sum results do not match."); + } + } +} + +template static void test_group_partition() { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 64; + + int* hPtr = NULL; + int* dPtr = NULL; + int* cpuPrefixSum = NULL; + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + hipMalloc(&dPtr, arrSize); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr); + hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + + cpuPrefixSum = new int[tileSz]; + serialScan(cpuPrefixSum, tileSz); + std::cout << "\nPrefix sum results on CPU\n"; + printResults(cpuPrefixSum, tileSz); + + std::cout << "\nPrefix sum results on GPU\n"; + printResults(hPtr, tileSz); + std::cout << "\n"; + verifyResults(hPtr, cpuPrefixSum, tileSz); + std::cout << "Results verified!\n"; + + delete[] cpuPrefixSum; + hipFree(hPtr); + hipFree(dPtr); +} + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + std::cout << "Testing static tiled_partition for different tile sizes" << std::endl; + /* Test static tile_partition */ + std::cout << "TEST 1:" << '\n' << std::endl; + test_group_partition<2>(); + std::cout << "TEST 2:" << '\n' << std::endl; + test_group_partition<4>(); + std::cout << "TEST 3:" << '\n' << std::endl; + test_group_partition<8>(); + std::cout << "TEST 4:" << '\n' << std::endl; + test_group_partition<16>(); + std::cout << "TEST 5:" << '\n' << std::endl; + test_group_partition<32>(); + passed(); +} + +/* Kogge-Stone algorithm */ \ No newline at end of file