[Issue]: Maybe we can refine the logic for whether atomicAdd can be used directly in nvshmemi_signal_op ?

[Lyapunova-cd]

How is this issue impacting you?

Application hang

Share Your Debug Logs

RT, when NVSHMEM_DISABLE_P2P=1 and ibgda enabled, we still use ibgda atomic op but we don't create qp for the local access? And I think that we can use atomicAdd directly for local HBM and atomicAdd for remote HBMs when nvlink is connected between mype and dstpe

Steps to Reproduce the Issue

No response

NVSHMEM Version

3.3.9+cuda12.2

Your platform details

No response

Error Message & Behavior

No response

[benhg]

And I think that we can use atomicAdd directly for local HBM and atomicAdd for remote HBMs when nvlink is connected between mype and dstpe

If the job has NVSHMEM_DISABLE_P2P=1, we can't use P2P (NVLink or PCIe load/store) in order to do the operation, because the user has disabled it.

RT, when NVSHMEM_DISABLE_P2P=1 and ibgda enabled, we still use ibgda atomic op but we don't create qp for the local access?

In the code you are showing, the nvshmemi_transfer_amo_nonfetch will end up doing an RDMA atomic operation. Are you seeing an application hang when using these RDMA atomic operations? Could you capture logs with NVSHMEM_DEBUG=INFO and attach them ?

[carrroo]

I wrote a simple test case for the nvshmem_int_atomic_add interface.
My NVSHMEM version is 3.3.24.
If I set both NVSHMEM_DISABLE_P2P=1 and NVSHMEM_IB_ENABLE_IBGDA=1, the program hangs.
Without these two settings the test passes and produces the expected output:

[PE 0] local value 5  (PASS)                                             
[PE 1] local value 5  (PASS)                                             
[PE 0] final value 8  (PASS)                                             
[PE 1] final value 8  (PASS)   

The test code is shown below.
The hang log is displayed after the code.

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <mpi.h>
#include <nvshmem.h>
#include <nvshmemx.h>
#include <cuda_runtime.h>

#define MPI_CHECK(stmt)                                                  \
    do {                                                                 \
        int _e = stmt;                                                   \
        if (_e != MPI_SUCCESS) {                                         \
            fprintf(stderr, "[%s:%d] MPI error %d\n", __FILE__, __LINE__, _e); \
            exit(EXIT_FAILURE);                                          \
        }                                                                \
    } while (0)

