The thread hierarchy abstractions of cooperative groups are depicted in the following figures: :ref:`grid hierarchy <coop_thread_top_hierarchy>` and :ref:`block hierarchy <coop_thread_bottom_hierarchy>`.
The **block** is the same as the :ref:`inherent_thread_model` block entity.
..note::
Explicit warp-level thread handling is absent from the Cooperative groups API. In order to exploit the known hardware SIMD width on which built-in functionality translates to simpler logic, you can use the group partitioning part of the API, such as ``tiled_partition``.
:alt:The new level between block thread and threads.
Cooperative group thread hierarchy in blocks.
The cooperative groups API introduce a new level between block thread and threads. The :ref:`thread-block tile <coop_thread_block_tile>` give the opportunity to have tiles in the thread block, while the :ref:`coalesced group <coop_coalesced_groups>` holds the active threads of the parent group. These groups further discussed in the :ref:`groups types <coop_group_types>` section.
For details on memory model, check the :ref:`memory model description <memory_hierarchy>`.
.._coop_group_types:
Group types
===========
Group types are based on the levels of synchronization and data sharing among threads.
Thread-block group
------------------
Represents an intra-block cooperative groups type where the participating threads within the group are the same threads that participated in the currently executing ``block``.
..code-block::cpp
classthread_block;
Constructed via:
..code-block::cpp
thread_blockg=this_thread_block();
The ``group_index()`` , ``thread_index()`` , ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()`` , ``sync()`` and ``group_dim()`` member functions are public of the thread_block class. For further details, check the :ref:`thread_block references <thread_block_ref>` .
Grid group
------------
Represents an inter-block cooperative groups type where the group's participating threads span multiple blocks running the same kernel on the same device. Use the cooperative launch API to enable synchronization across the grid group.
..code-block::cpp
classgrid_group;
Constructed via:
..code-block::cpp
grid_groupg=this_grid();
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()`` and ``sync()`` member functions
are public of the ``grid_group`` class. For further details, check the :ref:`grid_group references <grid_group_ref>`.
Multi-grid group
------------------
Represents an inter-device cooperative groups type where the participating threads within the group span multiple devices that run the same kernel on the devices. Use the cooperative launch API to enable synchronization across the multi-grid group.
..code-block::cpp
classmulti_grid_group;
Constructed via:
..code-block::cpp
// Kernel must be launched with the cooperative multi-device API
multi_grid_groupg=this_multi_grid();
The ``num_grids()`` , ``grid_rank()`` , ``thread_rank()``, ``size()``, ``cg_type()``, ``is_valid()`` ,
and ``sync()`` member functions are public of the ``multi_grid_group`` class. For
further details check the :ref:`multi_grid_group references <multi_grid_group_ref>` .
.._coop_thread_block_tile:
Thread-block tile
------------------
This constructs a templated class derived from ``thread_group``. The template defines the tile
size of the new thread group at compile time. This group type also supports sub-wave level intrinsics.
* Size must be a power of 2 and not larger than warp (wavefront) size.
*``shfl()`` functions support integer or float type.
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()``, ``sync()``, ``meta_group_rank()``, ``meta_group_size()``, ``shfl()``, ``shfl_down()``, ``shfl_up()``, ``shfl_xor()``, ``ballot()``, ``any()``, ``all()``, ``match_any()`` and ``match_all()`` member functions are public of the ``thread_block_tile`` class. For further details, check the :ref:`thread_block_tile references <thread_block_tile_ref>` .
Threads (64 threads on CDNA and 32 threads on RDNA) in a warp cannot execute different instructions simultaneously, so conditional branches are executed serially within the warp. When threads encounter a conditional branch, they can diverge, resulting in some threads being disabled if they do not meet the condition to execute that branch. The active threads are referred to as coalesced, and coalesced group represents an active thread group within a warp.
AMD GPUs do not support independent thread scheduling. Some CUDA application can rely on this feature and the ported HIP version on AMD GPUs can deadlock, when they try to make use of independent thread scheduling.
This group type also supports sub-wave level intrinsics.
..code-block::cpp
classcoalesced_group;
Constructed via:
..code-block::cpp
coalesced_groupactive=coalesced_threads();
..note::
``shfl()`` functions support integer or float type.
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()``, ``sync()``, ``meta_group_rank()``, ``meta_group_size()``, ``shfl()``, ``shfl_down()``, ``shfl_up()``, ``ballot()``, ``any()``, ``all()``, ``match_any()`` and ``match_all()`` member functions are public of the ``coalesced_group`` class. For more information, see :ref:`coalesced_group references <coalesced_group_ref>` .
Cooperative groups simple example
=================================
The difference to the original block model in the ``reduce_sum`` device function is the following.
..tab-set::
..tab-item:: Original Block
:sync:original-block
..code-block::cuda
__device__ int reduce_sum(int *shared, int val) {
// Thread ID
const unsigned int thread_id = threadIdx.x;
// Every iteration the number of active threads
// halves, until we processed all values
for(unsigned int i = blockDim.x / 2; i > 0; i /= 2) {
// Store value in shared memory with thread ID
shared[thread_id] = val;
// Synchronize all threads
__syncthreads();
// Active thread sum up
if(thread_id < i)
val += shared[thread_id + i];
// Synchronize all threads in the group
__syncthreads();
}
// ...
}
..tab-item:: Cooperative groups
:sync:cooperative-groups
..code-block::cuda
__device__ int reduce_sum(thread_group g,
int *shared,
int val) {
// Thread ID
const unsigned int group_thread_id = g.thread_rank();
// Every iteration the number of active threads
// halves, until we processed all values
for(unsigned int i = g.size() / 2; i > 0; i /= 2) {
// Store value in shared memroy with thread ID
shared[group_thread_id] = val;
// Synchronize all threads in the group
g.sync();
// Active thread sum up
if(group_thread_id < i)
val += shared[group_thread_id + i];
// Synchronize all threads in the group
g.sync();
}
// ...
}
The ``reduce_sum()`` function call and input data initialization difference to the original block model is the following.
At the device function, the input group type is the ``thread_group``, which is the parent class of all the cooperative groups type. With this, you can write generic functions, which can work with any type of cooperative groups.
.._coop_synchronization:
Synchronization
===============
With each group type, the synchronization requires using the correct cooperative groups launch API.
**Check the kernel launch capability**
..tab-set::
..tab-item:: Thread-block
:sync:thread-block
Do not need kernel launch validation.
..tab-item:: Grid
:sync:grid
Confirm the cooperative launch capability on the single AMD GPU:
..code-block::cpp
intdevice=0;
intsupports_coop_launch=0;
// Check support
// Use hipDeviceAttributeCooperativeMultiDeviceLaunch when launching across multiple devices
