4.18. 计算互连传输

4.18.1. 先决条件与范围

使用逻辑端点有以下要求：

• CUDA 驱动程序 API。逻辑端点 API 是底层CUDA 驱动程序 API的一部分。

• NVLink 互连网络（fabric）连通性与设备支持。参与的 GPU 必须能够通过 NVLink 互连网络（fabric）相互访问。设备是否支持逻辑端点取决于其架构、驱动程序和系统配置，因此应用程序应通过查询第 4.18.2.2 节中描述的设备属性来确定是否支持，而不是根据现有硬件进行推断。

• IMEX 守护进程与通道。在进程间共享端点使用互连网络（fabric）IPC 句柄，这要求 NVIDIA IMEX 守护进程正在运行，并且已配置 IMEX 通道。

端点可以与同一 NVLink 互连网络（fabric）内同一节点或不同节点上的进程共享。互连网络（fabric）IPC 句柄机制与 VMM API 使用的机制类似。有关互连网络（fabric）句柄和 IMEX 通道的详细信息，请参阅虚拟内存管理。

4.18.2. 预备知识

以下定义确立了本节后续内容中使用的关键对象和操作。

int unicastSupported = 0;
cuDeviceGetAttribute(&unicastSupported,
                     CU_DEVICE_ATTRIBUTE_LOGICAL_ENDPOINT_UNICAST_SUPPORTED,
                     device);
if (unicastSupported != 0) {
    // `device` supports unicast logical endpoints
}

int multicastSupported = 0;
cuDeviceGetAttribute(&multicastSupported,
                     CU_DEVICE_ATTRIBUTE_LOGICAL_ENDPOINT_MULTICAST_SUPPORTED,
                     device);
if (multicastSupported != 0) {
    // `device` supports multicast logical endpoints
}

int supportedHandleTypes = 0;
cuDeviceGetAttribute(
    &supportedHandleTypes,
    CU_DEVICE_ATTRIBUTE_LOGICAL_ENDPOINT_SUPPORTED_HANDLE_TYPES,
    device);
if ((supportedHandleTypes &
     CU_LOGICAL_ENDPOINT_IPC_HANDLE_TYPE_FABRIC) != 0) {
    // `device` supports fabric IPC handles for logical endpoints
}

int countedOpsSupported = 0;
cuDeviceGetAttribute(&countedOpsSupported,
                     CU_DEVICE_ATTRIBUTE_LOGICAL_ENDPOINT_COUNTED_OPS_SUPPORTED,
                     device);
if (countedOpsSupported != 0) {
    // `device` supports counted operations via logical endpoints
}

int ownerAccessSupported = 0;
cuDeviceGetAttribute(&ownerAccessSupported,
                     CU_DEVICE_ATTRIBUTE_LOGICAL_ENDPOINT_UNICAST_ACCESS_ON_OWNER_DEVICE_SUPPORTED,
                     device);
if (ownerAccessSupported != 0) {
    // `device` supports unicast logical endpoint access on the owner device
}

4.18.3. API 概述

逻辑端点 API 是底层CUDA Driver API的一部分，使用逻辑端点遵循明确定义的生命周期。下面的步骤概述了工作流程，后续各节将详细介绍每一步。这些步骤分为以下阶段：

• 设置阶段（步骤 1-11）：验证参与设备是同一互连组的成员，预留一个 id，描述并创建端点，与对等端共享端点，确认端点已就绪，然后在开始数据移动之前分配并绑定其后备内存。

• 使用阶段（步骤 12）：从设备发起针对该端点的互连网络操作。

• 清理阶段（步骤 13）：在所有使用均已完成后，解除端点的绑定并销毁端点，释放其 id，并释放后备内存。

1. 验证互连组成员资格。按照第 4.18.4 节所述，使用cuDeviceGetCliqueCount和cuDeviceGetCliqueInfo验证参与设备属于同一互连组。

2. 预留 id。使用cuLogicalEndpointIdReserve预留一个逻辑端点 id 范围。预留以进程为单位。id 的生命周期以及重用和释放规则在第 4.18.6 节中介绍。

3. 描述端点。在CUlogicalEndpointProp中填入端点类型（单播或多播）、单播的所有者设备或多播的组大小、大小、请求的 IPC 句柄类型（ipcHandleTypes），以及任何标志，例如使用CU_LOGICAL_ENDPOINT_FLAG_COUNTED_OPS请求计数操作支持（请参阅计数操作）。

4. 检查对齐和大小限制。针对拟使用的属性调用cuLogicalEndpointGetLimits，以获取所需的bindAlignment和maxSize。maxSize是该配置允许的最大端点大小。端点大小和所有绑定都必须满足这些限制。

5. 创建端点。调用cuLogicalEndpointCreate，将其中一个预留 id 与新创建的端点关联起来，该端点由上面填充的CUlogicalEndpointProp描述。