#define CUDA_CHECK(stmt)                                                 \
    do {                                                                 \
        cudaError_t _e = stmt;                                           \
        if (_e != cudaSuccess) {                                         \
            fprintf(stderr, "[%s:%d] CUDA error %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(_e));                             \
            exit(EXIT_FAILURE);                                          \
        }                                                                \
    } while (0)

/* ---------------- device kernel ---------------- */
__global__ void kernel_local(int *dest, int val)
{
    nvshmem_int_atomic_add(dest, val, nvshmem_my_pe());
}

__global__ void kernel_remote(int *dest, int val, int pe)
{
    nvshmem_int_atomic_add(dest, val, pe);
}

/* ---------------- main ---------------- */
int main(int argc, char *argv[])
{
#if 1
    setenv("NVSHMEM_DISABLE_P2P", "1", 1);
    setenv("NVSHMEM_IB_ENABLE_IBGDA", "1", 1);
    setenv("NVSHMEM_IBGDA_FORCE_NIC_BUF_MEMTYPE", "hostmem", 1);
    setenv("NVSHMEM_IBGDA_NIC_HANDLER", "cpu", 1);
#endif
    MPI_CHECK(MPI_Init(&argc, &argv));
    int rank, size;
    MPI_CHECK(MPI_Comm_rank(MPI_COMM_WORLD, &rank));
    MPI_CHECK(MPI_Comm_size(MPI_COMM_WORLD, &size));
    if (size != 2) {
        if (rank == 0)
            fprintf(stderr, "Please run with exactly 2 MPI ranks\n");
        MPI_CHECK(MPI_Finalize());
        return 0;
    }
    
    nvshmemx_init_attr_t attr = NVSHMEMX_INIT_ATTR_INITIALIZER;
    MPI_Comm comm = MPI_COMM_WORLD;
    attr.mpi_comm = &comm;
    if (nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM, &attr) != 0) {
        fprintf(stderr, "NVSHMEM init failed\n");
        exit(EXIT_FAILURE);
    }
    
    int mype = nvshmem_my_pe();
    int npes = nvshmem_n_pes();
    int node_pe = nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE);
    printf("[Host] PE %d of %d, node PE %d\n", mype, npes, node_pe);
    
    CUDA_CHECK(cudaSetDevice(node_pe));
    
    int *sym = (int *)nvshmem_malloc(sizeof(int));
    if (!sym) { fprintf(stderr, "nvshmem_malloc failed\n"); exit(1); }
    
    int h = 0;
    CUDA_CHECK(cudaMemcpy(sym, &h, sizeof(int), cudaMemcpyHostToDevice));
    
    /* local add 5 */
    kernel_local<<<1,1>>>(sym, 5);
    CUDA_CHECK(cudaDeviceSynchronize());
    nvshmem_barrier_all();
    
    CUDA_CHECK(cudaMemcpy(&h, sym, sizeof(int), cudaMemcpyDeviceToHost));
    int ok = (h == 5);
    printf("[PE %d] local value %d  (%s)\n", mype, h, ok ? "PASS" : "FAIL");
    
    /* remote add 3 */
    int peer = 1 - mype;
    kernel_remote<<<1,1>>>(sym, 3, peer);
    CUDA_CHECK(cudaDeviceSynchronize());
    nvshmem_barrier_all();
    
    CUDA_CHECK(cudaMemcpy(&h, sym, sizeof(int), cudaMemcpyDeviceToHost));
    ok &= (h == 8);
    printf("[PE %d] final value %d  (%s)\n", mype, h, ok ? "PASS" : "FAIL");
    
    nvshmem_free(sym);
    nvshmem_finalize();
    MPI_CHECK(MPI_Finalize());
    return ok ? 0 : 1;
}

root@NH-DC-NM129-I06-12U-GPU-246:/home/nvshmem_test/build# mpirun --allow-run-as-root -n 2 test_nvshmem_int_atomic_add                                                                                                            
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO PE distribution has been identified as NVSHMEMI_PE_DIST_BLOCK                                                                                                                     
NVSHMEM configuration:                                                                                                                                                                                                                     
  CUDA API                     12090                                                                                                                                                                                                       
  CUDA Runtime                 12090                                                                                                                                                                                                       
  CUDA Driver                  12030                                                                                                                                                                                                       
  Build Timestamp              Aug 19 2025 15:40:54                                                                                                                                                                                        
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [0] NVSHMEM INFO PE distribution has been identified as NVSHMEMI_PE_DIST_BLOCK                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [0] NVSHMEM INFO PE 1 (process) affinity to 2 CPUs:                                                                                                                                                
     1 65                                                                                                                                                                                                                                  
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [0] NVSHMEM INFO cudaDriverVersion 12030                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [0] NVSHMEM INFO NVSHMEM symmetric heap kind = DEVICE selected                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [0] NVSHMEM INFO nvshmemi_common_init failed, continuing                                                                                                                                           
[Host] PE 1 of 2, node PE 1                                                                                                                                                                                                                
  Build Variables                                                                                                                                                                                                                          
        NVSHMEM_DEBUG=OFF NVSHMEM_DEVEL=OFF NVSHMEM_DEFAULT_PMI2=OFF                                                                                                                                                                       
        NVSHMEM_DEFAULT_PMIX=OFF NVSHMEM_DEFAULT_UCX=OFF                                                                                                                                                                                   
        NVSHMEM_ENABLE_ALL_DEVICE_INLINING=OFF NVSHMEM_GPU_COLL_USE_LDST=OFF                                                                                                                                                               
        NVSHMEM_IBGDA_SUPPORT=ON NVSHMEM_IBGDA_SUPPORT_GPUMEM_ONLY=OFF                                                                                                                                                                     
        NVSHMEM_IBDEVX_SUPPORT=ON NVSHMEM_IBRC_SUPPORT=ON NVSHMEM_LIBFABRIC_SUPPORT=ON                                                                                                                                                     
        NVSHMEM_MPI_SUPPORT=ON NVSHMEM_NVTX=ON NVSHMEM_PMIX_SUPPORT=ON                                                                                                                                                                     
        NVSHMEM_SHMEM_SUPPORT=ON NVSHMEM_TEST_STATIC_LIB=OFF                                                                                                                                                                               
        NVSHMEM_TIMEOUT_DEVICE_POLLING=OFF NVSHMEM_TRACE=OFF NVSHMEM_UCX_SUPPORT=ON                                                                                                                                                        
        NVSHMEM_USE_DLMALLOC=OFF NVSHMEM_USE_NCCL=ON NVSHMEM_USE_GDRCOPY=ON                                                                                                                                                                
        NVSHMEM_VERBOSE=OFF                                                                                                                                                                                                                
        CUDA_HOME=/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/build/cuda                                                                                                                                                                      
        GDRCOPY_HOME=/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/gdrcopy_build                                                                                                                                                                
        LIBFABRIC_HOME=/data/tmp/libfabric_install MPI_HOME=/data/tmp/openmpi_install                                                                                                                                                      
        NCCL_HOME=/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/nccl/build                                                                                                                                                                      
        NVSHMEM_PREFIX=/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/install                                                                                                                                                       
        PMIX_HOME=/data/tmp/pmix_install SHMEM_HOME=/data/tmp/openmpi_install                                                                                                                                                              
        UCX_HOME=/dvs/p4/build/sw/gpgpu/openucx/build                                                                                                                                                                                      
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO PE 0 (process) affinity to 2 CPUs:                                                                                                                                                
     0 64                                                                                                                                                                                                                                  
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO cudaDriverVersion 12030                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEM symmetric heap kind = DEVICE selected                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO nvshmemi_common_init failed, continuing                                                                                                                                           
[Host] PE 0 of 2, node PE 0                                                                                                                                                                                                                
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEM symmetric heap kind = DEVICE selected                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO [1] nvshmemi_get_cucontext->cuCtxSynchronize->CUDA_SUCCESS) my_stream (nil)                                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO in get_cucontext, queried and saved context for device: 1 context: 0x55cda4391f10                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO [1] Created stream 0x55cda4c78e70 for device 1                                                                                                                                    
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO CUDA 64-bit stream memops support is not available                                                                                                                                
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVML library found. libnvidia-ml.so.1                                                                                                                                             
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEM symmetric heap kind = DEVICE selected                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO [0] nvshmemi_get_cucontext->cuCtxSynchronize->CUDA_SUCCESS) my_stream (nil)                                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO in get_cucontext, queried and saved context for device: 0 context: 0x563ad520a2a0                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO [0] Created stream 0x563ad5afb5b0 for device 0                                                                                                                                    
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO CUDA 64-bit stream memops support is not available                                                                                                                                
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVML library found. libnvidia-ml.so.1                                                                                                                                             
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVLS: cuMulticast is not supported on CUDA or disabled by user                                                                                                                    
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVLS: cuMulticast is not supported on CUDA or disabled by user                                                                                                                    
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO ALGO: FCOLLECT_ALGO set to 0                                                                                                                                                      
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO ALGO: BCAST_ALGO set to 0                                                                                                                                                         
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO ALGO: REDUCE_ALGO set to 0 (0 -> 1)                                                                                                                                               
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO ALGO: REDMAXLOC_ALGO set to 1                                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO ALGO: REDUCESCATTER_ALGO set to 0                                                                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO ALGO: FCOLLECT_ALGO set to 0                                                                                                                                                      
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO ALGO: BCAST_ALGO set to 0                                                                                                                                                         
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO ALGO: REDUCE_ALGO set to 0 (0 -> 1)                                                                                                                                               
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO ALGO: REDMAXLOC_ALGO set to 1                                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO ALGO: REDUCESCATTER_ALGO set to 0                                                                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO team psync mem req 509952 bytes, team mem total req 130548520 bytes, max teams 32                                                                                                 
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:94 NVSHMEM WARN P2P access was disabled in the environment                                              
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO team psync mem req 509952 bytes, team mem total req 130548520 bytes, max teams 32                                                                                                 
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:94 NVSHMEM WARN P2P access was disabled in the environment                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1712 Begin - Enumerating IB devices in the system ([]) -                                                         
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1712 Begin - Enumerating IB devices in the system ([]) -                                                         
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=0 (of 10), name=mlx5_cx4lx_5, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=1 (of 10), name=mlx5_cx4lx_4, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=2 (of 10), name=mlx5_cx6_3, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=0 (of 10), name=mlx5_cx4lx_5, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=3 (of 10), name=mlx5_cx6_2, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=1 (of 10), name=mlx5_cx4lx_4, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=4 (of 10), name=mlx5_cx4lx_3, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=2 (of 10), name=mlx5_cx6_3, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=5 (of 10), name=mlx5_cx4lx_2, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=6 (of 10), name=mlx5_cx4lx_1, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=3 (of 10), name=mlx5_cx6_2, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=7 (of 10), name=mlx5_cx4lx_0, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=4 (of 10), name=mlx5_cx4lx_3, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=8 (of 10), name=mlx5_cx6_1, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=5 (of 10), name=mlx5_cx4lx_2, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=9 (of 10), name=mlx5_cx6_0, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1818 End - Enumerating IB devices in the system                                                                                                  
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1823 Begin - Ordered list of devices for assignment (after processing user provdied env vars (if any))  -                                        
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=0 (of 5), device id=2, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=1 (of 5), device id=3, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=2 (of 5), device id=4, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=3 (of 5), device id=8, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=4 (of 5), device id=9, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1831 End - Ordered list of devices for assignment (after processing user provdied env vars (if any))                                             
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 215 /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:1871 Ib Alloc Size 2097152 pointer 0x7f97b0bff000  
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=6 (of 10), name=mlx5_cx4lx_1, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibgda/ibgda.cpp:3693: non-zero status: 3 NVSHMEM_IBGDA_NIC_HANDLER=cpu requires GDRCopy.                                                                       
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:nvshmemi_transport_init:282: init failed for transport: IBGDA                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO Successfully initialized the transport: ibrc                                                                                                                                      
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=7 (of 10), name=mlx5_cx4lx_0, num_ports=1                                                   
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=8 (of 10), name=mlx5_cx6_1, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1733 Enumerated IB devices in the system - device id=9 (of 10), name=mlx5_cx6_0, num_ports=1                                                     
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1818 End - Enumerating IB devices in the system                                                                                                  
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1823 Begin - Ordered list of devices for assignment (after processing user provdied env vars (if any))  -                                        
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=0 (of 5), device id=2, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=1 (of 5), device id=3, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=2 (of 5), device id=4, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=3 (of 5), device id=8, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1827 Ordered list of devices for assignment - idx=4 (of 5), device id=9, port_num=1                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 1831 End - Ordered list of devices for assignment (after processing user provdied env vars (if any))                                             
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 215 /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:1871 Ib Alloc Size 2097152 pointer 0x7f5764581000  
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibgda/ibgda.cpp:3693: non-zero status: 3 NVSHMEM_IBGDA_NIC_HANDLER=cpu requires GDRCopy.                                                                       
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:nvshmemi_transport_init:282: init failed for transport: IBGDA                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO Successfully initialized the transport: ibrc                                                                                                                                      
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEM_ENABLE_NIC_PE_MAPPING = 0, device 0 setting dev_id = 0                                                                                                                    
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEM_ENABLE_NIC_PE_MAPPING = 0, device 0 setting dev_id = 0                                                                                                                    
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 105 NVSHMEM_IB_ADDR_FAMILY set by environment to AF_INET                                                                          
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 128 NVSHMEM_IB_ADDR_RANGE set by environment to ::/0                                                                              
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.h 141 NET/IB: Ip address '::' is invalid for family AF_INET, ignoring address                                                       
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO [0] status 0 cudaErrorInvalidValue 1 cudaErrorInvalidSymbol 13 cudaErrorInvalidMemcpyDirection 21 cudaErrorNoKernelImageForDevice 209                                             
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO [1] status 0 cudaErrorInvalidValue 1 cudaErrorInvalidSymbol 13 cudaErrorInvalidMemcpyDirection 21 cudaErrorNoKernelImageForDevice 209                                             
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.cpp 106 ibv_reg_mr handle 0x7ffe79a260c0 handle->mr (nil)                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 215 /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:594 Ib Alloc Size 8 pointer 0x55cda6453000         
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_ib_common.cpp 106 ibv_reg_mr handle 0x7ffc39453cf0 handle->mr (nil)                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp 215 /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:594 Ib Alloc Size 8 pointer 0x563ad72d6000         
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO P2P list: 1                                                                                                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO P2P list: 0                                                                                                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEM_TEAM_SHARED: start=1, stride=1, size=1                                                                                                                                    
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEM_TEAM_SHARED: start=0, stride=1, size=1                                                                                                                                    
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEMX_TEAM_NODE: start=0, stride=1, size=2                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEMX_TEAM_NODE: start=0, stride=1, size=2                                                                                                                                     
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEMX_TEAM_SAME_MYPE_NODE: start=1, stride=2, size=1                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEMX_TEAM_SAME_MYPE_NODE: start=0, stride=2, size=1                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEMI_TEAM_SAME_GPU: start=1, stride=1, size=1                                                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEMI_TEAM_SAME_GPU: start=0, stride=1, size=1                                                                                                                                 
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO NVSHMEMI_TEAM_GPU_LEADERS: start=0, stride=1, size=2                                                                                                                              
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] NVSHMEM INFO NVSHMEMI_TEAM_GPU_LEADERS: start=0, stride=1, size=2                                                                                                                              
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 1                               
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 3                               
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 4                               
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 1                               
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 3                               
                                                                                                                                                                                                                                           
                                                                                                                                                                                                                                           
NH-DC-NM129-I06-12U-GPU-246:83320:83320 [0] /dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/team/team_internal.cpp:759 NVSHMEM WARN Skipping NVLINK SHARP resource initialized for team ID: 4                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:nvshmemt_ibrc_amo:1366: RMA/AMO verb 4 not implemented                                                                                           
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibrc/ibrc.cpp:nvshmemt_ibrc_amo:1366: RMA/AMO verb 4 not implemented 

