点对点同步

以下部分讨论了 NVSHMEM API，这些 API 提供了一种基于对称数据对象的值在两个 PE 之间进行同步的机制。点对点同步例程可以用于可移植地确保内存访问操作按照发起 PE 使用 nvshmem_fence 和 nvshmem_quiet 例程强制执行的顺序观察远程更新。

标准 AMO 类型包括 C §7.18.1.1 和 C §7.20.1.1 在 stdint.h 中定义的一些精确宽度整数类型。当 C 转换环境不提供带有 stdint.h 的精确宽度整数类型时，NVSHMEM 实现不需要提供对这些类型的支持。nvshmem_test_any 和 nvshmem_wait_until_any 例程需要 C §7.18.3 和 C §7.20.3 在 stdint.h 中定义的 SIZE_MAX 宏。

点对点同步接口提供了命名常量，这些常量的值是整数常量表达式，用于指定 NVSHMEM 同步例程使用的比较运算符。常量名称和关联操作在表 [p2p-consts] 中列出。

点对点比较常量
常量名称 比较
NVSHMEM_CMP_EQ 等于
NVSHMEM_CMP_NE 不等于
NVSHMEM_CMP_GT 大于
NVSHMEM_CMP_GE 大于或等于
NVSHMEM_CMP_LT 小于
NVSHMEM_CMP_LE 小于或等于

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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

void nvshmemx_TYPENAME_wait_on_stream(TYPE *ivar, TYPE cmp_value, cudaStream_t stream)
__device__ void nvshmem_TYPENAME_wait(TYPE *ivar, TYPE cmp_value)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivar [输入]
远程可访问的数据对象的对称地址。ivar 的类型应与概要部分中暗示的类型匹配。
cmp [输入]
将 ivar 与 cmp_value 进行比较的比较运算符。
cmp_value [输入]
要与 ivar 进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait 和 nvshmem_wait_until 操作会阻塞，直到调用 PE 上的对称数据对象 ivar 中包含的值满足等待条件为止。调用 PE 上的 ivar 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。

这些例程可用于实现 PE 之间或同一 PE 内线程之间的点对点同步。调用 nvshmem_wait 会阻塞，直到调用 PE 上 ivar 的值不等于 cmp_value。nvshmem_wait_until 的调用会阻塞，直到调用 PE 上 ivar 的值满足由比较运算符 cmp 和比较值 cmp_value 指定的等待条件。

实现必须确保 nvshmem_wait 和 nvshmem_wait_until 在 ivar 指示的内存更新完全完成之前不会返回。

返回值

无

NVSHMEM_WAIT_UNTIL_ALL