6. 共享端点。使用cuLogicalEndpointExport将端点导出为互连网络 IPC 句柄，通过任意 IPC 机制将该句柄传递给需要它的对等进程，并让每个对等进程使用cuLogicalEndpointImport将其导入。导出和导入端点并不要求端点已绑定内存。

7. 添加设备（仅限多播）。对于多播端点，必须使用cuLogicalEndpointAddDevice将每个参与设备添加到组中。每个进程在本地持有该端点之后添加自己的设备，因此对于多播，先共享端点，再添加设备。

8. 确认就绪。cuLogicalEndpointCreate和cuLogicalEndpointImport都是非阻塞的。使用前，请使用cuLogicalEndpointQuery确认就绪状态。对于本地创建的端点，应在绑定其内存之前确认；对于导入的端点，应在核函数发起针对它的操作之前确认。

9. 分配后备内存。使用虚拟内存管理 API（cuMemCreate）或通过cudaMallocAsync分配后备内存。如果端点将与另一个进程共享，所绑定的内存必须能够对外共享。通过导入的多播端点绑定内存时，也适用相同的要求。创建内存分配时，通过请求互连网络句柄类型（CU_MEM_HANDLE_TYPE_FABRIC）来将其配置为可对外共享，并设置分配大小，使其同时满足bindAlignment和分配粒度的要求。此步骤与端点创建相互独立，但两者都必须在绑定之前完成。

10. 绑定内存。使用cuLogicalEndpointBindMem（通过分配句柄）或cuLogicalEndpointBindAddr（通过映射地址）将物理内存与端点关联起来。不能将内存绑定到导入的单播端点，但可以将内存绑定到导入的多播端点。互连网络操作仅对已绑定的范围有效。

11. 绑定后同步。在任何进程发起互连网络操作之前，参与进程必须在绑定后进行同步，如第 4.18.8 节所述。

12. 使用端点。启动核函数，由核函数发起针对对等端点 id 的互连网络操作。

13. 清理。解除内存绑定，销毁端点，释放 id 范围，并释放后备内存。

4.18.4. 验证互连组成员资格

在发起任何计算互连传输操作之前，我们需要确保所有参与设备都是支持所需计算互连传输操作集的互连组的成员，并且对于特定的互连组类型，它们都是同一互连组的成员。否则，设备无法通过网络使用该特定类型的操作相互访问。

首先，我们查询该设备是几个互连组的成员：

  size_t cliqueCount = 0; // not modified if there are errors
  cuDeviceGetCliqueCount(&cliqueCount, cuDevice);

然后，我们查询该设备所属的所有互连组的信息，并找到将要使用的互连组的互连组 id。在此示例中，我们将使用类型为CU_CLIQUE_TYPE_UNICAST_LOGICAL_ENDPOINT的互连组。

  std::vector<CUcliqueInfo> cliqueInfos(cliqueCount);
  cuDeviceGetCliqueInfo(cliqueInfos.data(), &cliqueCount, cuDevice);
  unsigned int localCliqueId = 0;
  bool found = false;
  for (size_t idx = 0; idx < cliqueCount; ++idx) {
    if (cliqueInfos[idx].type == CU_CLIQUE_TYPE_UNICAST_LOGICAL_ENDPOINT) {
      localCliqueId = cliqueInfos[idx].id;
      found = true;
      break;
    }
  }

验证所有参与设备对于该互连组类型都具有相同的互连组 id。

  std::vector<unsigned int> allCliqueIds(numRanks);
  MPI_Allgather(&localCliqueId, sizeof(unsigned int), MPI_BYTE,
                          allCliqueIds.data(), sizeof(unsigned int), MPI_BYTE, MPI_COMM_WORLD);
  for (int i = 0; i < numRanks; i++) {
    if (allCliqueIds[i] != localCliqueId) {
      fprintf(stderr, "Rank %d: Clique id %u does not match expected %u\n", myRank, allCliqueIds[i], localCliqueId);
      return EXIT_FAILURE;
    }
  }

4.18.5. 创建端点

创建和导入端点首先要确定进程需要多少个逻辑端点 id。然后使用cuLogicalEndpointIdReserve预留一个 id 范围，该函数返回该范围的基准 id。选择足够大的数量，以满足进程将创建或导入的每个端点的需要。与该范围关联的端点可以具有不同的属性。例如，应用程序可以为每个 GPU 预留一个 id，再为一个多播端点额外预留一个 id。此示例为每个秩预留一个 id：

  CUlogicalEndpointId leId = 0;
  cuLogicalEndpointIdReserve(&leId, (uint32_t)numRanks);

