[Issue]: Failed to run reduction_latency across nodes using ibrc

[kwu130]

How is this issue impacting you?

Application hang

Share Your Debug Logs

Operation log as follows:
log.txt

Steps to Reproduce the Issue

Minimal Steps:

cd nvshmem-3.4.5-0
cmake -DCUDA_ARCHITECTURES="75" -B build -S .
cd build
make -j8
cd build-install/bin/perftest/device/coll
nvshmrun -n 4 -hosts node1,node2 ./reduction_latency -w 0 -n 1 -c 1 -t 32 -b 8 -e 16

Environment Details:

NVSHMEM_IBGDA_SUPPORT=1
NVSHMEM_DISABLE_MNNVL=true
NVSHMEM_BUILD_EXAMPLES=1
NVSHMEM_MPI_SUPPORT=0
NVSHMEM_IBGDA_NIC_HANDLER=gpu
NVSHMEM_IB_DISABLE_DMABUF=1
NVSHMEM_USE_GDRCOPY=0
NVSHMEM_NVTX=0
NVSHMEM_DISABLE_P2P=1
NVSHMEM_BUILD_TESTS=1
NVSHMEM_IB_ENABLE_IBGDA=0
NVSHMEM_DEBUG_SUBSYS=ALL
NVSHMEM_DEVICELIB_CUDA_HOME=/usr/local/cuda
NVSHMEM_PREFIX=/opt/nvshmem-3.4.5-0/build-install
NVSHMEM_IBRC_SUPPORT=1
NVSHMEM_DISABLE_CUDA_VMM=true
NVSHMEM_DEBUG=INFO

Intermittency: everytime
Previous Success: no, also failed in nvshmem-3.2.5

NVSHMEM Version

nvshmem-3.4.5-0 + cuda12.6

Your platform details

Error Message & Behavior

[nvshmem-3.4.5-0/perftest/device/coll/reduction_latency.cu:253] cuda failed with an illegal memory access was encountered
[nvshmem-3.4.5-0/perftest/device/coll/reduction_latency.cu:253] cuda failed with an illegal memory access was encountered

[kwu130]

Even when running this test case with ibgda, the same failure occurs.

[benhg]

Is there a reason you are building for SM75 (Turing) on an Ada system? Can you try again with CUDA_ARCHITECTURES=sm89? And could you run again with CUDA_LOG_FILE=stdout? That should dump whatever errors inside CUDA are causing the invalid access to get reported.

[kwu130]

CUDA_LOG_FILE=stdout

Because I only have two nodes for testing, one is l20, and the other is a800. They both support the SM80 architecture at most, and even if I rebuild nvshmem with CUDA_ARCHITECTURES=80, the same error occurs.
Even with CUDA_LOG_FILE=stdout set, there is no more information (the log file is as follows).
If I use one node, such as l20, to simulate cross-node communication (by setting NVSHMEM_DISABLE_P2P=1), the error will not occur.
Could you please try to reproduce this issue on your own environment?

log with sm80.txt

[benhg]

I'm not sure I can reproduce your environment exactly in-house.

If I use one node, such as l20, to simulate cross-node communication (by setting NVSHMEM_DISABLE_P2P=1), the error will not occur.

Hearing this makes me think there's a misconfiguration of something in your network. Have you verified via something like RDMA-Core Perftests that your IB/RoCE network is working the way you want it to?

[kwu130]

Hearing this makes me think there's a misconfiguration of something in your network. Have you verified via something like RDMA-Core Perftests that your IB/RoCE network is working the way you want it to?

My network environment should be fine, other test cases in nvshmem are also running normally, and the GPU's memory write bandwidth tested with ib_write_bw is also normal.

@benhg I continued to analyze it and finally found that the invalid address was accessed in nvshmemi_packLL. The output is as follows. I hope this helps you resolve the bug.