[benhg]

@carrroo , thank you for your test program. I haven't run it yet but from the logs I believe something a bit different is going on.

First, because you don't have GDRCopy set up on your system, IBGDA is actually failing to initialize and the system will fall back to using IBRC. Can you verify whether the behavior is the same on the IBRC transport (that is, leave P2P disabled via environment variable, but don't enable IBGDA). The relevant logs are these:

/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibgda/ibgda.cpp:3693: non-zero status: 3 NVSHMEM_IBGDA_NIC_HANDLER=cpu requires GDRCopy.                                                                       
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:nvshmemi_transport_init:282: init failed for transport: IBGDA                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO Successfully initialized the transport: ibrc       

Next, the error you're getting RMA/AMO verb 4 not implemented 4 is coming from IBRC, because any operation that isn't an RDMA read (nvshmem_get/g), RDMA_write (nvshmem_put) or RDMA_SEND_INLINE (nvshmem_p) is not supported by that transport (for RMA), while only SIGNAL_SET and FETCH_ADD are implemented for AMOs. You can see ibrc.cpp around L1100-1200 for the implementation details. Interestingly, op 4 is NVSHMEMI_AMO_ADD so I'm not sure exactly what's going on there.

[kwu130]

How is this issue impacting you?

Application hang

Share Your Debug Logs