预留范围之后，在创建或导入端点时关联其中的 id。对于本地创建的端点，使用cuLogicalEndpointCreate；当远端端点的句柄可用时，使用cuLogicalEndpointImport。对于进程创建的每个端点，在CUlogicalEndpointProp中描述其属性。单播端点指定将持有其所绑定内存的唯一所有者设备。设置端点类型、请求的 IPC 句柄类型（ipcHandleTypes）以及任何标志，仅在有请求时才启用计数操作：

  CUlogicalEndpointProp endpointProp {};
  endpointProp.type = CU_LOGICAL_ENDPOINT_TYPE_UNICAST;
  endpointProp.unicast = {.device = cuDevice};
  endpointProp.ipcHandleTypes = CU_LOGICAL_ENDPOINT_IPC_HANDLE_TYPE_FABRIC;
  endpointProp.flags = useCounted ? CU_LOGICAL_ENDPOINT_FLAG_COUNTED_OPS : CU_LOGICAL_ENDPOINT_FLAG_NONE;

端点大小在得知其对齐和最大大小限制后单独设置。有关如何查询这些限制以及如何选择大小，请参阅第 4.18.6.1 节。

通过调用cuLogicalEndpointCreate，将一个预留 id 与所描述的属性关联起来。创建本地端点并导入对等端点的进程通常为每个参与者预留一个 id，以便通过基准 id 加上对等端的秩来寻址该对等端。每个秩在leId + myRank处创建自己的端点：

  cuLogicalEndpointCreate(leId + myRank, &endpointProp);

cuLogicalEndpointCreate是非阻塞的，因此它返回时端点并不能立即使用。在将内存绑定到端点之前，使用cuLogicalEndpointQuery确认端点已完全构建，如第 4.18.7.1 节所述。

多播端点。多播端点遵循相同的步骤，区别在于其描述方式，以及它需要添加一组设备。描述时使用type = CU_LOGICAL_ENDPOINT_TYPE_MULTICAST，并将multicast.numDevices设置为多播组中的设备数量。创建并共享端点后，每个参与设备都必须使用cuLogicalEndpointAddDevice加入该组。进程只有在本地持有端点后才添加自己的设备，创建者在cuLogicalEndpointCreate之后添加，各对等端在cuLogicalEndpointImport之后添加。在多播组完整之前，端点尚未就绪，无法使用。在端点的整个生命周期内，成员资格是永久性的。

4.18.6. 逻辑端点 id

逻辑端点 id 是进程局部的名称，其生命周期独立于它所标识的端点的生命周期。端点 id 始终处于以下状态之一：

1. 未预留。该 id 不归该进程所有，不能与端点关联。cuLogicalEndpointIdReserve会将其预留。

2. 已预留。该 id 归该进程所有，但尚未指代任何对象。使用已预留的 id 调用cuLogicalEndpointCreate或cuLogicalEndpointImport会将其与一个端点关联，并使其进入已关联状态。cuLogicalEndpointIdRelease会使其回到未预留状态。

3. 已关联。该 id 指代一个端点，使用该 id 的互连网络操作以该端点为目标。cuLogicalEndpointDestroy会移除该关联，并使该 id 回到已预留状态。

cuLogicalEndpointIdReserve预留一个包含count个端点 id 的范围，并返回其基准baseLeId，因此预留的范围为[baseLeId, baseLeId + count)。基准是一个输出值：调用者不能选择获得哪些 id，并且如果没有包含count个 id 的可用范围，预留就会失败，例如，进程中的另一个库持有预留范围时就可能出现这种情况。

一个端点可以同时与多个 id 关联。例如，它可以通过cuLogicalEndpointCreate与所有者的 id 关联，并通过cuLogicalEndpointImport与其他 id 关联。这些 id 是同一端点的别名，针对其中任意一个 id 发起的操作都会到达同一目标。

cuLogicalEndpointDestroy移除一个 id 的关联。此后，通过后续的创建或导入操作，该 id 可以与另一个端点关联，无需再次预留。销毁操作还会解除该进程通过此 id 绑定的所有内存的绑定。只有在端点的最后一个别名被销毁时，端点本身才会被释放。

cuLogicalEndpointIdRelease(baseLeId, count)释放范围[baseLeId, baseLeId + count)中最多count个 id。该范围中的每个 id 都必须已经预留，并且与该范围中的 id 关联的每个端点都必须先被销毁。

由于预留以进程为单位，因此在每个导入进程中导入端点时，不必使用相同的 id。一种方便的约定是，每个进程预留一个大小相同的范围，为每个参与者预留一个 id，并在本地以baseLeId + peer来寻址给定对等端的端点。每个进程都有自己预留的baseLeId，它在各进程之间不必相同。如果在不同进程中使用不同 id 导入同一个对等端点，CUDA 不会检测或报告这种不匹配。应用程序有责任维护正确的本地对等端到端点 id 的映射。

cuuint64_t bindAlignment = 0;
cuuint64_t maxSize = 0;
cuLogicalEndpointGetLimits(&bindAlignment, &maxSize, &prop);

4.18.7. 共享端点