__device__ void nvshmem_TYPENAME_wait_until_all(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_all 例程等待，直到 ivars 和 status 指定的等待集中的所有条目都满足调用 PE 上的等待条件为止。调用 PE 上的 ivars 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 nelems 为 0，则等待集为空，并且此例程立即返回。此例程根据调用 PE 上的比较运算符 cmp，将等待集中 ivars 数组的每个元素与值 cmp_value 进行比较。此例程在语义上类似于第 NVSHMEM_WAIT_UNTIL 节中的 nvshmem_wait_until，但增加了对涉及对称数据对象数组的点对点同步的支持。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程立即返回。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_all 在 ivars 指示的内存更新完全完成之前不会返回。

返回值

无。

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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_any 例程等待，直到 ivars 和 status 指定的等待集中的任何一个条目满足调用 PE 上的等待条件为止。调用 PE 上的 ivars 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。此例程根据调用 PE 上的比较运算符 cmp，将等待集中 ivars 数组的每个元素与值 cmp_value 进行比较。等待这些元素的顺序未指定。如果等待集中的 ivars 中的条目 \(i\) 满足等待条件，则一系列对 nvshmem_wait_until_any 的调用最终必须返回 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程返回 SIZE_MAX。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_any 在 ivars 指示的内存更新完全完成之前不会返回。

返回值

nvshmem_wait_until_any 返回 ivars 数组中满足等待条件的元素的索引。如果等待集为空，则此例程返回 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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
indices [输出]
长度至少为 nelems 的索引数组的本地地址，指向满足等待条件的 ivars 中的索引。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_some 例程等待，直到 ivars 和 status 指定的等待集中至少有一个条目满足调用 PE 上的等待条件为止。调用 PE 上的 ivars 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。此例程根据调用 PE 上的比较运算符 cmp，将等待集中 ivars 数组的每个元素与值 cmp_value 进行比较。此例程至少测试等待集中的 ivars 的所有元素一次，并且等待元素的顺序未指定。

返回后，indices 数组包含等待集中至少一个元素的索引，该元素在调用 nvshmem_wait_until_some 期间满足了等待条件。nvshmem_wait_until_some 的返回值等于这些满足条件的元素的总数。对于给定的返回值 \(N\)，indices 数组的前 \(N\) 个元素包含满足等待条件的那些唯一索引。 indices 的前 \(N\) 个元素可能相对于 ivars 的相应索引是无序的。indices 指向的数组的长度必须至少为 nelems。如果等待集中的 ivars 中的条目 \(i\) 满足等待条件，则一系列对 nvshmem_wait_until_some 的调用最终必须将 \(i\) 包含在 indices 数组中。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程返回 0。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars、indices 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_some 在 ivars 指示的内存更新完全完成之前不会返回。

返回值

nvshmem_wait_until_some 返回在 indices 数组中返回的索引数。如果等待集为空，则此例程返回 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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_all_vector 例程等待，直到 ivars 和 status 指定的等待集中的所有条目都满足调用 PE 上的等待条件为止。调用 PE 上的 ivars 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 nelems 为 0，则等待集为空，并且此例程立即返回。此例程根据调用 PE 上的比较运算符 cmp，将等待集中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程立即返回。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_all_vector 在 ivars 指示的内存更新完全完成之前不会返回。

返回值

无。

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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_any_vector 例程等待，直到 ivars 和 status 指定的等待集中的任何一个条目满足调用 PE 上的等待条件为止。调用 PE 上的 ivars 对象可能会被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。此例程根据调用 PE 上的比较运算符 cmp，将等待集中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。等待这些元素的顺序未指定。如果等待集中的 ivars 中的条目 \(i\) 满足等待条件，则一系列对 nvshmem_wait_until_any_vector 的调用最终必须返回 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程返回 SIZE_MAX。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_any_vector 在由 ivars 指示的内存更新完全完成之前不会返回。

返回值

nvshmem_wait_until_any_vector 返回 ivars 数组中满足等待条件的元素的索引。如果等待集合为空，则此例程返回 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)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
indices [输出]
长度至少为 nelems 的索引数组的本地地址，指向满足等待条件的 ivars 中的索引。
status [输入]
长度为 nelems 的可选掩码数组的本地地址，指示 ivars 中哪些元素从等待集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_wait_until_some_vector 例程等待直到由 ivars 和 status 指定的等待集合中的至少一个条目在调用 PE 上满足等待条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。此例程根据调用 PE 上的比较运算符 cmp，将等待集合中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。此例程至少测试等待集合中 ivars 的所有元素一次，并且等待元素的顺序未指定。

返回时，indices 数组包含等待集合中至少一个元素的索引，该元素在调用 nvshmem_wait_until_some_vector 期间满足等待条件。 nvshmem_wait_until_some_vector 的返回值等于这些满足条件的元素的总数。对于给定的返回值 \(N\)，indices 数组的前 \(N\) 个元素包含那些满足等待条件的唯一索引。 indices 的前 \(N\) 个元素可能相对于 ivars 的相应索引是无序的。 indices 指向的数组的长度必须至少为 nelems。如果等待集合中 ivars 中的条目 \(i\) 满足等待条件，则对 nvshmem_wait_until_some_vector 的一系列调用最终必须在 indices 数组中包含 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从等待集中排除。设置为 0 的 status 元素将包含在等待集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则等待集为空，并且此例程返回 0。如果 status 是空指针，则将其忽略，并且 ivars 中的所有元素都包含在等待集中。ivars、indices 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_wait_until_some_vector 在由 ivars 指示的内存更新完全完成之前不会返回。

返回值

nvshmem_wait_until_some_vector 返回在 indices 数组中返回的索引数。如果等待集合为空，则此例程返回 0。

NVSHMEM_TEST

__device__ int nvshmem_TYPENAME_test(TYPE *ivar, int cmp, TYPE cmp_value)

其中 TYPE 是标准 AMO 类型之一，并具有表 标准 AMO 类型和名称 指定的相应 TYPENAME，或者 TYPE 是信号类型之一，并具有表 信号类型和名称 指定的相应 TYPENAME。

ivar [输入]
远程可访问的数据对象的对称地址。ivar 的类型应与概要部分中暗示的类型匹配。
cmp [输入]
比较运算符，用于比较 ivar 和 cmp_value。
cmp_value [输入]
将与 ivar 指向的对象进行比较的值。 cmp_value 的类型应与 SYNOPSIS 部分中暗示的类型匹配。