RT, when NVSHMEM_DISABLE_P2P=1 and ibgda enabled, we still use ibgda atomic op but we don't create qp for the local access? And I think that we can use atomicAdd directly for local HBM and atomicAdd for remote HBMs when nvlink is connected between mype and dstpe

### Steps to Reproduce the Issue _No response_

NVSHMEM Version

3.3.9+cuda12.2

Your platform details

No response

Error Message & Behavior

No response

I think we can change if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) to if (peer_base_addr != NULL) , because peer_base_addr != NULLimplies that peer_base_addr can be accessed directly in local memory rather than through ibgda.

I'm not sure if this understanding is correct, and I look forward to any feedback.

[kwu130]

@carrroo , thank you for your test program. I haven't run it yet but from the logs I believe something a bit different is going on.

First, because you don't have GDRCopy set up on your system, IBGDA is actually failing to initialize and the system will fall back to using IBRC. Can you verify whether the behavior is the same on the IBRC transport (that is, leave P2P disabled via environment variable, but don't enable IBGDA). The relevant logs are these:首先，由于您的系统上未设置 GDRCopy，IBGDA 实际上无法初始化，系统将回退到使用 IBRC。您能否验证 IBRC 传输上的行为是否相同（即通过环境变量禁用 P2P，但不启用 IBGDA）。相关日志如下：

