.. _subsec:team: Team Management --------------- The PEs in an NVSHMEM program communicate using either point-to-point routines—such as RMA and AMO routines—that specify the PE number of the target PE, or collective routines that operate over a set of PEs. NVSHMEM teams allow programs to group a set of PEs for communication. Team-based collective operations include all PEs in a valid team. Point-to-point communication can make use of team-relative PE numbering through PE number translation. Predefined and Application-Defined Teams ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ An NVSHMEM team may be predefined (i.e., provided by the NVSHMEM library) or defined by the NVSHMEM application. An application-defined team is created by “splitting” a parent team into one or more new teams—each with some subset of PEs of the parent team—via one of the ``nvshmem_team_split_`` routines. All predefined teams are valid for the duration of the NVSHMEM portion of an application. Any team successfully created by a ``nvshmem_team_split_`` routine is valid until it is destroyed. All valid teams have a least one member. Team Handles ~~~~~~~~~~~~ A *team handle* is an opaque object with type ``nvshmem_team_t`` that is used to reference a team. Team handles are not remotely accessible objects. The predefined teams may be accessed via the team handles listed in Section :ref:`subsec:library_handles`. NVSHMEM communication routines that do not accept a team handle argument operate on the world team, which may be accessed through the ``NVSHMEM_TEAM_WORLD`` handle. The world team encompasses the set of all PEs in the NVSHMEM program, and a given PE’s number in the world team is equal to the value returned by ``nvshmem_my_pe``. A team handle may be initialized to or assigned the value ``NVSHMEM_TEAM_INVALID`` to indicate that handle does not reference a valid team. When managed in this way, applications can use an equality comparison to test whether a given team handle references a valid team. Thread Safety ~~~~~~~~~~~~~ When it is allowed by the threading model provided by the NVSHMEM library, a team may be used concurrently in non-collective operations (e.g., ``nvshmem_team_my_pe``) by multiple threads within the PE where it was created. A team may not be used concurrently by multiple threads in the same PE for collective operations. However, multiple collective operations on different teams may be performed in parallel. Collective Ordering ~~~~~~~~~~~~~~~~~~~ In NVSHMEM, a team object encapsulates resources used to communicate between PEs in collective operations. When calling multiple subsequent collective operations on a team, the collective operations—along with any relevant team based resources—are matched across the PEs in the team based on ordering of collective routine calls. It is the responsibility of the user to ensure that team-based collectives occur in the same program order across all PEs in a team. For a full discussion of collective semantics, see Section :ref:`subsec:coll`. Team Creation ~~~~~~~~~~~~~ Team creation is a collective operation on the parent team object. New teams result from a ``nvshmem_team_split_`` routine, which takes a parent team and other arguments and produces new teams that contain a subset of the PEs that are members of the parent team. All PEs in a parent team must participate in a split operation to create new teams. If a PE from the parent team is not a member of any resulting new teams, it will receive a value of ``NVSHMEM_TEAM_INVALID`` as the value for the new team handle. Teams that are created by a ``nvshmem_team_split_`` routine may be provided a configuration argument that specifies attributes of each new team. This configuration argument is of type ``nvshmem_team_config_t``, which is detailed further in Section :ref:`subsec:shmem_team_config_t`. PEs in a newly created team are consecutively numbered starting with PE number 0. PEs are ordered by their PE number in the parent team. Team relative PE numbers can be used for point-to-point operations through using the translation routine ``nvshmem_team_translate_pe``. Split operations are collective and are subject to the constraints on team-based collectives specified in Section :ref:`subsec:coll`. In particular, in multithreaded executions, threads at a given PE must not perform simultaneous split operations on the same parent team. Team creation operations are matched across participating PEs based on the order in which they are performed. Thus, team creation events must also occur in the same order on all PEs in the parent team. Upon completion of a team creation operation, the parent and any resulting child teams will be immediately usable for any team-based operations, including creating new child teams, without any intervening synchronization. .. _subsec:shmem_team_my_pe: **SHMEM_TEAM_MY_PE** ~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_my\_pe(shmem_team_t team) .. c:function:: __device__ int nvshmem\_team\_my\_pe(shmem_team_t team) *team [IN]* An NVSHMEM team handle. **Description** When ``team`` specifies a valid team, the ``nvshmem_team_my_pe`` routine returns the number of the calling PE within the specified team. The number is an integer between :math:`0` and :math:`N-1` for a team containing :math:`N` PEs. Each member of the team has a unique number. If ``team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then the value ``-1`` is returned. If ``team`` is otherwise invalid, the behavior is undefined. **Returns** The number of the calling PE within the specified team, or the value ``-1`` if the team handle compares equal to ``NVSHMEM_TEAM_INVALID``. **Notes** For the world team, this routine will return the same value as ``nvshmem_my_pe``. .. _subsec:shmem_team_n_pes: **SHMEM_TEAM_N_PES** ~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_n\_pes(shmem_team_t team) .. c:function:: __device__ int nvshmem\_team\_n\_pes(shmem_team_t team) *team [IN]* An NVSHMEM team handle. **Description** When ``team`` specifies a valid team, the ``nvshmem_team_n_pes`` routine returns the number of PEs in the team. This will always be a value between :math:`1` and :math:`N`, where :math:`N` is the total number of PEs running in the NVSHMEM program. If ``team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then the value ``-1`` is returned. If ``team`` is otherwise invalid, the behavior is undefined. **Returns** The number of PEs in the specified team, or the value ``-1`` if the team handle compares equal to ``NVSHMEM_TEAM_INVALID``. **Notes** For the world team, this routine will return the same value as ``nvshmem_n_pes``. .. _subsec:shmem_team_config_t: **SHMEM_TEAM_CONFIG_T** ~~~~~~~~~~~~~~~~~~~~~~~ .. code-block:: c typedef struct { int num_contexts; } nvshmem_team_config_t; None. **Description** A team configuration object is provided as an argument to ``nvshmem_team_split_`` routines. It specifies the requested capabilities of the team to be created. The ``num_contexts`` member is reserved for future use. When using the configuration structure to create teams, a mask parameter controls which fields may be accessed by the NVSHMEM library. Any configuration parameter value that is not indicated in the mask will be ignored, and the default value will be used instead. Therefore, a program must only set the fields for which it does not want the default value. A configuration mask is created through a bitwise OR operation of the following library constants. A configuration mask value of ``0`` indicates that the team should be created with the default values for all configuration parameters. *None* Reserved for future use. The default values for configuration parameters are: ``num_contexts`` = ``0`` .. _subsec:shmem_team_get_config: **SHMEM_TEAM_GET_CONFIG** ~~~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_get\_config(shmem_team_t team, long config_mask, shmem_team_config_t *config) *team [IN]* An NVSHMEM team handle. *config_mask [IN]* The bitwise mask representing the set of configuration parameters to fetch from the given team. *config [OUT]* A pointer to the configuration parameters for the given team. **Description** ``nvshmem_team_get_config`` returns through the ``config`` argument the configuration parameters as described by the mask, which were assigned according to input configuration parameters when the team was created. If ``team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then no operation is performed. If ``team`` is otherwise invalid, the behavior is undefined. **Returns** If ``team`` does not compare equal to ``NVSHMEM_TEAM_INVALID``, then ``nvshmem_team_get_config`` returns ``0``; otherwise, it returns nonzero. .. _subsec:shmem_team_translate_pe: **SHMEM_TEAM_TRANSLATE_PE** ~~~~~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_translate\_pe(shmem_team_t src_team, int src_pe, shmem_team_t dest_team) *src_team [IN]* An NVSHMEM team handle. *src_pe [IN]* A PE number in ``src_team``. *dest_team [IN]* An NVSHMEM team handle. **Description** The ``nvshmem_team_translate_pe`` routine will translate a given PE number in one team into the corresponding PE number in another team. Specifically, given the ``src_pe`` in ``src_team``, this routine returns that PE’s number in ``dest_team``. If ``src_pe`` is not a member of both ``src_team`` and ``dest_team``, a value of ``-1`` is returned. If at least one of ``src_team`` and ``dest_team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then ``-1`` is returned. If either of the ``src_team`` or ``dest_team`` handles are otherwise invalid, the behavior is undefined. **Returns** The specified PE’s number in the ``dest_team``, or a value of ``-1`` if any team handle arguments are invalid or the ``src_pe`` is not in both the source and destination teams. **Notes** If ``NVSHMEM_TEAM_WORLD`` is provided as the ``dest_team`` parameter, this routine acts as a global PE number translator and will return the corresponding ``NVSHMEM_TEAM_WORLD`` number. .. _subsec:shmem_team_split_strided: **SHMEM_TEAM_SPLIT_STRIDED** ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_split\_strided(shmem_team_t parent_team, int start, int stride, int size, const shmem_team_config_t *config, long config_mask, shmem_team_t *new_team) *parent_team [IN]* An NVSHMEM team. *start [IN]* The lowest PE number of the subset of PEs from the parent team that will form the new team. *stride [IN]* The stride between team PE numbers in the parent team that comprise the subset of PEs that will form the new team. *size [IN]* The number of PEs from the parent team in the subset of PEs that will form the new team. ``size`` must be a positive integer. *config [IN]* A pointer to the configuration parameters for the new team. *config_mask [IN]* The bitwise mask representing the set of configuration parameters to use from ``config``. *new_team [OUT]* An NVSHMEM team handle. Upon successful creation, it references an NVSHMEM team that contains the subset of all PEs in the parent team specified by the PE triplet provided. **Description** The ``nvshmem_team_split_strided`` routine is a collective routine. It creates a new NVSHMEM team from an existing parent team, where the PE subset of the resulting team is defined by the triplet of arguments (``start``, ``stride``, ``size``). A valid triplet is one such that: .. math:: start + stride \cdot i \in \mathbb{Z}_{N-1} \; \forall \; i \in \mathbb{Z}_{size-1} where :math:`\mathbb{Z}` is the set of natural numbers (:math:`0, 1, \dots`), :math:`N` is the number of PEs in the parent team and :math:`size` is a positive number indicating the number of PEs in the new team. The index :math:`i` specifies the number of the given PE in the new team. Thus, PEs in the new team remain in the same relative order as in the parent team. This routine must be called by all PEs in the parent team. All PEs must provide the same values for the PE triplet. On successful creation of the new team: - The ``new_team`` handle will reference a valid team for the subset of PEs in the parent team that are members of the new team. - Those PEs in the parent team that are not members of the new team will have ``new_team`` assigned to ``NVSHMEM_TEAM_INVALID``. - ``nvshmem_team_split_strided`` will return zero to all PEs in the parent team. If the new team cannot be created or an invalid PE triplet is provided, then ``new_team`` will be assigned the value ``NVSHMEM_TEAM_INVALID`` and ``nvshmem_team_split_strided`` will return a nonzero value on all PEs in the parent team. The ``config`` argument specifies team configuration parameters, which are described in Section :ref:`subsec:shmem_team_config_t`. The ``config_mask`` argument is a bitwise mask representing the set of configuration parameters to use from ``config``. A ``config_mask`` value of ``0`` indicates that the team should be created with the default values for all configuration parameters. See Section :ref:`subsec:shmem_team_config_t` for field mask names and default configuration parameters. If ``parent_team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then no new team will be created and ``new_team`` will be assigned the value ``NVSHMEM_TEAM_INVALID``. If ``parent_team`` is otherwise invalid, the behavior is undefined. **Returns** Zero on successful creation of ``new_team``; otherwise, nonzero. **Notes** The ``nvshmem_team_split_strided`` operation uses an arbitrary ``stride`` argument, whereas the ``logPE_stride`` argument to the active set collective operations only permits strides that are a power of two. Arbitrary strides allow a greater number of PE subsets to be expressed and can support a broader range of usage models. See the description of team handles and predefined teams in Section :ref:`subsec:team` for more information about team handle semantics and usage. .. _subsec:shmem_team_split_2d: **SHMEM_TEAM_SPLIT_2D** ~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: int nvshmem\_team\_split\_2d(shmem_team_t parent_team, int xrange, const shmem_team_config_t *xaxis_config, long xaxis_mask, shmem_team_t *xaxis_team, const shmem_team_config_t *yaxis_config, long yaxis_mask, shmem_team_t *yaxis_team) *parent_team [IN]* A valid NVSHMEM team. Any predefined teams, such as ``NVSHMEM_TEAM_WORLD``, may be used, or any team created by the user. *xrange [IN]* A positive integer representing the number of elements in the first dimension. *xaxis_config [IN]* A pointer to the configuration parameters for the new ``x``-axis team. *xaxis_mask [IN]* The bitwise mask representing the set of configuration parameters to use from ``xaxis_config``. *xaxis_team [OUT]* A new PE team handle representing a PE subset consisting of all the PEs that have the same coordinate along the ``y``-axis as the calling PE. *yaxis_config [IN]* A pointer to the configuration parameters for the new ``y``-axis team. *yaxis_mask [IN]* The bitwise mask representing the set of configuration parameters to use from ``yaxis_config``. *yaxis_team [OUT]* A new PE team handle representing a PE subset consisting of all the PEs that have the same coordinate along the ``x``-axis as the calling PE. **Description** The ``nvshmem_team_split_2d`` routine is a collective operation. It returns two new teams to the calling PE by splitting an existing parent team into subsets based on a 2D Cartesian space. The user provides the size of the ``x`` dimension, which is then used to derive the size of the ``y`` dimension based on the size of the parent team. The size of the ``y`` dimension will be equal to :math:`\lceil N \div xrange \rceil`, where ``N`` is the size of the parent team. In other words, :math:`xrange \times yrange \geq N`, so that every PE in the parent team has a unique ``(x,y)`` location in the 2D Cartesian space. The resulting ``xaxis_team`` and ``yaxis_team`` correspond to the calling PE’s row and column, respectively, in the 2D Cartesian space. The mapping of PE number to coordinates is :math:`(x, y) = ( pe \mod xrange, \lfloor pe \div xrange \rfloor )`, where :math:`pe` is the PE number in the parent team. For example, if :math:`xrange = 3`, then the first 3 PEs in the parent team will form the first ``xteam``, the second three PEs in the parent team form the second ``xteam``, and so on. Thus, after the split operation, each of the new ``xteam`` s will contain all PEs that have the same coordinate along the ``y``-axis as the calling PE. Each of the new ``yteam`` s will contain all PEs with the same coordinate along the ``x``-axis as the calling PE. The PEs are numbered in the new teams based on the coordinate of the PE along the given axis. As a result, the value returned by ``nvshmem_team_my_pe(xteam)`` is the ``x``-coordinate and the value returned by ``nvshmem_team_my_pe(yteam)`` is the ``y``-coordinate of the calling PE. Any valid NVSHMEM team can be used as the parent team. This routine must be called by all PEs in the parent team. The value of ``xrange`` must be positive and all PEs in the parent team must pass the same value for ``xrange``. When ``xrange`` is greater than the size of the parent team, ``nvshmem_team_split_2d`` behaves as though ``xrange`` were equal to the size of the parent team. The ``xaxis_config`` and ``yaxis_config`` arguments specify team configuration parameters for the ``x``- and ``y``-axis teams, respectively. These parameters are described in Section :ref:`subsec:shmem_team_config_t`. All PEs that will be in the same resultant team must specify the same configuration parameters. The PEs in the parent team *do not* have to all provide the same parameters for new teams. The ``xaxis_mask`` and ``yaxis_mask`` arguments are a bitwise masks representing the set of configuration parameters to use from ``xaxis_config`` and ``yaxis_config``, respectively. A mask value of ``0`` indicates that the team should be created with the default values for all configuration parameters. See Section :ref:`subsec:shmem_team_config_t` for field mask names and default configuration parameters. If ``parent_team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then no new teams will be created and both ``xaxis_team`` and ``yaxis_team`` will be assigned the value ``NVSHMEM_TEAM_INVALID``. If ``parent_team`` is otherwise invalid, the behavior is undefined. If any ``xaxis_team`` or ``yaxis_team`` on any PE in ``parent_team`` cannot be created, then both team handles on all PEs in ``parent_team`` will be assigned the value ``NVSHMEM_TEAM_INVALID`` and ``nvshmem_team_split_2d`` will return a nonzero value. **Returns** Zero on successful creation of all ``xaxis_team`` s and ``yaxis_team`` s; otherwise, nonzero. **Notes** Since the split may result in a 2D space with more points than there are members of the parent team, there may be a final, incomplete row of the 2D mapping of the parent team. This means that the resultant ``yteam`` s may vary in size by up to 1 PE, and that there may be one resultant ``xteam`` of smaller size than all of the other ``xteam`` s. The following grid shows the 12 teams that would result from splitting a parent team of size 10 with ``xrange`` of 3. The numbers in the grid cells are the PE numbers in the parent team. The rows are the ``xteam`` s. The columns are the ``yteam``\ s. ========== ===== ===== ===== yteam yteam yteam x=0 x=1 x=2 xteam, y=0 0 1 2 xteam, y=1 3 4 5 xteam, y=2 6 7 8 xteam, y=3 9 ========== ===== ===== ===== It would be legal, for example, if PEs 0, 3, 6, 9 specified a different value for ``yaxis_config`` than all of the other PEs, as long as the configuration parameters match for all PEs in each of the new teams. See the description of team handles and predefined teams in Section :ref:`subsec:team` for more information about team handle semantics and usage. .. _subsec:shmem_team_destroy: **SHMEM_TEAM_DESTROY** ~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: void nvshmem\_team\_destroy(shmem_team_t team) *team [IN]* An NVSHMEM team handle. **Description** The ``nvshmem_team_destroy`` routine is a collective operation that destroys the team referenced by the team handle argument ``team``. Upon return, the referenced team is invalid. If ``team`` compares equal to ``NVSHMEM_TEAM_WORLD`` or any other predefined team, the behavior is undefined. If ``team`` compares equal to ``NVSHMEM_TEAM_INVALID``, then no operation is performed. If ``team`` is otherwise invalid, the behavior is undefined. **Returns** None.