与对等方共享端点遵循导出、传输、导入的交换过程：所有者将端点导出为互连网络（fabric）IPC 句柄，将该句柄传输给需要它的对等方，然后每个对等方导入端点，并将其与自身预留的某个尚未关联端点的逻辑端点 ID 关联。共享使用互连网络（fabric）IPC 句柄，因此端点必须绑定到可对外共享的内存，并且环境必须满足第 4.18.1 节中描述的互连网络（fabric）IPC 要求。

使用cuLogicalEndpointExport导出端点，请求与描述该端点时所用类型相同的 IPC 句柄类型（ipcHandleTypes）。这会生成一个CUlogicalEndpointFabricHandle，即代表该逻辑端点的互连网络（fabric）IPC 句柄。使用应用程序自行选择的进程间通信机制将其传输给对等方，然后让每个对等方使用cuLogicalEndpointImport导入它。导入 ID 在本地选择，使用cuLogicalEndpointImport导入它。导入方可以选择用于导入端点的逻辑端点 ID。此 ID 必须已成功预留，且不得与任何其他逻辑端点关联，但不必与所有者用于创建端点的 ID 相同。每个设备都可以在不同的 ID 处导入同一端点，甚至可以在多个不同的 ID 处导入它（即允许多次导入同一端点）。

下面的单播点对点示例在环上共享端点。每个秩导出自己的端点，并通过一次MPI_Sendrecv将句柄发送给接收邻居，同时收集发送邻居的句柄。它将该句柄导入为leId + sendRank，随后使用此 ID 针对所导入端点的绑定内存发出互连网络（fabric）操作：

  // Export our own endpoint to a fabric IPC handle, hand it to the peers that
  // need it, and import each peer's handle to a locally reserved id.
  CUlogicalEndpointFabricHandle localHandle;
  CUlogicalEndpointFabricHandle sendRankHandle;
  cuLogicalEndpointExport(&localHandle, leId + myRank, CU_LOGICAL_ENDPOINT_IPC_HANDLE_TYPE_FABRIC);
  MPI_Sendrecv(
      &localHandle, sizeof(CUlogicalEndpointFabricHandle), MPI_BYTE, recvRank, 0,
      &sendRankHandle, sizeof(CUlogicalEndpointFabricHandle), MPI_BYTE, sendRank, 0,
      MPI_COMM_WORLD, MPI_STATUS_IGNORE);
  cuLogicalEndpointImport(leId + sendRank, &sendRankHandle, CU_LOGICAL_ENDPOINT_IPC_HANDLE_TYPE_FABRIC);

4.18.8. 绑定内存

端点不拥有任何内存。绑定将端点偏移空间中的一个范围与物理分配关联，互连网络（fabric）操作仅对已绑定的范围有效。绑定按设备建立，因此单播绑定适用于所有者设备，而多播绑定适用于多播团队中的单个设备。

根据应用程序持有分配的方式，可以使用以下两种方式之一绑定内存：

• cuLogicalEndpointBindMem按分配句柄进行绑定，使用的是cuMemCreate返回的CUmemGenericAllocationHandle。

• cuLogicalEndpointBindAddr按映射的虚拟地址进行绑定。示例包括cudaMallocAsync返回的指针，以及使用cuMemMap映射的范围内的地址。有关当前支持的分配类型，请参阅 CUDA 内存管理 API。

两者都接受端点 ID、绑定适用的设备、端点偏移空间内的偏移量、要绑定的分配以及绑定的大小。cuLogicalEndpointBindMem还接受分配内的memOffset，从而允许绑定该句柄的一个子范围。端点偏移量、大小以及（对于cuLogicalEndpointBindMem）memOffset都必须是bindAlignment的整数倍，并且绑定范围必须位于端点大小所限定的范围内（请参阅第 4.18.6.1 节）。

  // Bind the whole backing allocation at endpoint offset 0, by allocation handle.
  cuLogicalEndpointBindMem(leId + myRank, cuDevice, 0, exportHandle, 0, exportSize, 0);
  // No rank may issue an operation until every destination has been bound.
  MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));

建立初始绑定后，参与的进程必须先同步，然后任何进程才能发出互连网络（fabric）操作。对于单播端点，每个导入方都必须在所有者绑定其支撑内存后与导出进程同步。对于多播端点，所有参与的进程都必须等待，直到每个参与方都已绑定其支撑内存。应用程序可以使用任何能够建立这种顺序的进程间同步机制。

由于对等方导入的是端点，而不是所有者的分配，因此所有者可以替换端点背后的内存，而无需任何对等方重新导入端点。使用cuLogicalEndpointUnbind解除当前范围的绑定，并通过同一个端点 ID 将新的分配绑定到端点偏移空间中的同一范围。所有者及其对等方必须在重新绑定前后进行同步，以确保在端点未绑定期间，没有对等方针对该端点发出操作。发出落在当前绑定范围之外的操作会产生未定义行为。

4.18.9. 互连网络（fabric）操作