/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/common/transport_gdr_common.cpp 26 GDRCopy library not found. disabling GDRCopy.                                                                               
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/modules/transport/ibgda/ibgda.cpp:3693: non-zero status: 3 NVSHMEM_IBGDA_NIC_HANDLER=cpu requires GDRCopy.                                                                       
                                                                                                                                                                                                                                           
/dvs/p4/build/sw/rel/gpgpu/toolkit/r12.9/main_nvshmem/src/host/transport/transport.cpp:nvshmemi_transport_init:282: init failed for transport: IBGDA                                                                                       
NH-DC-NM129-I06-12U-GPU-246:83321:83321 [1] NVSHMEM INFO Successfully initialized the transport: ibrc       

Next, the error you're getting RMA/AMO verb 4 not implemented 4 is coming from IBRC, because any operation that isn't an RDMA read (nvshmem_get/g), RDMA_write (nvshmem_put) or RDMA_SEND_INLINE (nvshmem_p) is not supported by that transport (for RMA), while only SIGNAL_SET and FETCH_ADD are implemented for AMOs. You can see ibrc.cpp around L1100-1200 for the implementation details. Interestingly, op 4 is NVSHMEMI_AMO_ADD so I'm not sure exactly what's going on there.接下来，您收到的错误 RMA/AMO verb 4 not implemented 4 来自 IBRC，因为该传输（对于 RMA）不支持任何非 RDMA 读取 (nvshmem_get/g)、RDMA 写入 (nvshmem_put) 或 RDMA_SEND_INLINE (nvshmem_p) 的操作，而 AMO 仅实现了 SIGNAL_SET 和 FETCH_ADD。您可以在 L1100-1200 左右查看 ibrc.cpp 的实现细节。有趣的是，操作 4 是 NVSHMEMI_AMO_ADD ，所以我不确定那里到底发生了什么。

