RDMA读写数据放置/可见性语义及GPU Direct适配问题咨询

1. RDMA Read 完成语义相关 #

• 常规主机内存场景：当完成队列（CQ）上报RDMA READ完成通知时，读取的数据已经完全写入本地缓冲区，可以直接由CPU访问，无额外可见性问题。这是IBTA规范明确要求的语义：RDMA READ的完成事件上报的前提是所有数据已经落地到本地目标地址，不存在端侧缓存未刷新的情况。
• GPU Direct RDMA（GDR）映射GPU显存场景：该行为和主机内存场景不完全一致。原因是部分GPU架构的显存控制器存在DMA写入队列，网卡通过PCIe写入GPU显存的数据可能还滞留在GPU的PCIe接收缓冲区/显存写入队列中，并没有真正写入显存地址对应的物理存储单元，此时GPU内核直接访问该地址可能读到旧数据。
  • 保障数据可见性的操作：可以在GPU侧执行一次cudaDeviceSynchronize()，或者针对对应显存区域调用cudaStreamWaitEvent()绑定同步事件；也可以在申请GDR显存时开启CU_MEM_ALLOC_GPU_DIRECT_RDMA_WRITE_MEMORY_BARRIER标志（针对安培及以后架构的GPU），强制网卡DMA写入完成后自动触发显存写屏障，保证完成事件上报时数据对GPU可见。

2. 带立即数的RDMA Write/Write+Send语义相关 #

• 常规主机内存场景：当远端接收队列（RQ）观测到带立即数的RDMA Write完成事件（或者Write + Send组合的Send完成事件）时，所有对应写入的数据已经完全落地到远端目标主机内存，可以直接校验数据存在性和正确性。IBTA规范要求带立即数的RDMA Write的立即数投递，必须晚于所有写入数据写入目标内存之后，不会出现立即数先到、数据还未写完的乱序情况。
• GDR映射GPU显存场景：该行为逻辑和主机内存场景同样存在差异。和RDMA Read的GDR场景原因一致，网卡写入GPU显存的数据可能滞留在GPU的内部缓冲区，此时即使立即数已经上报到接收队列，GPU侧直接访问对应显存地址可能读到不完整或者旧的数据。
  • 处理方式和GDR场景下的RDMA Read一致：通过CUDA提供的同步原语做显式同步，或者开启对应显存分配的屏障标志，保证数据可见性。

内容的提问来源于stack exchange，提问作者user718134