线程通过执行接受计算互连传输句柄的互连网络（fabric）操作来访问端点资源，该句柄是一个逻辑端点 ID 与偏移量的组合。互连网络（fabric）操作是异步的，也就是说，程序必须显式等待它们完成，才能观察到它们的效果。互连网络（fabric）操作带有try_前缀，以指示它们可能失败。也就是说，在等待完成之后，线程必须检查该操作的完成机制（请参阅第 4.18.10 节），以确定操作是否成功。这使互连网络（fabric）操作有别于普通指针操作：基于指针的访问所发生的故障，是通过销毁上下文或整个应用程序来向应用程序报告的。想要处理这些错误的应用程序，例如要从指针故障中恢复的应用程序，只能在上下文或进程级别这样做。互连网络（fabric）操作允许应用程序在程序线程级别处理错误。

设备代码通过以下cuda::ptx指令发出互连网络（fabric）操作：

上述操作（try_atom除外）以 16 字节为单位移动数据，因此传输的两端都必须按 16 字节对齐：源指针和端点内的目标偏移量都必须是 16 字节的整数倍。

该示例发出互连网络（fabric）put 操作。put 从共享内存缓冲区读取数据，并将其写入目标句柄。多个线程将输入数据写入共享内存，然后使用__syncthreads与发出 put 的线程同步。put 通过异步代理读取共享内存，因此发出操作的线程调用cuda::ptx::fence_proxy_async(cuda::ptx::space_shared)，以将通用代理的共享内存写入同步到异步代理。put 接受目标句柄、源缓冲区基址、字节数以及用于跟踪完成情况的mbarrier：

    // Make the staged shared-memory source visible to the async proxy before the put reads it.
    cuda::ptx::fence_proxy_async(cuda::ptx::space_shared);
    cuda::ptx::fabric_try_put(cuda::ptx::space_shared, cuda::ptx::sem_relaxed,
                              cuda::ptx::scope_sys, putLeId, offset, smemBuf,
                              FABRIC_CHUNK_SIZE, &putBar);

该调用设置了若干指令限定符，其完整集合由 PTX ISA 规定，包括每条指令接受的访问大小、内存排序模式和作用域（请参阅互连网络（fabric）指令）。这里值得指出该操作的两个方面，因为它们将该操作与程序的其余部分联系起来：如何跟踪其完成情况，以及如何指定内存排序。

事务计数。使用mbarrier来指示完成的互连网络（fabric）操作遵循与使用 TMA 传输数据中描述的指针寻址cuda::ptx::cp_async_bulk复制相同的mbarrier事务记账模型。随着 put 完成，每交付 16 字节，它就将其mbarrier的事务计数减 1。在所有预期的到达都已发生且事务计数达到零后，屏障阶段完成。put 不会设置预期事务计数。在提交该操作之后发出的单独一个cuda::ptx::mbarrier_arrive_expect_tx会记录一次到达，并将事务计数增加字节数除以 16 的值，因为这就是 put 将减少的计数值。取回操作则使用基于字节的事务记账。cuda::ptx::fabric_try_get和cuda::ptx::fabric_try_pullred每交付一个字节就将事务计数减 1，因此cuda::ptx::mbarrier_arrive_expect_tx会加上完整的字节数。第 4.18.10 节介绍了 mbarrier 的布局要求，以及如何读回互连网络（fabric）状态。

带计数的 put 还会在其数据到达时递增目标端的字节计数器，因此接收方可以在访问本地数据之前轮询该计数器，而不必交换单独的完成消息。下面的代码使用cuda::ptx::fabric_try_put_counted（普通 put 的.counted::bytes变体）来实现这一点，并传入第二个目标偏移量，即计数器的端点偏移量：

    // Make the staged shared-memory source visible to the async proxy before the put reads it.
    cuda::ptx::fence_proxy_async(cuda::ptx::space_shared);
    cuda::ptx::fabric_try_put_counted(cuda::ptx::space_shared, cuda::ptx::sem_relaxed,
                                      cuda::ptx::scope_sys, putLeId, offset, counterOffset,
                                      smemBuf, FABRIC_CHUNK_SIZE, &putBar);

带计数的操作要求端点在创建时支持带计数操作（CU_LOGICAL_ENDPOINT_FLAG_COUNTED_OPS，请参阅第 4.18.2.2 节），并且目标计数器必须在端点内存中按 256 字节边界对齐。应用程序在布局端点的支撑内存时，通过在适当对齐的区域中为计数器预留空间来满足这一要求。

尽管该示例仅发出 put，但 get、归约和拉取归约操作使用第 4.18.10 节中描述的相同提交、状态报告和等待序列。如上所述，它们的事务计数单位取决于具体操作。