This error message is inaccurate. The actual situation is that the ibv interface does not support 4-byte atomic operations.

[benhg]

I think we can change if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) to if (peer_base_addr != NULL) , because peer_base_addr != NULLimplies that peer_base_addr can be accessed directly in local memory rather than through ibgda.

We could do that, but that would be using NVLink/PCIe (since NVSHMEMI_JOB_GPU_LDST is allowed in a multi-gpu program with NVLink/PCIe support) , which the user would have explicitly disabled via NVSHMEM_DISABLE_P2P=1, so it's best not to.

Given your last response from this morning (at least in my time zone), I believe what we should do is add an additional check for the unsupported operations and give a more helpful error message.

[kwu130]

I think we can change if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) to if (peer_base_addr != NULL) , because peer_base_addr != NULLimplies that peer_base_addr can be accessed directly in local memory rather than through ibgda.

We could do that, but that would be using NVLink/PCIe (since NVSHMEMI_JOB_GPU_LDST is allowed in a multi-gpu program with NVLink/PCIe support) , which the user would have explicitly disabled via NVSHMEM_DISABLE_P2P=1, so it's best not to.

Given your last response from this morning (at least in my time zone), I believe what we should do is add an additional check for the unsupported operations and give a more helpful error message.

Assuming we enable ibgda and disable p2p, in nvshmemi_signal_op, consider two cases:

1. pe != mype, then peer_base_addr == NULL, so,if (peer_base_addr != NULL)and if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) are both false, at this time, nvshmemi_transfer_amo_nonfetch is called, the two if statements are equivalent.
2. pe == mype, then peer_base_addr != NULL, so, if (peer_base_addr != NULL) is true, and if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) is also false, at this time, the former calls atomicAdd_system, and the latter calls nvshmemi_transfer_amo_nonfetch, but in this case, it is sufficient to directly call atomicAdd_system, as it is a local memory operation.

In summary, I believe we can completely replace if (nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST) with if (peer_base_addr != NULL).

Looking forward to your reply again.

[benhg]

For case 1:
Yes, they end up being equivalent, but I think from a readability point of view it's easier to understand what's going on with the current implementation - JOB_GPU_LDST is more descriptive than the base_addr being NULL.

For case 2:
Similar reasoning as case 1, except that in this case we also prefer to have the wrapper functions to be more modular.

[Lyapunova-cd]

For case 1: Yes, they end up being equivalent, but I think from a readability point of view it's easier to understand what's going on with the current implementation - JOB_GPU_LDST is more descriptive than the base_addr being NULL.

For case 2: Similar reasoning as case 1, except that in this case we also prefer to have the wrapper functions to be more modular.

we need to check if the current op is SIGNAL_ADD extraly? job connectivity <= GPU_LDST dosen't indicate that we can use atomicAdd_system directly

[benhg]

@Lyapunova-cd yes, that's true too. GPU_LDST and GPU_LDST_ATOMICS are different. For example, PCI Express doesn't support all atomics.

[kwu130]

@benhg Hi,
I have another question. I noticed that in the nvshmem_defines header, for atomic operations such as nvshmem_##Name##_atomic_fetch_add, the implementation decides whether to use local atomics or remote atomics solely based on the condition:

nvshmemi_device_state_d.job_connectivity <= NVSHMEMI_JOB_GPU_LDST_ATOMICS

With this logic, if multiple transport mechanisms are initialized (e.g., both P2P and IBGDA), then job_connectivity will necessarily be greater than NVSHMEMI_JOB_GPU_LDST_ATOMICS. As a result, even when the target PE resides on the same node, the code still invokes the remote atomic path.

Wouldn’t this be suboptimal for performance? Would it be possible (or more efficient) to optimize this by checking whether the target PE is on the same node and whether native atomic operations are supported between the two PEs, and in that case directly using local atomics instead of the remote atomic path?

Looking forward to your reply. Thank you.