描述

nvshmem_test 根据比较运算符 cmp，测试由 ivar 指向的对称对象与值 cmp_value 的数值比较。调用 PE 上的 ivar 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。

实现必须确保 nvshmem_test 在由 ivar 指示的内存更新完全完成之前不会返回 1。

返回值

如果根据比较运算符 cmp，由 ivar 指向的对称对象与值 cmp_value 的比较结果为真，则 nvshmem_test 返回 1；否则，返回 0。

NVSHMEM_TEST_ALL

__device__ int nvshmem_TYPENAME_test_all(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_all 例程指示由 ivars 和 status 指定的测试集中的所有条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中的并非所有条目都满足测试条件，则此例程不会阻塞并返回零。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与值 cmp_value 进行比较。

如果 nelems 为 0，则测试集为空，并且此例程返回 1。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 0。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_all 在由 ivars 指示的内存更新完全完成之前不会返回 1。

返回值

如果 ivars 中的所有变量都满足测试条件，或者如果 nelems 为 0，则 nvshmem_test_all 返回 1，否则此例程返回 0。

NVSHMEM_TEST_ANY

__device__ size_t nvshmem_TYPENAME_test_any(TYPE *ivars, size_t nelems, const int *status, int cmp, TYPE cmp_value)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_any 例程指示由 ivars 和 status 指定的测试集中的任何条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中没有条目满足测试条件，则此例程不会阻塞并返回 SIZE_MAX。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与值 cmp_value 进行比较。测试这些元素的顺序未指定。如果测试集中 ivars 中的条目 \(i\) 满足测试条件，则对 nvshmem_test_any 的一系列调用最终必须返回 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 SIZE_MAX。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_any 在由相应的 ivars 元素指示的内存更新完全完成之前不会返回索引。

返回值

nvshmem_test_any 返回 ivars 数组中满足测试条件的元素的索引。如果测试集为空或测试集中没有满足条件的条件，则此例程返回 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)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
indices [输出]
索引数组的本地地址，长度至少为 nelems，指向 ivars 中满足测试条件的元素。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_value 进行比较。
cmp_value [输入]
要与 ivars 指向的对象进行比较的值。cmp_value 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_some 例程指示由 ivars 和 status 指定的测试集中的至少一个条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中没有条目满足测试条件，则此例程不会阻塞并返回零。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与值 cmp_value 进行比较。此例程至少测试测试集中 ivars 的所有元素一次，并且测试元素的顺序未指定。如果测试集中 ivars 中的条目 \(i\) 满足测试条件，则对 nvshmem_test_some 的一系列调用最终必须返回 \(i\)。

返回时，indices 数组包含测试集中在调用 nvshmem_test_some 期间满足测试条件的元素的索引。 nvshmem_test_some 的返回值等于这些满足条件的元素的总数。如果返回值是 \(N\)，则 indices 数组的前 \(N\) 个元素包含那些满足测试条件的唯一索引。 indices 的前 \(N\) 个元素可能相对于 ivars 的相应索引是无序的。 indices 指向的数组的长度必须至少为 nelems。如果测试集中 ivars 中的条目 \(i\) 满足测试条件，则对 nvshmem_test_some 的一系列调用最终必须在 indices 数组中包含 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 0。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_some 在由相应的 ivars 元素指示的内存更新完全完成之前不会返回索引。

返回值

nvshmem_test_some 返回在 indices 数组中返回的索引数。如果测试集为空，则此例程返回 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)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_all_vector 例程指示由 ivars 和 status 指定的测试集中的所有条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中的并非所有条目都满足测试条件，则此例程不会阻塞并返回零。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。如果 nelems 为 0，则测试集为空，并且此例程返回 1。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 0。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_all_vector 在由 ivars 指示的内存更新完全完成之前不会返回 1。

返回值

nvshmem_test_all_vector 返回 1，如果 ivars 中的所有变量都满足测试条件，或者如果 nelems 为 0，否则此例程返回 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)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_any_vector 例程指示由 ivars 和 status 指定的测试集中的任何条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中没有条目满足测试条件，则此例程不会阻塞并返回 SIZE_MAX。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。测试这些元素的顺序未指定。如果测试集中 ivars 中的条目 \(i\) 满足测试条件，则对 nvshmem_test_any_vector 的一系列调用最终必须返回 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 SIZE_MAX。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_any_vector 在由相应的 ivars 元素指示的内存更新完全完成之前不会返回索引。

返回值

nvshmem_test_any_vector 返回 ivars 数组中满足测试条件的元素的索引。如果测试集为空或测试集中没有满足条件的条件，则此例程返回 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)