内存排序。该示例以sys作用域的relaxed排序发出其互连网络（fabric）put 操作。一次互连网络（fabric）put 会跨越多个代理：它通过互连网络（fabric）代理访问远程数据（对于带计数的 put，还会访问计数器），通过异步代理读取其.shared::cta源，并通过通用代理更新用于指示完成的mbarrier。由于 mbarrier 通过通用代理更新，因此只需使用线程块作用域的操作在屏障处等待，就能观察到完成。然而，与cuda::ptx::cp_async_bulk操作不同，观察到互连网络（fabric）操作完成，并不会使互连网络（fabric）代理的访问相对于通用代理有序。在使用指针通过通用代理访问数据之前，必须使用cuda::ptx::fence_proxy_generic_fabric_alias将互连网络（fabric）代理与通用代理同步。

仅有端点 ID 并不能提供对远程目标的指针访问。导入端点不会将其绑定的资源映射到导入方的虚拟地址空间。因此，下面的示例使用由发出操作的线程所在设备拥有的单播端点。leId是端点 ID，而绑定在dataOffset处的内存也映射到了该设备的虚拟地址空间中。dataPtr指向同一位置。

发出 put 后，线程等待其完成。在使用dataPtr通过通用代理读取目标之前，线程必须先使用cuda::ptx::sem_release发出cuda::ptx::fence_proxy_generic_fabric_alias，然后再使用cuda::ptx::sem_acquire发出。仅等待完成，并不能使互连网络（fabric）代理的写入排序在通用代理的指针加载之前。waitForSuccessfulCompletion代表第 4.18.10 节中描述的mbarrier等待和状态检查。目标的初始值为 0，smemSrcData是一个 16 字节的共享内存源，其第一个uint32_t元素保存着 42。

namespace ptx = cuda::ptx;

ptx::fabric_try_put(
    ptx::space_shared,
    ptx::sem_relaxed,
    ptx::scope_sys,
    leId, dataOffset,
    smemSrcData, 16, &mBar);
ptx::fabric_submit();
ptx::mbarrier_arrive_expect_tx(
    ptx::sem_acquire,
    ptx::scope_cta,
    ptx::space_shared,
    &mBar, 1);
waitForSuccessfulCompletion(&mBar);

// Order the subsequent generic-proxy load after the fabric put.
ptx::fence_proxy_generic_fabric_alias(ptx::sem_release);
ptx::fence_proxy_generic_fabric_alias(ptx::sem_acquire);

assert(*dataPtr == 42);

mbarrier 等待返回后，put 已完成，但互连网络（fabric）代理的写入相对于通用代理的加载仍然是无序的。释放栅栏和获取栅栏建立了这种顺序，因此指针加载会观察到值 42。

当一个线程使用互连网络（fabric）操作发布数据，而端点所有者设备上的另一个线程通过指针使用绑定目标中的数据时，也需要相同的释放/获取代理栅栏。在下面的消息传递示意示例中，发送方通过互连网络（fabric）代理使用互连网络（fabric）操作写入数据和信号标志。接收方使用通用代理通过指针轮询标志并读取数据。dataPtr和flagPtr对绑定到接收方所拥有端点的目标内存进行寻址。初始时，目标数据和标志均为 0。smemSrcData和smemSrcFlag是由四个uint32_t元素组成的 16 字节共享内存源。smemSrcData的第一个元素保存着 42，smemSrcFlag的第一个元素保存着 1。

namespace ptx = cuda::ptx;

ptx::fabric_try_put(
    ptx::space_shared,
    ptx::sem_relaxed,
    ptx::scope_sys,
    leId, dataOffset,
    smemSrcData, 16, &mBar);
ptx::fabric_submit();
ptx::mbarrier_arrive_expect_tx(
    ptx::sem_acquire,
    ptx::scope_cta,
    ptx::space_shared,
    &mBar, 1);
waitForSuccessfulCompletion(&mBar);

ptx::fence_proxy_generic_fabric_alias(
    ptx::sem_release);

ptx::fabric_try_put(
    ptx::space_shared,
    ptx::sem_relaxed,
    ptx::scope_sys,
    leId, flagOffset,
    smemSrcFlag, 16, &mBar);
ptx::fabric_submit();
ptx::mbarrier_arrive_expect_tx(
    ptx::sem_acquire,
    ptx::scope_cta,
    ptx::space_shared,
    &mBar, 1);
waitForSuccessfulCompletion(&mBar);

namespace ptx = cuda::ptx;

cuda::atomic_ref<uint32_t,
cuda::thread_scope_system> flag(*flagPtr);

while (flag.load(cuda::memory_order_relaxed) != 1);

// Order the data load after the observed flag update.
ptx::fence_proxy_generic_fabric_alias(
    ptx::sem_acquire);

assert(*dataPtr == 42);