template <typename T, threadgroup_t SCOPE, int UNROLL>
__device__ NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmemi_packLL(T *psync, const T *source,
                                                              size_t nelems, uint32_t ll_flag,
                                                              nvshmemi_team_t *teami, int pe_count,
                                                              int team_offset) {
    const size_t num_subelems = (nelems * sizeof(T)) / sizeof(uint32_t);
    const int myIdx = nvshmemi_thread_id_in_threadgroup<SCOPE>();
    const int groupSize = nvshmemi_threadgroup_size<SCOPE>();
    /* each thread will write 2 subelements (8 bytes) */
    const size_t element_start = myIdx * UNROLL * 2;
    const size_t element_stride = groupSize * UNROLL * 2;
    size_t element_offset;
    uint32_t regs[2 * UNROLL];

    int current_global_pe_index;
    uint32_t *current_dest_address;

    /* We need to be sure to fill each unrolled loop */
    assert(num_subelems % UNROLL * 2 == 0);
    assert((uintptr_t)psync % 16 == 0);

    for (element_offset = element_start; element_offset < num_subelems;
         element_offset += element_stride) {
#pragma unroll
        for (int unroll = 0; unroll < UNROLL * 2; unroll += 2) {
            if (sizeof(T) == sizeof(uint64_t)) {
                asm("ld.global.v2.u32 {%0,%1}, [%2];"
                    : "=r"(regs[unroll]), "=r"(regs[unroll + 1])
                    : "l"((uint32_t *)source + element_offset + unroll));
            } else {
                regs[unroll] = *((uint32_t *)source + element_offset + unroll);
                regs[unroll + 1] = *((uint32_t *)source + element_offset + unroll + 1);
            }
        }
        for (int pe = 0; pe < pe_count; pe++) {
            current_global_pe_index =
                nvshmemi_team_translate_pe_to_team_world_wrap(teami, team_offset + pe);
            current_dest_address =
                (uint32_t *)nvshmemi_ptr(psync, current_global_pe_index) + element_offset * 2;

#pragma unroll
            for (int unroll = 0; unroll < UNROLL * 2; unroll += 2) {
                printf("<%s:%d> PE[%d], current_dest_address: %p\n",
                        __FUNCTION__, __LINE__, nvshmemi_device_state_d.mype, current_dest_address);
                asm volatile("st.volatile.global.v4.u32 [%0], {%1,%2,%3,%4};" ::"l"(
                                 (int4 *)current_dest_address),
                             "r"(regs[unroll]), "r"(ll_flag), "r"(regs[unroll + 1]), "r"(ll_flag));
                current_dest_address += sizeof(int4) / sizeof(uint32_t);
            }
        }
    }
}

[benhg]

It's possible that your system doesn't have proper support for the LL/LL128 protocol. That may show up this way. Could you try forcing the LL protocol off ? By the way, I think L20 (which is a PCIe GPU) probably does not support LL protocol.

[kwu130]

It's possible that your system doesn't have proper support for the LL/LL128 protocol. That may show up this way. Could you try forcing the LL protocol off ? By the way, I think L20 (which is a PCIe GPU) probably does not support LL protocol.

@benhg Thank you for your reply. Following your suggestion, I re-ran the experiment on two H20 nodes that fully support the LL/LL128 protocol. Unfortunately, the results remained unchanged. I also tried setting the environment variable NVSHMEM_FCOLLECT_ALGO=2 (corresponding to FCOLLECT_ONESHOT) to bypass the LL protocol execution path, so that the program executes nvshmemi_fcollect_allpush_threadgroup instead of nvshmemi_fcollect_allpush_ll_threadgroup. However, under this configuration, the program hung during execution.

As a side note, I obtained the same results even when building the code with CUDA_ARCHITECTURES=sm89 on H20.

Additionally, it is worth noting that the problem only occurs when the SCOPE is set to block; both thread and warp scopes function as expected.

I would like to know whether, in your environment, executing nvshmrun -n 4 -hosts node1,node2 ./reduction_latency completes successfully without triggering the error “cuda failed with an illegal memory access was encountered”.