SWDEV-450109 - add missing coop group operations
any, all, ballot, match_any, match_all for coalesced_group and tiled_group. add binary_partition. Change-Id: I8ed0e654f36d788653352a0617beb5b8d44749bc
Этот коммит содержится в:
Jatin Chaudhary
коммит произвёл
Jatin Jaikishan Chaudhary
Jatin Jaikishan Chaudhary
родитель
1566ff7639
Коммит
36d58c6518
@@ -79,6 +79,7 @@ class thread_group {
|
||||
struct _tiled_info tiled_info;
|
||||
} coalesced_info;
|
||||
|
||||
friend __CG_QUALIFIER__ thread_group this_thread();
|
||||
friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
|
||||
unsigned int tile_size);
|
||||
friend class thread_block;
|
||||
@@ -175,6 +176,7 @@ class grid_group : public thread_group {
|
||||
__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__ dim3 group_dim() const { return internal::workgroup::block_dim(); }
|
||||
};
|
||||
|
||||
/** @brief User exposed API interface to construct grid cooperative group type
|
||||
@@ -306,6 +308,8 @@ class tiled_group : public thread_group {
|
||||
}
|
||||
};
|
||||
|
||||
template <unsigned int size, class ParentCGTy> class thread_block_tile;
|
||||
|
||||
/** \brief The coalesced_group cooperative group type
|
||||
*
|
||||
* \details Represents a active thread group in a wavefront.
|
||||
@@ -318,6 +322,10 @@ class coalesced_group : public thread_group {
|
||||
friend __CG_QUALIFIER__ coalesced_group coalesced_threads();
|
||||
friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size);
|
||||
friend __CG_QUALIFIER__ coalesced_group tiled_partition(const coalesced_group& parent, unsigned int tile_size);
|
||||
friend __CG_QUALIFIER__ coalesced_group binary_partition(const coalesced_group& cgrp, bool pred);
|
||||
template <unsigned int fsize, class fparent>
|
||||
friend __CG_QUALIFIER__ coalesced_group
|
||||
binary_partition(const thread_block_tile<fsize, fparent>& tgrp, bool pred);
|
||||
|
||||
__CG_QUALIFIER__ coalesced_group new_tiled_group(unsigned int tile_size) const {
|
||||
const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
|
||||
@@ -465,6 +473,34 @@ class coalesced_group : public thread_group {
|
||||
|
||||
return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE);
|
||||
}
|
||||
#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
__CG_QUALIFIER__ unsigned long long ballot(int pred) const {
|
||||
return internal::helper::adjust_mask(
|
||||
coalesced_info.member_mask,
|
||||
__ballot_sync<unsigned long long>(coalesced_info.member_mask, pred));
|
||||
}
|
||||
|
||||
__CG_QUALIFIER__ int any(int pred) const {
|
||||
return __any_sync(static_cast<unsigned long long>(coalesced_info.member_mask), pred);
|
||||
}
|
||||
|
||||
__CG_QUALIFIER__ int all(int pred) const {
|
||||
return __all_sync(static_cast<unsigned long long>(coalesced_info.member_mask), pred);
|
||||
}
|
||||
|
||||
template <typename T> __CG_QUALIFIER__ unsigned long long match_any(T value) const {
|
||||
return internal::helper::adjust_mask(
|
||||
coalesced_info.member_mask,
|
||||
__match_any_sync(static_cast<unsigned long long>(coalesced_info.member_mask), value));
|
||||
}
|
||||
|
||||
template <typename T> __CG_QUALIFIER__ unsigned long long match_all(T value, int& pred) const {
|
||||
return internal::helper::adjust_mask(
|
||||
coalesced_info.member_mask,
|
||||
__match_all_sync(static_cast<unsigned long long>(coalesced_info.member_mask), value,
|
||||
&pred));
|
||||
}
|
||||
#endif
|
||||
};
|
||||
|
||||
/** \brief User exposed API to create coalesced groups.
|
||||
@@ -625,6 +661,17 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
|
||||
"Tile size is either not a power of 2 or greater than the wavefront size");
|
||||
using tile_base<size>::numThreads;
|
||||
|
||||
template <unsigned int fsize, class fparent>
|
||||
friend __CG_QUALIFIER__ coalesced_group
|
||||
binary_partition(const thread_block_tile<fsize, fparent>& tgrp, bool pred);
|
||||
|
||||
#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
__CG_QUALIFIER__ unsigned long long build_mask() const {
|
||||
unsigned long long mask = ~0ull >> (64 - numThreads);
|
||||
return mask << ((internal::workgroup::thread_rank() / numThreads) * numThreads);
|
||||
}
|
||||
#endif
|
||||
|
||||
public:
|
||||
__CG_STATIC_QUALIFIER__ void sync() {
|
||||
internal::tiled_group::sync();
|
||||
@@ -649,7 +696,29 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
|
||||
static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
|
||||
return (__shfl_xor(var, laneMask, numThreads));
|
||||
}
|
||||
|
||||
#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
__CG_QUALIFIER__ unsigned long long ballot(int pred) const {
|
||||
const auto mask = build_mask();
|
||||
return internal::helper::adjust_mask(mask, __ballot_sync(mask, pred));
|
||||
}
|
||||
|
||||
__CG_QUALIFIER__ int any(int pred) const { return __any_sync(build_mask(), pred); }
|
||||
|
||||
__CG_QUALIFIER__ int all(int pred) const { return __all_sync(build_mask(), pred); }
|
||||
|
||||
template <typename T> __CG_QUALIFIER__ unsigned long long match_any(T value) const {
|
||||
const auto mask = build_mask();
|
||||
return internal::helper::adjust_mask(mask, __match_any_sync(mask, value));
|
||||
}
|
||||
|
||||
template <typename T> __CG_QUALIFIER__ unsigned long long match_all(T value, int& pred) const {
|
||||
const auto mask = build_mask();
|
||||
return internal::helper::adjust_mask(mask, __match_all_sync(mask, value, &pred));
|
||||
}
|
||||
#endif
|
||||
};
|
||||
|
||||
/** \brief User exposed API that captures the state of the parent group pre-partition
|
||||
*/
|
||||
template <unsigned int tileSize, typename ParentCGTy>
|
||||
@@ -727,6 +796,10 @@ class thread_block_tile_type<tileSize, void> : public thread_block_tile_base<til
|
||||
*/
|
||||
};
|
||||
|
||||
__CG_QUALIFIER__ thread_group this_thread() {
|
||||
thread_group g(internal::group_type::cg_coalesced_group, 1, __ockl_activelane_u32());
|
||||
return g;
|
||||
}
|
||||
|
||||
/** \brief User exposed API to partition groups.
|
||||
*
|
||||
@@ -763,8 +836,6 @@ __CG_QUALIFIER__ coalesced_group tiled_partition(const coalesced_group& parent,
|
||||
return (parent.new_tiled_group(tile_size));
|
||||
}
|
||||
|
||||
template <unsigned int size, class ParentCGTy> class thread_block_tile;
|
||||
|
||||
namespace impl {
|
||||
template <unsigned int size, class ParentCGTy> class thread_block_tile_internal;
|
||||
|
||||
@@ -829,6 +900,34 @@ __CG_QUALIFIER__ thread_block_tile<size, ParentCGTy> tiled_partition(const Paren
|
||||
"Tiled partition with size > wavefront size. Currently not supported ");
|
||||
return impl::tiled_partition_internal<size, ParentCGTy>(g);
|
||||
}
|
||||
|
||||
#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
|
||||
/** \brief Binary partition
|
||||
*
|
||||
* \details This splits the input thread group into two partitions determined by predicate
|
||||
*/
|
||||
__CG_QUALIFIER__ coalesced_group binary_partition(const coalesced_group& cgrp, bool pred) {
|
||||
auto mask = __ballot_sync<unsigned long long>(cgrp.coalesced_info.member_mask, pred);
|
||||
|
||||
if (pred) {
|
||||
return coalesced_group(mask);
|
||||
} else {
|
||||
return coalesced_group(cgrp.coalesced_info.member_mask ^ mask);
|
||||
}
|
||||
}
|
||||
|
||||
template <unsigned int size, class parent>
|
||||
__CG_QUALIFIER__ coalesced_group binary_partition(const thread_block_tile<size, parent>& tgrp,
|
||||
bool pred) {
|
||||
auto mask = __ballot_sync<unsigned long long>(tgrp.build_mask(), pred);
|
||||
|
||||
if (pred) {
|
||||
return coalesced_group(mask);
|
||||
} else {
|
||||
return coalesced_group(tgrp.build_mask() ^ mask);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
} // namespace cooperative_groups
|
||||
|
||||
#endif // __cplusplus
|
||||
|
||||
@@ -102,6 +102,33 @@ typedef enum {
|
||||
* on Windows.
|
||||
*
|
||||
*/
|
||||
namespace helper {
|
||||
/**
|
||||
* @brief Create output mask from input_mask at places where base_mask is set
|
||||
*
|
||||
* Example: base_mask = 0101'0101, input_mask = 1111'0000
|
||||
* Output mask: 1100
|
||||
* Explaination:
|
||||
* | | | | | | | |
|
||||
* base: 0|1|0|1|'0|1|0|1| // Which bits are set
|
||||
* input: 1|1|1|1|'0|0|0|0| // Which values are picked
|
||||
* | | | | | | | |
|
||||
* output: 1 1 0 0
|
||||
*/
|
||||
__CG_STATIC_QUALIFIER__ unsigned long long adjust_mask(
|
||||
unsigned long long base_mask, unsigned long long input_mask) {
|
||||
unsigned long long out = 0;
|
||||
for (unsigned int i = 0, index = 0; i < __AMDGCN_WAVEFRONT_SIZE; i++) {
|
||||
auto lane_active = base_mask & (1ull << i);
|
||||
if (lane_active) {
|
||||
auto result = input_mask & (1ull << i);
|
||||
out |= ((result ? 1ull : 0ull) << index);
|
||||
index++;
|
||||
}
|
||||
}
|
||||
return out;
|
||||
}
|
||||
} // namespace helper
|
||||
/**
|
||||
*
|
||||
* @brief Functionalities related to multi-grid cooperative group type
|
||||
|
||||
Ссылка в новой задаче
Block a user