数据 put 完成后，发送方先发出cuda::ptx::fence_proxy_generic_fabric_alias(cuda::ptx::sem_release)，然后再发出标志 put。当接收方的标志加载观察到该 put 写入的值 1 时，接收方就知道发送方已发布数据。在使用数据之前，接收方必须发出cuda::ptx::fence_proxy_generic_fabric_alias(cuda::ptx::sem_acquire)。此获取栅栏防止后续对数据的指针加载被排序在所观察到的标志更新之前。发送方的释放栅栏与接收方的获取栅栏共同使发送方的数据 put 排序在接收方的数据加载之前。因此，在没有任何其他数据写入的情况下，该加载会观察到值 42。

4.18.10. 完成状态与错误报告

互连网络（fabric）操作通过其完成对象报告状态。此处描述的操作使用共享内存中的mbarrier来表示完成，因此同一个mbarrier既发出完成信号，也记录状态。通过使用cuda::ptx::layout_v1调用cuda::ptx::mbarrier_init来初始化此mbarrier。这种布局为该对象扩展了一个保存互连网络（fabric）操作状态的字段。默认的cuda::ptx::layout_v0只跟踪到达计数和事务计数，无法承载该状态，因此不能用于此报告机制：

    cuda::ptx::mbarrier_init(cuda::ptx::layout_v1, &putBar, 1);

cuda::ptx::fabric_try_put接受跟踪该操作的 mbarrier。随着操作的数据送达，互连网络（fabric）引擎会在屏障上完成相应的事务，并在那里记录操作的状态。因此，传输已送达的信号来自 put 操作本身，而不是发出操作的线程。

put 操作会送达其字节，但不会告知屏障应预期多少字节。对于此处所示的逐线程 put 操作，发出操作的线程会在单独的步骤中提供该计数。它首先使用fabric_submit提交操作，以便互连网络（fabric）引擎最终开始处理该操作，然后调用cuda::ptx::mbarrier_arrive_expect_tx记录一次到达并增加预期事务计数，即该操作将送达的字节数，以 16 字节为单位计数。每个互连网络（fabric）操作都必须在跟踪其完成情况的屏障阶段推进之前提交。提交操作后，发出操作的线程要么执行屏障阶段推进所需的屏障操作（例如cuda::ptx::mbarrier_arrive_expect_tx或cuda::ptx::mbarrier_expect_tx），要么与执行该操作的另一个线程同步。发出操作的线程不必亲自在屏障上等待。另一个线程可以执行该等待。一旦所有预期的到达均已发生，且事务计数降为零，屏障阶段就会推进，等待随之返回。程序必须在网格退出之前等待每个互连网络（fabric）操作完成。网格退出时仍有操作未完成属于未定义行为。

以下序列提交该 put 操作并等待其完成：

    cuda::ptx::fabric_submit();
    cuda::ptx::mbarrier_arrive_expect_tx(cuda::ptx::sem_relaxed, cuda::ptx::scope_cta,
                                         cuda::ptx::space_shared, &putBar,
                                         (FABRIC_CHUNK_SIZE / FABRIC_TX_GRANULARITY_BYTES));
    // Spin while the barrier's phase hasn't flipped yet. The report outputs are
    // only valid once the barrier completes, so inspect them after.
    bool isReportSeen = false;
    uint8_t reportValue = 0;
    bool complete = waitWithDebugTimeout(
        [&]() {
          return !cuda::ptx::mbarrier_try_wait_parity(cuda::ptx::mbarrier_phase_primary, cuda::ptx::sem_relaxed,
                                                      cuda::ptx::scope_cta, isReportSeen, reportValue, &putBar,
                                                      /*phaseParity=*/0);
        },
        "cftTryPutKernel: mbarrier try_wait parity");
    if (!complete) { __trap(); }
    if (isReportSeen) {
      reportFabricError(&reportValue, offset);
    }

cuda::ptx::mbarrier_try_wait_parity的phase_type::primary形式除了具有完成谓词之外，还具有报告谓词和报告值目标操作数。如果完成谓词为假，则报告谓词和报告值包含未指定的值。当try_wait观察到阶段完成时，即完成谓词为真时，报告谓词和报告值共同包含有关该屏障阶段所跟踪操作的附加信息。如果该屏障阶段跟踪的所有操作都是互连网络（fabric）操作，则为假的报告谓词表示所有互连网络（fabric）操作均已成功。否则，如果报告谓词为真，则报告值可能包含有关失败的更多信息，程序可以使用cudaFabricOpErrorStatusCount和cudaFabricOpErrorStatusGet检查这些信息。

inline __device__ void reportFabricError(uint8_t *reportValue, uint64_t offset) {
  const unsigned long long off = offset;
  unsigned int errCount = 0;
  if (cudaFabricOpErrorStatusCount(reportValue, cudaFabricOpStatusSourceMbarrierV1, &errCount) != cudaSuccess) {
    printf("cudaFabricOpErrorStatusCount failed at offset %llu\n", off);
    return;
  }
  for (unsigned int i = 0; i < errCount; ++i) {
    cudaFabricOpStatusInfo info = cudaFabricOpStatusInfoSuccess;
    if (cudaFabricOpErrorStatusGet(reportValue, cudaFabricOpStatusSourceMbarrierV1, i, &info) == cudaSuccess) {
      printf("fabric put error[%u/%u] at offset %llu: status=%d\n", i, errCount, off, (int)info);
    } else {
      printf("cudaFabricOpErrorStatusGet failed for error[%u/%u] at offset %llu\n", i, errCount, off);
    }
  }
}

