Skip to content

Commit

Permalink
Prepare for 2nd round review
Browse files Browse the repository at this point in the history
  • Loading branch information
neon60 committed Jul 10, 2024
1 parent bfe4aff commit a1c70e1
Show file tree
Hide file tree
Showing 2 changed files with 19 additions and 21 deletions.
19 changes: 11 additions & 8 deletions docs/how-to/cooperative_groups.rst
Original file line number Diff line number Diff line change
Expand Up @@ -30,10 +30,10 @@ The **multi grid** is an abstraction of potentially multiple simultaneous launch

For further information, check the :ref:`inherent thread model <inherent_thread_model>`. For details on memory model, check the :ref:`memory model description <memory_hierarchy>`

Cooperative group simple example
================================
Cooperative groups simple example
=================================

You'll find the code difference to the original block model in the following examples.
The difference to the original block model in the ``reduce_sum`` device function is the following.

.. tab-set::
.. tab-item:: Original Block
Expand Down Expand Up @@ -98,6 +98,8 @@ You'll find the code difference to the original block model in the following exa
// ...
}
The ``reduce_sum`` function call and input data initialization difference to the original block model is the following.

.. tab-set::
.. tab-item:: Original Block
:sync: original-block
Expand Down Expand Up @@ -133,6 +135,7 @@ You'll find the code difference to the original block model in the following exa
// ...
// Initialize the thread_block
thread_block thread_block_group = this_thread_block();
// Perform reduction
output = reduce_sum(thread_block_group, workspace, input);
Expand All @@ -151,7 +154,7 @@ Group types are based on the levels of synchronization and data sharing among th
Thread-block group
------------------

Represents an intra-workgroup cooperative group type where the participating threads within the group are the same threads that participated in the currently executing ``workgroup``.
Represents an intra-workgroup cooperative groups type where the participating threads within the group are the same threads that participated in the currently executing ``workgroup``.

.. code-block:: cpp
Expand All @@ -168,7 +171,7 @@ The ``group_index()`` , ``thread_index()`` , ``thread_rank()`` , ``size()``, ``c
Grid group
------------

Represents an inter-workgroup cooperative group type where the group's participating threads span multiple workgroups running the same kernel on the same device. Use the cooperative launch API to synchronize across the grid.
Represents an inter-workgroup cooperative groups type where the group's participating threads span multiple workgroups running the same kernel on the same device. Use the cooperative launch API to synchronize across the grid.

.. code-block:: cpp
Expand All @@ -186,7 +189,7 @@ are public of the ``grid_group`` class. For further details, check the :ref:`gri
Multi-grid group
------------------

Represents an inter-device cooperative group type where the participating threads within the group span multiple devices that run the same kernel on the devices. All the multi-grid group APIs require that you have used the appropriate launch API.
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. All the multi-grid group APIs require that you have used the appropriate launch API.

.. code-block:: cpp
Expand All @@ -207,7 +210,7 @@ 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.
size of the new thread group at compile time. This group type also supports sub-wave level intrinsics.

.. code-block:: cpp
Expand Down Expand Up @@ -254,7 +257,7 @@ The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()``, ``sync()``, `
Synchronization
===============

With each group type, the synchronization requires using the correct cooperative group
With each group type, the synchronization requires using the correct cooperative groups
launch API.

**Check the kernel launch capability**
Expand Down
21 changes: 8 additions & 13 deletions docs/reference/cooperative_groups_reference.rst
Original file line number Diff line number Diff line change
Expand Up @@ -22,10 +22,10 @@ The following host side functions used for cooperative kernel launches.

.. doxygenfunction:: hipModuleLaunchCooperativeKernelMultiDevice

Cooperative group classes
=========================
Cooperative groups classes
==========================

The following cooperative group classes can be used on the device side.
The following cooperative groups classes can be used on the device side.

.. _thread_group_ref:

Expand All @@ -47,12 +47,7 @@ The following cooperative group classes can be used on the device side.
.. doxygenclass:: cooperative_groups::multi_grid_group
:members:

.. _tiled_group_ref:

.. doxygenclass:: cooperative_groups::tiled_group
:members:

.. _thread_block_tile_type_ref:
.. _thread_block_tile_ref:

.. doxygenclass:: cooperative_groups::thread_block_tile
:members:
Expand All @@ -62,8 +57,8 @@ The following cooperative group classes can be used on the device side.
.. doxygenclass:: cooperative_groups::coalesced_group
:members:

Cooperative group construct functions
=====================================
Cooperative groups construct functions
======================================

The following functions are used to construct different group types instances on the device side.

Expand All @@ -83,8 +78,8 @@ The following functions are used to construct different group types instances on

.. doxygenfunction:: cooperative_groups::binary_partition(const thread_block_tile<size, parent>& tgrp, bool pred)

Cooperative group exposed API functions
=======================================
Cooperative groups exposed API functions
========================================

The following functions are the exposed API for different group types instances on the device side.

Expand Down

0 comments on commit a1c70e1

Please sign in to comment.