Point-To-Point Synchronization¶
The following section discusses NVSHMEMAPIs that provide a mechanism for
synchronization between two PEs based on the value of a symmetric data
object. The point-to-point synchronization routines can be used to
portably ensure that memory access operations observe remote updates in
the order enforced by the initiator PE using the nvshmem_fence
and
nvshmem_quiet
routines.
The standard AMO types include some of the exact-width integer types
defined in stdint.h
by C §7.18.1.1 and C §7.20.1.1. When the C
translation environment does not provide exact-width integer types with
stdint.h
, an NVSHMEM implemementation is not required to provide
support for these types. The nvshmem_test_any
and
nvshmem_wait_until_any
routines require the SIZE_MAX
macro
defined in stdint.h
by C §7.18.3 and C §7.20.3.
TYPE | TYPENAME |
---|---|
short | short |
int | int |
long | long |
long long | longlong |
unsigned short | ushort |
unsigned int | uint |
unsigned long | ulong |
unsigned long long | ulonglong |
int32_t | int32 |
int64_t | int64 |
uint32_t | uint32 |
uint64_t | uint64 |
size_t | size |
ptrdiff_t | ptrdiff |
The point-to-point synchronization interface provides named constants whose values are integer constant expressions that specify the comparison operators used by NVSHMEM synchronization routines. The constant names and associated operations are presented in Table 1.
Constant Name | Comparison |
---|---|
NVSHMEM_CMP_EQ |
Equal |
NVSHMEM_CMP_NE |
Not equal |
NVSHMEM_CMP_GT |
Greater than |
NVSHMEM_CMP_GE |
Greater than or equal to |
NVSHMEM_CMP_LT |
Less than |
NVSHMEM_CMP_LE |
Less than or equal to |
NVSHMEM_WAIT_UNTIL¶
-
void
nvshmemx_TYPENAME_wait_until_on_stream
(TYPE *ivar, int cmp, TYPE cmp_value, cudaStream_t stream)¶
-
__device__ void
nvshmem_TYPENAME_wait_until
(TYPE *ivar, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
-
void
nvshmemx_TYPENAME_wait_on_stream
(TYPE *ivar, TYPE cmp_value, cudaStream_t stream)¶
-
__device__ void
nvshmem_TYPENAME_wait
(TYPE *ivar, TYPE cmp_value)¶
where TYPE is one of {short
, int
, long
, long long
} and
has a corresponding TYPENAME specified by
Table Point-to-Point Synchronization Types and Names, or TYPE is one of the signal
types and has a corresponding TYPENAME specified by
Table Signal Types and Names.
- ivar [IN]
- Symmetric address of a remotely accessible data object. The type of
ivar
should match that implied in the SYNOPSIS section. - cmp [IN]
- The compare operator that compares
ivar
withcmp_value
. - cmp_value [IN]
- The value to be compared with
ivar
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait
and nvshmem_wait_until
operations block until
the value contained in the symmetric data object, ivar
, at the
calling PE satisfies the wait condition. The ivar
object at the
calling PE may be updated by an AMO performed by a thread located within
the calling PE or within another PE.
These routines can be used to implement point-to-point synchronization
between PEs or between threads within the same PE. A call to
nvshmem_wait
blocks until the value of ivar
at the calling PE is
not equal to cmp_value
. A call to nvshmem_wait_until
blocks
until the value of ivar
at the calling PE satisfies the wait
condition specified by the comparison operator, cmp
, and comparison
value, cmp_value
.
Implementations must ensure that nvshmem_wait
and
nvshmem_wait_until
do not return before the update of the memory
indicated by ivar
is fully complete.
Returns
None
NVSHMEM_WAIT_UNTIL_ALL¶
-
__device__ void
nvshmem_TYPENAME_wait_until_all
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_all
routine waits until all entries in the
wait set specified by ivars
and status
have satisfied the wait
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. If nelems
is 0, the wait set is empty and this
routine returns immediately. This routine compares each element of the
ivars
array in the wait set with the value cmp_value
according
to the comparison operator cmp
at the calling PE. This routine is
semantically similar to nvshmem_wait_until
in
Section NVSHMEM_WAIT_UNTIL, but adds support for
point-to-point synchronization involving an array of symmetric data
objects.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns immediately. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the wait set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_all
does not
return before the update of the memory indicated by ivars
is fully
complete.
Returns
None.
NVSHMEM_WAIT_UNTIL_ANY¶
-
__device__ size_t
nvshmem_TYPENAME_wait_until_any
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_any
routine waits until any one entry in the
wait set specified by ivars
and status
satisfies the wait
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine compares each element of the
ivars
array in the wait set with the value cmp_value
according
to the comparison operator cmp
at the calling PE. The order in which
these elements are waited upon is unspecified. If an entry in
ivars
within the wait set satisfies the wait condition, a series of
calls to nvshmem_wait_until_any
must eventually return .
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns SIZE_MAX
. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the wait set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_any
does not
return before the update of the memory indicated by ivars
is fully
complete.
Returns
nvshmem_wait_until_any
returns the index of an element in the
ivars
array that satisfies the wait condition. If the wait set is
empty, this routine returns SIZE_MAX
.
NVSHMEM_WAIT_UNTIL_SOME¶
-
__device__ size_t
nvshmem_TYPENAME_wait_until_some
(TYPE *ivars, size_t nelems, size_t *indices, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - indices [OUT]
- Local address of an array of indices of length at least
nelems
intoivars
that satisfied the wait condition. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_some
routine waits until at least one entry
in the wait set specified by ivars
and status
satisfies the wait
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine compares each element of the
ivars
array in the wait set with the value cmp_value
according
to the comparison operator cmp
at the calling PE. This routine tests
all elements of ivars
in the wait set at least once, and the order
in which the elements are waited upon is unspecified.
Upon return, the indices
array contains the indices of at least one
element in the wait set that satisfied the wait condition during the
call to nvshmem_wait_until_some
. The return value of
nvshmem_wait_until_some
is equal to the total number of these
satisfied elements. For a given return value , the first
elements of the
indices
array contain those unique indices
that satisfied the wait condition. These first elements of
indices
may be unordered with respect to the corresponding indices
of ivars
. The array pointed to by indices
must be at least
nelems
long. If an entry in
ivars
within the wait set
satisfies the wait condition, a series of calls to
nvshmem_wait_until_some
must eventually include in the
indices
array.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the wait set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_some
does not
return before the update of the memory indicated by ivars
is fully
complete.
Returns
nvshmem_wait_until_some
returns the number of indices returned in
the indices
array. If the wait set is empty, this routine returns 0.
NVSHMEM_WAIT_UNTIL_ALL_VECTOR¶
-
__device__ void
nvshmem_TYPENAME_wait_until_all_vector
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_all_vector
routine waits until all entries in
the wait set specified by ivars
and status
have satisfied the
wait conditions at the calling PE. The ivars
objects at the calling
PE may be updated by an AMO performed by a thread located within the
calling PE or within another PE. If nelems
is 0, the wait set is
empty and this routine returns immediately. This routine compares each
element of the ivars
array in the wait set with each respective
value in cmp_values
according to the comparison operator cmp
at
the calling PE.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns immediately. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the wait set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_all_vector
does
not return before the update of the memory indicated by ivars
is
fully complete.
Returns
None.
NVSHMEM_WAIT_UNTIL_ANY_VECTOR¶
-
__device__ size_t
nvshmem_TYPENAME_wait_until_any_vector
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_any_vector
routine waits until any one entry
in the wait set specified by ivars
and status
satisfies the wait
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine compares each element of the
ivars
array in the wait set with each respective value in
cmp_values
according to the comparison operator cmp
at the
calling PE. The order in which these elements are waited upon is
unspecified. If an entry in
ivars
within the wait set
satisfies the wait condition, a series of calls to
nvshmem_wait_until_any_vector
must eventually return .
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns SIZE_MAX
. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the wait set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_any_vector
does
not return before the update of the memory indicated by ivars
is
fully complete.
Returns
nvshmem_wait_until_any_vector
returns the index of an element in the
ivars
array that satisfies the wait condition. If the wait set is
empty, this routine returns SIZE_MAX
.
NVSHMEM_WAIT_UNTIL_SOME_VECTOR¶
-
__device__ size_t
nvshmem_TYPENAME_wait_until_some_vector
(TYPE *ivars, size_t nelems, size_t *indices, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - indices [OUT]
- Local address of an array of indices of length at least
nelems
intoivars
that satisfied the wait condition. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the wait set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_wait_until_some_vector
routine waits until at least one
entry in the wait set specified by ivars
and status
satisfies
the wait condition at the calling PE. The ivars
objects at the
calling PE may be updated by an AMO performed by a thread located within
the calling PE or within another PE. This routine compares each element
of the ivars
array in the wait set with each respective value in
cmp_values
according to the comparison operator cmp
at the
calling PE. This routine tests all elements of ivars
in the wait set
at least once, and the order in which the elements are waited upon is
unspecified.
Upon return, the indices
array contains the indices of at least one
element in the wait set that satisfied the wait condition during the
call to nvshmem_wait_until_some_vector
. The return value of
nvshmem_wait_until_some_vector
is equal to the total number of these
satisfied elements. For a given return value , the first
elements of the
indices
array contain those unique indices
that satisfied the wait condition. These first elements of
indices
may be unordered with respect to the corresponding indices
of ivars
. The array pointed to by indices
must be at least
nelems
long. If an entry in
ivars
within the wait set
satisfies the wait condition, a series of calls to
nvshmem_wait_until_some_vector
must eventually include in
the
indices
array.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the wait set. Elements of
status
set to 0 will be included in the wait set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the wait set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the wait set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_wait_until_some_vector
does
not return before the update of the memory indicated by ivars
is
fully complete.
Returns
nvshmem_wait_until_some_vector
returns the number of indices
returned in the indices
array. If the wait set is empty, this
routine returns 0.
NVSHMEM_TEST¶
-
__device__ int
nvshmem_TYPENAME_test
(TYPE *ivar, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivar [IN]
- Symmetric address of a remotely accessible data object. The type of
ivar
should match that implied in the SYNOPSIS section. - cmp [IN]
- The comparison operator that compares
ivar
withcmp_value
. - cmp_value [IN]
- The value against which the object pointed to by
ivar
will be compared. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
nvshmem_test
tests the numeric comparison of the symmetric object
pointed to by ivar
with the value cmp_value
according to the
comparison operator cmp
. The ivar
object at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE.
Implementations must ensure that nvshmem_test
does not return 1
before the update of the memory indicated by ivar
is fully complete.
Returns
nvshmem_test
returns 1 if the comparison of the symmetric object
pointed to by ivar
with the value cmp_value
according to the
comparison operator cmp
evaluates to true; otherwise, it returns 0.
NVSHMEM_TEST_ALL¶
-
__device__ int
nvshmem_TYPENAME_test_all
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_all
routine indicates whether all entries in the
test set specified by ivars
and status
have satisfied the test
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine does not block and returns zero if
not all entries in ivars
satisfied the test condition. This routine
compares each element of the ivars
array in the test set with the
value cmp_value
according to the comparison operator cmp
at the
calling PE.
If nelems
is 0, the test set is empty and this routine returns 1.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the test set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_all
does not return 1
before the update of the memory indicated by ivars
is fully
complete.
Returns
nvshmem_test_all
returns 1 if all variables in ivars
satisfy the
test condition or if nelems
is 0, otherwise this routine returns 0.
NVSHMEM_TEST_ANY¶
-
__device__ size_t
nvshmem_TYPENAME_test_any
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_any
routine indicates whether any entry in the test
set specified by ivars
and status
has satisfied the test
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine does not block and returns
SIZE_MAX
if no entries in ivars
satisfied the test condition.
This routine compares each element of the ivars
array in the test
set with the value cmp_value
according to the comparison operator
cmp
at the calling PE. The order in which these elements are tested
is unspecified. If an entry in
ivars
within the test set
satisfies the test condition, a series of calls to nvshmem_test_any
must eventually return .
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns SIZE_MAX
. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the test set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_any
does not return an
index before the update of the memory indicated by the corresponding
ivars
element is fully complete.
Returns
nvshmem_test_any
returns the index of an element in the ivars
array that satisfies the test condition. If the test set is empty or no
conditions in the test set are satisfied, this routine returns
SIZE_MAX
.
NVSHMEM_TEST_SOME¶
-
__device__ size_t
nvshmem_TYPENAME_test_some
(TYPE *ivars, size_t nelems, size_t *indices, const int *status, int cmp, TYPE cmp_value)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - indices [OUT]
- Local address of an array of indices of length at least
nelems
intoivars
that satisfied the test condition. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
withcmp_value
. - cmp_value [IN]
- The value to be compared with the objects pointed to by
ivars
. The type ofcmp_value
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_some
routine indicates whether at least one entry
in the test set specified by ivars
and status
satisfies the test
condition at the calling PE. The ivars
objects at the calling PE may
be updated by an AMO performed by a thread located within the calling PE
or within another PE. This routine does not block and returns zero if no
entries in ivars
satisfied the test condition. This routine compares
each element of the ivars
array in the test set with the value
cmp_value
according to the comparison operator cmp
at the
calling PE. This routine tests all elements of ivars
in the test set
at least once, and the order in which the elements are tested is
unspecified. If an entry in
ivars
within the test set
satisfies the test condition, a series of calls to nvshmem_test_some
must eventually return .
Upon return, the indices
array contains the indices of the elements
in the test set that satisfied the test condition during the call to
nvshmem_test_some
. The return value of nvshmem_test_some
is
equal to the total number of these satisfied elements. If the return
value is , then the first
elements of the
indices
array contain those unique indices that satisfied the test condition.
These first elements of
indices
may be unordered with
respect to the corresponding indices of ivars
. The array pointed to
by indices
must be at least nelems
long. If an entry
in
ivars
within the test set satisfies the test condition, a series
of calls to nvshmem_test_some
must eventually include in
the
indices
array.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the test set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_some
does not return
indices before the updates of the memory indicated by the corresponding
ivars
elements are fully complete.
Returns
nvshmem_test_some
returns the number of indices returned in the
indices
array. If the test set is empty, this routine returns 0.
NVSHMEM_TEST_ALL_VECTOR¶
-
__device__ int
nvshmem_TYPENAME_test_all_vector
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_all_vector
routine indicates whether all entries in
the test set specified by ivars
and status
have satisfied the
test condition at the calling PE. The ivars
objects at the calling
PE may be updated by an AMO performed by a thread located within the
calling PE or within another PE. This routine does not block and returns
zero if not all entries in ivars
satisfied the test conditions. This
routine compares each element of the ivars
array in the test set
with each respective value in cmp_values
according to the comparison
operator cmp
at the calling PE. If nelems
is 0, the test set is
empty and this routine returns 1.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the test set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_all_vector
does not
return 1 before the update of the memory indicated by ivars
is fully
complete.
Returns
nvshmem_test_all_vector
returns 1 if all variables in ivars
satisfy the test conditions or if nelems
is 0, otherwise this
routine returns 0.
NVSHMEM_TEST_ANY_VECTOR¶
-
__device__ size_t
nvshmem_TYPENAME_test_any_vector
(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_any_vector
routine indicates whether any entry in
the test set specified by ivars
and status
has satisfied the
test condition at the calling PE. The ivars
objects at the calling
PE may be updated by an AMO performed by a thread located within the
calling PE or within another PE. This routine does not block and returns
SIZE_MAX
if no entries in ivars
satisfied the test condition.
This routine compares each element of the ivars
array in the test
set with each respective value in cmp_values
according to the
comparison operator cmp
at the calling PE. The order in which these
elements are tested is unspecified. If an entry in
ivars
within the test set satisfies the test condition, a series of calls to
nvshmem_test_any_vector
must eventually return .
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns SIZE_MAX
. If status
is a null pointer, it is ignored and
all elements in ivars
are included in the test set. The ivars
and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_any_vector
does not
return an index before the update of the memory indicated by the
corresponding ivars
element is fully complete.
Returns
nvshmem_test_any_vector
returns the index of an element in the
ivars
array that satisfies the test condition. If the test set is
empty or no conditions in the test set are satisfied, this routine
returns SIZE_MAX
.
NVSHMEM_TEST_SOME_VECTOR¶
-
__device__ size_t
nvshmem_TYPENAME_test_some_vector
(TYPE *ivars, size_t nelems, size_t *indices, const int *status, int cmp, TYPE *cmp_values)¶
where TYPE is one of the standard AMO types and has a corresponding TYPENAME specified by Table Standard AMO Types and Names, or TYPE is one of the signal types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- ivars [IN]
- Symmetric address of an array of remotely accessible data objects.
The type of
ivars
should match that implied in the SYNOPSIS section. - nelems [IN]
- The number of elements in the
ivars
array. - indices [OUT]
- Local address of an array of indices of length at least
nelems
intoivars
that satisfied the test condition. - status [IN]
- Local address of an optional mask array of length
nelems
that indicates which elements inivars
are excluded from the test set. - cmp [IN]
- A comparison operator from Table 1 that compares
elements of
ivars
with elements ofcmp_values
. - cmp_values [IN]
- Local address of an array of length
nelems
containing values to be compared with the respective objects inivars
. The type ofcmp_values
should match that implied in the SYNOPSIS section.
Description
The nvshmem_test_some_vector
routine indicates whether at least one
entry in the test set specified by ivars
and status
satisfies
the test condition at the calling PE. The ivars
objects at the
calling PE may be updated by an AMO performed by a thread located within
the calling PE or within another PE. This routine does not block and
returns zero if no entries in ivars
satisfied the test condition.
This routine compares each element of the ivars
array in the test
set with each respective value in cmp_values
according to the
comparison operator cmp
at the calling PE. This routine tests all
elements of ivars
in the test set at least once, and the order in
which the elements are tested is unspecified.
Upon return, the indices
array contains the indices of the elements
in the test set that satisfied the test condition during the call to
nvshmem_test_some_vector
. The return value of
nvshmem_test_some_vector
is equal to the total number of these
satisfied elements. If the return value is , then the first
elements of the
indices
array contain those unique indices
that satisfied the test condition. These first elements of
indices
may be unordered with respect to the corresponding indices
of ivars
. The array pointed to by indices
must be at least
nelems
long. If an entry in
ivars
within the test set
satisfies the test condition, a series of calls to
nvshmem_test_some_vector
must eventually include in the
indices
array.
The optional status
is a mask array of length nelems
where each
element corresponds to the respective element in ivars
and indicates
whether the element is excluded from the test set. Elements of
status
set to 0 will be included in the test set, and elements set
to a nonzero value will be ignored. If all elements in status
are
nonzero or nelems
is 0, the test set is empty and this routine
returns 0. If status
is a null pointer, it is ignored and all
elements in ivars
are included in the test set. The ivars
,
indices
, and status
arrays must not overlap in memory.
Implementations must ensure that nvshmem_test_some_vector
does not
return indices before the updates of the memory indicated by the
corresponding ivars
elements are fully complete.
Returns
nvshmem_test_some_vector
returns the number of indices returned in
the indices
array. If the test set is empty, this routine returns 0.