---
GPU内からROCEの通信（GPU Direct）はどのように行われているのか？
https://deepwiki.com/search/ib-verbs_2ced64b3-c70c-4606-abe3-71e71393876c
GPU内からRoCE（RDMA over Converged Ethernet）の通信（GPU Direct）について、NCCLのコードベースから具体的な実装を説明します。
## GPU DirectとRoCEの統合
NCCLでは、GPU DirectとRoCEの組み合わせにより、GPU内のデータを直接ネットワーク経由で送信することが可能です。これは主に`src/transport/net_ib.cc`で実装されています。
## GDR（GPU Direct RDMA）サポートの検出
まず、システムがGDRをサポートしているかを確認します： [1](#2-0)
この関数では、nv_peer_memモジュールやnvidia_peermemモジュールの存在を確認してGDRサポートを判定しています。
## DMA-BUFサポートの検出
CUDA 11.7以降では、DMA-BUFを使用したより効率的なGPU Directが利用可能です： [2](#2-1)
## RoCEデバイスの識別
RoCEとInfiniBandを区別するため、リンク層の情報を取得します： [3](#2-2)
この定義により、`IBV_LINK_LAYER_ETHERNET`の場合はRoCEとして識別されます。
## GPU Directメモリ登録
GPU内のメモリをRoCEネットワークで使用するため、メモリ登録が行われます： [4](#2-3)
DMA-BUFがサポートされている場合は、より効率的な登録方法が使用されます： [5](#2-4)
## トポロジーベースのGDR判定
システムトポロジーに基づいて、GPU DirectとRoCEの組み合わせが有効かを判定します： [6](#2-5)
この関数では、GPUとNICの距離、C2Cシステムでの特別な処理、PCI経由のアクセスなどを考慮してGDRの使用可否を決定します。
## RoCE接続の確立
RoCE接続では、GIDとEthernetレイヤーの情報が使用されます： [7](#2-6)
## GPU Directの実際の通信
GPU内のデータは、RDMA書き込み操作を通じて直接ネットワーク経由で送信されます： [8](#2-7)
## Notes
GPU DirectとRoCEの組み合わせにより、GPU内のデータをCPUを経由せずに直接ネットワーク経由で送信できます。これは、nv_peer_memモジュールやDMA-BUFサポートを通じて実現され、システムトポロジーに基づいて最適な方法が選択されます。RoCEの場合は、Ethernetレイヤー上でRDMA操作が実行され、高いスループットと低いレイテンシを実現しています。
Wiki pages you might want to explore:
- [Transport Layer (NVIDIA/nccl)](/wiki/NVIDIA/nccl#2.2)
- [Topology System (NVIDIA/nccl)](/wiki/NVIDIA/nccl#2.3)
```
ncclResult_t ncclIbMultiSend(struct ncclIbSendComm* comm, int slot, void* pHandle) {
struct ncclIbRequest** reqs = comm->fifoReqs[slot];
volatile struct ncclIbSendFifo* slots = comm->fifo[slot];
int nreqs = slots[0].nreqs;
if (nreqs > NCCL_NET_IB_MAX_RECVS) return ncclInternalError;
uint64_t wr_id = 0ULL;
for (int r=0; r<nreqs; r++) {
struct ibv_send_wr* wr = comm->wrs+r;
memset(wr, 0, sizeof(struct ibv_send_wr));
struct ibv_sge* sge = comm->sges+r;
sge->addr=(uintptr_t)reqs[r]->send.data;
wr->opcode = IBV_WR_RDMA_WRITE;
wr->send_flags = 0;
wr->wr.rdma.remote_addr = slots[r].addr;
wr->next = wr + 1;
wr_id += (reqs[r] - comm->base.reqs) << (r*8);
#ifdef NCCL_ENABLE_NET_PROFILING
reqs[r]->pInfo[0].nEventHandles = 0;
#endif
}
// Write size as immediate data. In the case of multi-send, only write
// 0 or 1 as size to indicate whether there was data sent or received.
uint32_t immData = 0;
if (nreqs == 1) {
immData = reqs[0]->send.size;
} else {
int* sizes = comm->remSizesFifo.elems[slot];
for (int r=0; r<nreqs; r++) sizes[r] = reqs[r]->send.size;
comm->remSizesFifo.sge.addr = (uint64_t)sizes;
comm->remSizesFifo.sge.length = nreqs*sizeof(int);
}
```