解码后的状态指出操作失败的原因，使应用程序能够恢复，而不是导致进程发生故障。对于诸如数据包丢失或链路重置之类的瞬时互连网络（fabric）错误，可以通过重试、重新路由或应用程序专用的回退方式来处理。该示例打印解码后的错误，而通信库通常会针对这些错误采取行动，以保持向前进展。

4.18.11. 清理

清理需要在可能向某个端点发出操作的任何进程与已向该端点绑定资源的任何进程之间进行同步。对于单播端点，这意味着，只要导入方仍有可能向该端点发出操作，所有者就不得解除绑定或销毁该端点，也不得释放其后备资源。对于多播端点，每个参与进程都拥有其绑定的副本，因此，只要另一个进程仍有可能向该端点发出操作，任何参与者就都不得解除绑定或释放自己的副本。重要的是顺序，而不是强制执行该顺序的机制：只要进程间同步能在移除绑定之前建立这一顺序，就足够了。

  // ---- Clean up endpoints and free the backing allocation. ----
  // Once every rank's stream has synchronized, all fabric puts
  // have completed and no rank will issue another operation. Making that global
  // with a barrier ensures no peer still targets an endpoint we clean up below.
  MPI_Barrier(MPI_COMM_WORLD);

  // Destroy the imported endpoint (the peer target our puts addressed). An import
  // is only a local reference to the peer's endpoint id and maps no peer memory,
  // so once our puts have drained this just drops that reference.
  cuLogicalEndpointDestroy(leId + sendRank);

  // Unbind and destroy our owned endpoint, then release the reserved id range.
  cuLogicalEndpointUnbind(leId + myRank, cuDevice, 0, exportSize);
  cuLogicalEndpointDestroy(leId + myRank);
  cuLogicalEndpointIdRelease(leId, numRanks);

  // Free the backing allocation in VMM cleanup order: unmap the range, release
  // the physical handle, then free the reserved virtual address range.
  cuMemUnmap(exportPtr, exportSize);
  cuMemRelease(exportHandle);
  cuMemAddressFree(exportPtr, exportSize);

等待互连网络（fabric）操作和所有发出方停止。在释放跟踪互连网络（fabric）操作的完成对象之前，必须在该对象上等待。对于网格发出的互连网络（fabric）操作，网格中的一个线程必须在网格退出之前等待完成。未执行此等待就退出属于未定义行为。在该示例中，每个核函数都在返回之前等待完成，而流同步则确认这些核函数及其等待均已完成。

另一个进程发出的、以绑定到同一逻辑端点的资源为目标的互连网络（fabric）操作，可能仍在进行中。因此，在移除单播绑定之前，所有者必须等到每个导入方都已停止发出操作。在移除任何多播绑定之前，每个拥有已绑定副本的参与者都必须知道，所有进程均已停止发出操作。该示例在所有 rank 同步各自的 CUDA 流之后，使用一个屏障实现这一点。如果能够保持所要求的顺序，应用程序可以采用范围更小的同步方案。

销毁导入的端点 ID。每个导入进程通过调用cuLogicalEndpointDestroy销毁与其导入端点关联的逻辑端点 ID。这只会移除该本地别名，不会解除所有者上的任何绑定，因此销毁操作本身不需要跨进程协调。

销毁所有者的端点 ID。销毁所有者的逻辑端点 ID 也会解除此进程通过该 ID 绑定的任何内存的绑定，因此先调用cuLogicalEndpointUnbind是可选的，但建议这样做。一旦所有别名，即所有者的 ID 和任何导入的别名，都已销毁，端点的资源就会被释放。

单播端点和多播端点适用相同的处理流程。由于多播成员关系是永久性的，因此更改多播组需要销毁端点并创建一个新端点。只要端点仍然存在，设备就无法离开该组。

重用或释放端点 ID。使用cuLogicalEndpointDestroy销毁端点关联后，其 ID 仍处于保留状态，但不再有关联，可以立即重用。要将 ID 归还系统，请对之前使用cuLogicalEndpointIdReserve保留的范围调用cuLogicalEndpointIdRelease。

释放后备内存分配。端点并不拥有其后备内存，因此该内存分配需要单独释放。应使用与内存分配方式相对应的释放机制。对于 VMM 分配，通常的顺序是先使用cuMemUnmap移除映射，再使用cuMemRelease释放物理句柄，最后使用cuMemAddressFree归还保留的虚拟地址范围。程序不得释放仍绑定到端点的内存分配。