libfabric transport with the verbs provider hangs at the first atomic when the symmetric heap is in device memory (aarch64 / GH200) #128

Environment: 2 nodes, 1× GH200 (96 GB) + 1× ConnectX-7 (InfiniBand, NDR) each, aarch64, CUDA 13.2, driver 595.58, rdma-core 2601.0.7 (DOCA-OFED 26.01), libfabric 2.7.0 built from source with --enable-verbs --with-cuda, NVSHMEM devel @ 2dca770 built with NVSHMEM_LIBFABRIC_SUPPORT=ON, MPI bootstrap, 2 PEs (1 per node). nvidia_peermem loaded, no gdrcopy.

Repro:

NVSHMEM_REMOTE_TRANSPORT=libfabric NVSHMEM_LIBFABRIC_PROVIDER=verbs \
mpirun -np 2 ... perftest/device/pt-to-pt/shmem_put_latency -b 8 -e 8 -n 2 -w 1

nvshmem_init completes (Selected provider verbs;ofi_rxm, ..., hmem yes, dmabuf fd obtained for the 512 MB heap chunk, rxm connreq exchanged), then the first barrier never returns. With NVSHMEM_DEBUG=INFO the initiator's proxy reports

process_channel_amo/proxy_update_processed processed 40 complete 0

and the target's libfabric core prints

ofi_cudaMemcpy():239<warn> Either dst or src buffer of cudaMemcpy is not allocated or registered by cuda device. cudaDeviceSynchronize() will be performed to ensure data consistency.

after which there is no further output; the job has to be killed.

What does not change the outcome: NVSHMEM_LIBFABRIC_ENABLE_DMABUF=1, NVSHMEM_LIBFABRIC_DISABLE_BATCH_RMA=1, exporting MLX5_SCATTER_TO_CQE=0 before launch, FI_LOG_LEVEL=debug (no atomic-related messages from rxm).

What works: the same command with NVSHMEM_HEAP_KIND=SYSMEM (heap in Grace memory) completes normally (8 B put latency reported, 12 sizes). The IBRC, IBDevX, IBGDA and staged transports all pass the same perftests on this pair with the device heap.

Reading of the code: nvshmemt_libfabric_amo posts fi_atomicmsg(..., FI_INJECT) (src/modules/transport/libfabric/libfabric.cpp:2438-2444) and requests FI_ATOMIC for verbs (:3603). The verbs provider advertises FI_ATOMIC, but with ofi_rxm on top, atomics targeting an FI_HMEM buffer are emulated in software at the target, which is where the ofi_cudaMemcpy warning comes from; the completion never reaches the initiator. So either the transport should not select the rxm software-atomic path for device-memory targets with verbs, or the documentation should say that NVSHMEM_LIBFABRIC_PROVIDER=verbs requires a host-memory heap (or gdrcopy) for atomics.

Happy to run further experiments on this pair (e.g. with a libfabric build that logs rxm atomic handling) if that helps.