32 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H 33 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H 38 namespace cooperative_groups {
57 __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size,
58 uint64_t mask = (uint64_t)0) {
68 __CG_QUALIFIER__ uint32_t size()
const {
72 __CG_QUALIFIER__ uint32_t thread_rank()
const;
74 __CG_QUALIFIER__
bool is_valid()
const;
76 __CG_QUALIFIER__
void sync()
const;
85 class multi_grid_group :
public thread_group {
88 friend __CG_QUALIFIER__ multi_grid_group this_multi_grid();
92 explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size)
93 : thread_group(internal::cg_multi_grid, size) { }
98 __CG_QUALIFIER__ uint32_t num_grids() {
99 return internal::multi_grid::num_grids();
103 __CG_QUALIFIER__ uint32_t grid_rank() {
104 return internal::multi_grid::grid_rank();
106 __CG_QUALIFIER__ uint32_t thread_rank()
const {
107 return internal::multi_grid::thread_rank();
109 __CG_QUALIFIER__
bool is_valid()
const {
110 return internal::multi_grid::is_valid();
112 __CG_QUALIFIER__
void sync()
const {
113 internal::multi_grid::sync();
124 __CG_QUALIFIER__ multi_grid_group
126 return multi_grid_group(internal::multi_grid::size());
135 class grid_group :
public thread_group {
138 friend __CG_QUALIFIER__ grid_group this_grid();
142 explicit __CG_QUALIFIER__ grid_group(uint32_t size)
143 : thread_group(internal::cg_grid, size) { }
146 __CG_QUALIFIER__ uint32_t thread_rank()
const {
147 return internal::grid::thread_rank();
149 __CG_QUALIFIER__
bool is_valid()
const {
150 return internal::grid::is_valid();
152 __CG_QUALIFIER__
void sync()
const {
153 internal::grid::sync();
164 __CG_QUALIFIER__ grid_group
166 return grid_group(internal::grid::size());
172 __CG_QUALIFIER__ uint32_t thread_group::thread_rank()
const {
173 switch (this->_type) {
174 case internal::cg_multi_grid: {
175 return (static_cast<const multi_grid_group*>(
this)->thread_rank());
177 case internal::cg_grid: {
178 return (static_cast<const grid_group*>(
this)->thread_rank());
186 __CG_QUALIFIER__
bool thread_group::is_valid()
const {
187 switch (this->_type) {
188 case internal::cg_multi_grid: {
189 return (static_cast<const multi_grid_group*>(
this)->is_valid());
191 case internal::cg_grid: {
192 return (static_cast<const grid_group*>(
this)->is_valid());
200 __CG_QUALIFIER__
void thread_group::sync()
const {
201 switch (this->_type) {
202 case internal::cg_multi_grid: {
203 static_cast<const multi_grid_group*
>(
this)->sync();
206 case internal::cg_grid: {
207 static_cast<const grid_group*
>(
this)->sync();
215 #endif // __cplusplus 216 #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H Device side implementation of cooperative group feature.