31 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H 32 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H 36 #include <hip/hcc_detail/device_functions.h> 38 #if !defined(__align__) 39 #define __align__(x) __attribute__((aligned(x))) 42 #if !defined(__CG_QUALIFIER__) 43 #define __CG_QUALIFIER__ __device__ __forceinline__ 46 #if !defined(__CG_STATIC_QUALIFIER__) 47 #define __CG_STATIC_QUALIFIER__ __device__ static __forceinline__ 50 #if !defined(WAVEFRONT_SIZE) 51 #define WAVEFRONT_SIZE 64 54 namespace cooperative_groups {
69 namespace multi_grid {
71 __CG_STATIC_QUALIFIER__ uint32_t num_grids() {
72 return (uint32_t)__ockl_multi_grid_num_grids();
75 __CG_STATIC_QUALIFIER__ uint32_t grid_rank() {
76 return (uint32_t)__ockl_multi_grid_grid_rank();
79 __CG_STATIC_QUALIFIER__ uint32_t size() {
80 return (uint32_t)__ockl_multi_grid_size();
83 __CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
84 return (uint32_t)__ockl_multi_grid_thread_rank();
87 __CG_STATIC_QUALIFIER__
bool is_valid() {
88 return (
bool)__ockl_multi_grid_is_valid();
91 __CG_STATIC_QUALIFIER__
void sync() {
92 __ockl_multi_grid_sync();
102 __CG_STATIC_QUALIFIER__ uint32_t size() {
103 return (uint32_t)((hipBlockDim_z * hipGridDim_z) *
104 (hipBlockDim_y * hipGridDim_y) *
105 (hipBlockDim_x * hipGridDim_x));
108 __CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
111 (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) +
112 (hipBlockIdx_y * hipGridDim_x) +
117 uint32_t num_threads_till_current_workgroup =
118 (uint32_t)(blkIdx * (hipBlockIdx_x * hipBlockIdx_y * hipBlockIdx_z));
121 uint32_t local_thread_rank =
122 (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
123 (hipThreadIdx_y * hipBlockDim_x) +
126 return (num_threads_till_current_workgroup + local_thread_rank);
129 __CG_STATIC_QUALIFIER__
bool is_valid() {
130 return (
bool)__ockl_grid_is_valid();
133 __CG_STATIC_QUALIFIER__
void sync() {
143 #endif // __cplusplus 144 #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language ...