AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-256723 - Added group partitioning feature in ROCm CG.

Browse Source 
Change-Id: Ie54046feef3baba857a7068972ec1fc0a60c2df9
This commit is contained in:
Sourabh Betigeri
2021-02-09 09:38:26 -08:00
committed by Sourabh Betigeri
parent dab4f5bdba
commit 0a0b026304
5 changed files with 1150 additions and 136 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/hip_cooperative_groups.h
+287 -81
View File
@@ -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<const thread_block*>(this)->thread_rank());
}
case internal::cg_tiled_group: {
return (static_cast<const tiled_group*>(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<const thread_block*>(this)->is_valid());
}
case internal::cg_tiled_group: {
return (static_cast<const tiled_group*>(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<const thread_block*>(this)->sync();
break;
}
case internal::cg_tiled_group: {
static_cast<const tiled_group*>(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 <class CGTy>
__CG_QUALIFIER__ uint32_t group_size(CGTy const &g) {
return g.size();
}
template <class CGTy> __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); }
template <class CGTy>
__CG_QUALIFIER__ uint32_t thread_rank(CGTy const &g) {
template <class CGTy> __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) {
return g.thread_rank();
}
template <class CGTy>
__CG_QUALIFIER__ bool is_valid(CGTy const &g) {
return g.is_valid();
template <class CGTy> __CG_QUALIFIER__ bool is_valid(CGTy const& g) { return g.is_valid(); }
template <class CGTy> __CG_QUALIFIER__ void sync(CGTy const& g) { g.sync(); }
template <unsigned int tileSize> 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 <unsigned int size> class thread_block_tile_base : public tile_base<size> {
static_assert(is_valid_tile_size<size>::value,
"Tile size is either not a power of 2 or greater than the wavefront size");
using tile_base<size>::numThreads;
public:
__CG_STATIC_QUALIFIER__ void sync() {
// enforce ordering for memory instructions
__builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
}
template <class T> __CG_QUALIFIER__ T shfl(T var, int srcRank) const {
static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
return (__shfl(var, srcRank, numThreads));
}
template <class T> __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const {
static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
return (__shfl_down(var, lane_delta, numThreads));
}
template <class T> __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const {
static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
return (__shfl_up(var, lane_delta, numThreads));
}
template <class T> __CG_QUALIFIER__ T shfl_xor(T var, unsigned int laneMask) const {
static_assert(is_valid_type<T>::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 <unsigned int tileSize, class ParentCGTy = void>
class thread_block_tile_type : public thread_block_tile_base<tileSize>, public tiled_group {
_CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
friend class thread_block_tile_type<tileSize, ParentCGTy>;
typedef thread_block_tile_base<numThreads> 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<const tiled_group*>(&parent);
return cg->new_tiled_group(tile_size);
} else {
const thread_block* tb = static_cast<const thread_block*>(&parent);
return tb->new_tiled_group(tile_size);
}
}
template <class CGTy>
__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 <unsigned int size, class ParentCGTy> class thread_block_tile;
namespace impl {
template <unsigned int size, class ParentCGTy> class thread_block_tile_internal;
template <unsigned int size, class ParentCGTy>
class thread_block_tile_internal : public thread_block_tile_type<size, ParentCGTy> {
protected:
template <unsigned int tbtSize, class tbtParentT>
__CG_QUALIFIER__ thread_block_tile_internal(
const thread_block_tile_internal<tbtSize, tbtParentT>& g)
: thread_block_tile_type<size, ParentCGTy>() {}
__CG_QUALIFIER__ thread_block_tile_internal(const thread_block& g)
: thread_block_tile_type<size, ParentCGTy>() {}
};
} // namespace impl
template <unsigned int size, class ParentCGTy>
class thread_block_tile : public impl::thread_block_tile_internal<size, ParentCGTy> {
protected:
__CG_QUALIFIER__ thread_block_tile(const ParentCGTy& g)
: impl::thread_block_tile_internal<size, ParentCGTy>(g) {}
public:
__CG_QUALIFIER__ operator thread_block_tile<size, void>() const {
return thread_block_tile<size, void>(*this);
}
};
template <unsigned int size>
class thread_block_tile<size, void> : public impl::thread_block_tile_internal<size, void> {
template <unsigned int, class ParentCGTy> friend class thread_block_tile;
protected:
public:
template <class ParentCGTy>
__CG_QUALIFIER__ thread_block_tile(const thread_block_tile<size, ParentCGTy>& g)
: impl::thread_block_tile_internal<size, void>(g) {}
};
template <unsigned int size, class ParentCGTy = void> class thread_block_tile;
namespace impl {
template <unsigned int size, class ParentCGTy = void> struct tiled_partition_internal;
template <unsigned int size>
struct tiled_partition_internal<size, thread_block> : public thread_block_tile<size, thread_block> {
__CG_QUALIFIER__ tiled_partition_internal(const thread_block& g)
: thread_block_tile<size, thread_block>(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 <unsigned int size, class ParentCGTy>
__CG_QUALIFIER__ thread_block_tile<size, ParentCGTy> tiled_partition(const ParentCGTy& g) {
static_assert(is_valid_tile_size<size>::value,
"Tiled partition with size > wavefront size. Currently not supported ");
return impl::tiled_partition_internal<size, ParentCGTy>(g);
}
} // namespace cooperative_groups
#endif // __cplusplus
#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+53 -55
View File
@@ -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 <unsigned int size>
using is_power_of_2 = std::integral_constant<bool, (size & (size - 1)) == 0>;
template <unsigned int size>
using is_valid_wavefront = std::integral_constant<bool, (size <= WAVEFRONT_SIZE)>;
template <unsigned int size>
using is_valid_tile_size =
std::integral_constant<bool, is_power_of_2<size>::value && is_valid_wavefront<size>::value>;
template <typename T>
using is_valid_type =
std::integral_constant<bool, std::is_integral<T>::value || std::is_floating_point<T>::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