其中 *TYPE* 是标准 AMO 类型之一，并且具有表 标准 AMO 类型和名称 指定的相应 *TYPENAME*，或者 *TYPE* 是信号类型之一，并且具有表 信号类型和名称 指定的相应 *TYPENAME*。

ivars [输入]
远程可访问的数据对象数组的对称地址。ivars 的类型应与概要部分中暗示的类型匹配。
nelems [输入]
ivars 数组中的元素数量。
indices [输出]
索引数组的本地地址，长度至少为 nelems，指向 ivars 中满足测试条件的元素。
status [输入]
可选掩码数组的本地地址，长度为 nelems，指示 ivars 中的哪些元素从测试集中排除。
cmp [输入]
来自表 [p2p-consts] 的比较运算符，用于将 ivars 的元素与 cmp_values 的元素进行比较。
cmp_values [输入]
长度为 nelems 的数组的本地地址，该数组包含要与 ivars 中相应对象进行比较的值。cmp_values 的类型应与概要部分中暗示的类型匹配。

描述

nvshmem_test_some_vector 例程指示由 ivars 和 status 指定的测试集中的至少一个条目是否在调用 PE 上满足测试条件。调用 PE 上的 ivars 对象可能被位于调用 PE 内或另一个 PE 内的线程执行的 AMO 更新。如果 ivars 中没有条目满足测试条件，则此例程不会阻塞并返回零。此例程根据调用 PE 上的比较运算符 cmp，将测试集中 ivars 数组的每个元素与 cmp_values 中的每个相应值进行比较。此例程至少测试测试集中 ivars 的所有元素一次，并且测试元素的顺序未指定。

返回时，indices 数组包含测试集中在调用 nvshmem_test_some_vector 期间满足测试条件的元素的索引。 nvshmem_test_some_vector 的返回值等于这些满足条件的元素的总数。如果返回值是 \(N\)，则 indices 数组的前 \(N\) 个元素包含那些满足测试条件的唯一索引。 indices 的前 \(N\) 个元素可能相对于 ivars 的相应索引是无序的。 indices 指向的数组的长度必须至少为 nelems。如果测试集中 ivars 中的条目 \(i\) 满足测试条件，则对 nvshmem_test_some_vector 的一系列调用最终必须在 indices 数组中包含 \(i\)。

可选的 status 是长度为 nelems 的掩码数组，其中每个元素对应于 ivars 中的相应元素，并指示该元素是否从测试集中排除。设置为 0 的 status 元素将包含在测试集中，而设置为非零值的元素将被忽略。如果 status 中的所有元素均为非零值，或者 nelems 为 0，则测试集为空，并且此例程返回 0。如果 status 是空指针，则它将被忽略，并且 ivars 中的所有元素都将包含在测试集中。 ivars 和 status 数组在内存中不得重叠。

实现必须确保 nvshmem_test_some_vector 在由相应的 ivars 元素指示的内存更新完全完成之前不会返回索引。

返回值

nvshmem_test_some_vector 返回在 indices 数组中返回的索引数。如果测试集为空，则此例程返回 0。

NVSHMEM_SIGNAL_WAIT_UNTIL

uint64_t nvshmemx_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp, uint64_t cmp_val, cudaStream_t stream)
__device__ uint64_t nvshmem_signal_wait_until(uint64_t *sig_addr, int cmp, uint64_t cmp_val)
sig_addr [输入]
源信号变量的本地地址。
cmp [输入]
比较运算符，用于比较 sig_addr 和 cmp_val。
cmp_val [输入]
将与 sig_addr 指向的对象进行比较的值。

描述

nvshmem_signal_wait_until 操作会阻塞，直到调用 PE 上的信号数据对象 sig_addr 中包含的值满足等待条件。在具有单线程或多线程 PE 的 NVSHMEM 程序中，调用 PE 上的 sig_addr 对象预计仅作为信号进行更新，通过 NVSHMEM_PUT_SIGNAL 节和 NVSHMEM_PUT_SIGNAL_NBI 节中提供的信号操作进行更新。

此例程可用于实现 PE 之间或同一 PE 内的线程之间的点对点同步。对该例程的调用会阻塞，直到调用 PE 上 sig_addr 的值满足由比较运算符 cmp 和比较值 cmp_val 指定的等待条件。

实现必须确保 nvshmem_signal_wait_until 在由 sig_addr 指示的内存更新完全完成之前不会返回。

返回值

返回调用 PE 上满足等待条件的信号数据对象 sig_addr 的内容。