zeMemGetAddressRange error with alltoallv and zero-sized buffers

[rickybalin]

Summary

I am encountering the error below when running alltoallv on Aurora with oneCCL. This use of alltoallv intentionally sets only two non-zero buffers for each rank, whereas all other buffers are zero, thus performing a communication only with neighboring ranks. The error appears when setting CCL_ALLTOALLV=topo and when the non-zero buffers are greater than 8KB in size, but not with other alltoallv algorithms and smaller buffer sizes.

terminate called after throwing an instance of 'ccl::v1::exception'
  what():  oneCCL: ze_call.cpp:43 do_call: EXCEPTION: ze error at zeMemGetAddressRange, code: ZE_RESULT_ERROR_UNKNOWN

Version and environment

oneCCL version 2021.15
oneapi verison 2025.0.5
MPICH from pmodels/mpich@6037a7a
SUSE Linux Enterprise Server 15 SP4 v15.4
Aurora configuration for Intel Data Center Max 1550 GPU

Reproducer

alltoallv_ccl.cpp

#include <sycl/sycl.hpp>
#include <mpi.h>
#include <cmath>
#include <chrono>
#include "oneapi/ccl.hpp"
#include <algorithm>

// Get the nearest neighbors
std::vector<int> get_nearest_neighbors(int rank, int size)
{
    std::vector<int> neighbors;
    int num_neighbors = 1;

    if (size == 2) {
        int other_rank = (rank - 1 + size) % size;
        neighbors.push_back(other_rank);
    } else {
        for (int i = 0; i < num_neighbors; i++) {
            int left_rank = (rank - (1 + i) + size) % size;
            int right_rank = (rank + (1 + i)) % size;
            neighbors.push_back(left_rank);
            neighbors.push_back(right_rank);
        }
    }

    if (size < 100) {
        std::cout << "Rank " << rank << " neighbor list: ";
        for (int n : neighbors) {
            std::cout << n << " ";
        }
        std::cout << std::endl;
    }
    return neighbors;
}

bool contains(const std::vector<int>& vec, int value) {
    return std::find(vec.begin(), vec.end(), value) != vec.end();
}

int main(int argc, char** argv)
{
    int rank, size;
    double t1, t2, t3, t4, init_timer;

    ccl::init();

    std::chrono::time_point<std::chrono::high_resolution_clock> start = std::chrono::high_resolution_clock::now();
    MPI_Init(&argc, &argv);
    std::chrono::time_point<std::chrono::high_resolution_clock> end = std::chrono::high_resolution_clock::now();
    MPI_Barrier( MPI_COMM_WORLD );

    MPI_Comm_size(MPI_COMM_WORLD, &size);
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);

    /* create sycl queue */
    std::vector<sycl::device> gpu_devices;
    for (const auto& dev : sycl::device::get_devices()) {
        if (dev.is_gpu()) {
            gpu_devices.push_back(dev);
        }
    }
    if (gpu_devices.empty()) {
        std::cerr << "No GPU devices found!\n";
        MPI_Abort(MPI_COMM_WORLD, 1);
    }
    sycl::device selected_device = gpu_devices[rank % gpu_devices.size()];
    sycl::queue Q(selected_device);
    //sycl::queue Q(sycl::gpu_selector_v);

    /* create kvs */
    ccl::shared_ptr_class<ccl::kvs> kvs;
    ccl::kvs::address_type main_addr;
    if (rank == 0) {
        kvs = ccl::create_main_kvs();
        main_addr = kvs->get_address();
        MPI_Bcast((void *)main_addr.data(), main_addr.size(), MPI_BYTE, 0, MPI_COMM_WORLD);
    }
    else {
        MPI_Bcast((void *)main_addr.data(), main_addr.size(), MPI_BYTE, 0, MPI_COMM_WORLD);
        kvs = ccl::create_kvs(main_addr);
    }

    /* create communicator */
    auto dev = ccl::create_device(Q.get_device());
    auto ctx = ccl::create_context(Q.get_context());
    auto comm = ccl::create_communicator(size, rank, dev, ctx, kvs);

    /* create stream */
    auto stream = ccl::create_stream(Q);

    if (size < 100) {
        std::cout << "Rank " << rank << " running on " << Q.get_device().get_info<sycl::info::device::name>()  << std::endl;
    }
    int  elements_per_proc;
    int elements_per_proc_other;
    if (argc == 3)
    {
       elements_per_proc = atoi(argv[1])/4;
       elements_per_proc_other = atoi(argv[2])/4;
    }
    else
    {
        elements_per_proc = 1048576;
        elements_per_proc_other = 2;
    }

    // Get the neighboring ranks
    std::vector<int> neighbors;
    neighbors = get_nearest_neighbors(rank, size);
    MPI_Barrier( MPI_COMM_WORLD );

    // Initialize arrays
    std::vector<float> send_buff;
    std::vector<unsigned long> send_counts(size,0);
    std::vector<unsigned long> send_displs(size,0);
    std::vector<unsigned long> rcv_counts(size,0);
    std::vector<unsigned long> rcv_displs(size,0);

    // Fill in the send counts, displacements and buffers
    int global_send_elements = 0;
    for (int i=0; i<size; i++)
    {
        if (contains(neighbors,i)) {
            send_counts[i] = elements_per_proc;
            //global_send_elements += elements_per_proc;
            send_displs[i] = global_send_elements;
            global_send_elements += elements_per_proc;
            for (int n = 0; n < elements_per_proc; n++) {
                send_buff.push_back(rank);
            }
        } else {
            send_counts[i] = elements_per_proc_other;
            //global_send_elements += elements_per_proc;
            send_displs[i] = global_send_elements;
            global_send_elements += elements_per_proc_other;
            for (int n = 0; n < elements_per_proc_other; n++) {
                send_buff.push_back(-1.0);
            }
        }
    }
    if (size < 100) {
        std::cout << "Rank " << rank << " sending " << global_send_elements << " elements" << std::endl;
        fflush(stdout);
    }
    MPI_Barrier( MPI_COMM_WORLD );

    // Get the received data
    int global_rcv_elements = 0;
    MPI_Alltoall(send_counts.data(), 1, MPI_UNSIGNED_LONG,
                 rcv_counts.data(), 1, MPI_UNSIGNED_LONG,
                 MPI_COMM_WORLD);
    for (int i = 0; i < size; i++) {
        if (rcv_counts[i] != 0 and size < 100) {
                std::cout << "Rank " << rank << " receives " << rcv_counts[i] <<
                          " elements from rank " << i << std::endl;
        }
        //global_rcv_elements += rcv_counts[i];
        rcv_displs[i] = global_rcv_elements;
        global_rcv_elements += rcv_counts[i];
    }
    std::vector<float> rcv_buff(global_rcv_elements, -99.0);

    // Move the send and receive buffers to the GPU
    float *dsend_buff  = sycl::malloc_device<float>(global_send_elements,Q);
    float *drcv_buff  = sycl::malloc_device<float>(global_rcv_elements,Q);
    Q.memcpy((void *) dsend_buff, (void *) send_buff.data(), global_send_elements*sizeof(float));
    Q.memcpy((void *) drcv_buff, (void *) rcv_buff.data(), global_rcv_elements*sizeof(float));
    Q.wait();
    MPI_Barrier( MPI_COMM_WORLD );

    int iters = 10;
    std::vector<double> elapsed(iters);
    for (int i = 0; i < iters; i++)
    {
        t3 = MPI_Wtime();
        //MPI_Alltoallv(dsend_buff, send_counts.data(), send_displs.data(), MPI_FLOAT,
        //             drcv_buff, rcv_counts.data(), rcv_displs.data(), MPI_FLOAT,
        //             MPI_COMM_WORLD);
        ccl::alltoallv(dsend_buff, send_counts,
                       drcv_buff, rcv_counts,
                       comm, stream).wait();
        MPI_Barrier( MPI_COMM_WORLD );
        t4 = MPI_Wtime();
        if ( rank == 0 )    elapsed[i]=( t4 - t3 ) * 1e3;
    }

    double avg = 0.0;
    int skip = 4;
    if ( rank == 0 )
    {
        for (int i = skip; i < iters; i++)
        {
            avg = avg + elapsed[i];
            //std::cout<<elapsed[i]<<std::endl;
        }
        avg = avg / (iters - skip);
        std::cout << "Average all2all time: " << avg << " ms" << std::endl;
    }

    /*
    Q.memcpy(rcv_buff.data(), drcv_buff, global_rcv_elements*sizeof(float)).wait();
    if (rank == 0) {
        std::cout << "Rank 0 received: " << std::endl;
        for (int i=0; i<rcv_buff.size(); i++) {
            std::cout << i << " " << rcv_buff[i] << std::endl;
        }
    }
    */

    MPI_Finalize();
    return 0;
}

On Aurora, build with

module load frameworks
mpicxx -o all2allv_ccl all2allv_ccl.cpp -fsycl -lmpi \
    -I/opt/aurora/24.347.0/oneapi/ccl/2021.14/include \
    -L/opt/aurora/24.347.0/oneapi/ccl/2021.14/lib -lccl

and run with

#!/bin/bash

module load frameworks

# Use the latest oneCCL
unset CCL_ROOT
export CCL_CONFIGURATION_PATH=""
export CCL_CONFIGURATION=cpu_gpu_dpcpp
export CCL_ROOT="/lus/flare/projects/Aurora_deployment/datascience/software/ccl_2021.15/oneCCL/build_2021p15/"
export LD_LIBRARY_PATH=${CCL_ROOT}/lib:$LD_LIBRARY_PATH
export CPATH=${CCL_ROOT}/include:$CPATH
export LIBRARY_PATH=${CCL_ROOT}/lib:$LIBRARY_PATH

export CCL_PROCESS_LAUNCHER=pmix
export CCL_ATL_TRANSPORT=mpi
export CCL_KVS_MODE=mpi
export CCL_ENABLE_SYCL_KERNELS=1
export CCL_ALLTOALLV=topo
#export CCL_LOG_LEVEL=debug

EXE=/flare/datascience/balin/oneCCL/dl_scaling/aurora_frameworks_scaling/c++-sycl-oneccl-all2allv/all2allv_ccl
NNODES=`wc -l < $PBS_NODEFILE`
RANKS_PER_NODE=12
NRANKS=$(( NNODES * RANKS_PER_NODE ))
CPU_BINDING=list:1-4:8-11:16-19:24-27:32-35:40-43:53-56:60-63:68-71:76-79:84-87:92-95
EXT_ENV="--env FI_CXI_DEFAULT_CQ_SIZE=1048576 --env CCL_ALLTOALLV_MONOLITHIC_KERNEL=0"
OTHER_BUF_SIZE=0

echo Using MPI from:
ldd $EXE | grep libmpi
echo
echo Using oneCCL from:
ldd $EXE | grep ccl
echo


BUF_SIZE=16384
mpiexec ${EXT_ENV} --np ${NRANKS} -ppn ${RANKS_PER_NODE}  --cpu-bind  $CPU_BINDING \
        $EXE ${BUF_SIZE} ${OTHER_BUF_SIZE}

Logs

Can add if requested, very long log file.

Expected behavior

The expected behavior is for the error to not appear.

Observed behavior

The observed behavior is the error reposted above.

Existing workarounds

Workarounds are to change the alltoallv algorithm to the other options, but with a drop in performance.

Affected projects

This bug impacts scalability of ML applications using alltoallv on Aurora.

[no-ponomarev]

Thanks for posting an issue. zeMemGetAddressRange is a part of get_mem_info API which is called during IPC handlers creation in ze_handle_exchange_entry::create_local_ipc_handles. Looks like collective didn't even started. Hard to tell what is the problem, I have to reproduce it first and then can post an update.

[colleeneb]

On Aurora, I did a little bit of looking into this with a LeveL Zero tracer, and it looks like basically for this case some code in oneCCL is trying to touch memory on the GPU outside of what it allocated.

From the test code, in lines

    float *dsend_buff  = sycl::malloc_device<float>(global_send_elements,Q);
    float *drcv_buff  = sycl::malloc_device<float>(global_rcv_elements,Q);

for each rank 65536 bytes are allocated (sizeof(float)*8192) *2 (for send and receive). The 8192 is since, as printed to the screen, each rank sends/receives 4096 floats to two other ranks.

In the level zero trace of this run we see indeed exactly 65536 bytes are allocated:

lttng_ust_ze:zeMemAllocDevice_entry: { hContext: 0x0000000006be2160, device_desc: 0x00007ffdb07673d8, size: **65536**, alignment: 0, hDevice: 0x00000000058bd6a8, pptr: 0x00007ffdb07674f8, device_desc_val: { stype: ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC, pNext: 0x0000000000000000, flags: [  ], ordinal: 0 } }
lttng_ust_ze:zeMemAllocDevice_exit: { zeResult: ZE_RESULT_SUCCESS, pptr_val: 0xff03fffffffe0000 }

And we see that it’s at address 0xff03fffffffe0000, allocated for 65536 bytes, so from 0xff03fffffffe0000 to 0xff03fffffffeffff.

However, inside the alltoallv call, the trace shows:

lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe0000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetIpcHandle_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe0000, pIpcHandle: 0x00007ffdb0762428 }
lttng_ust_ze:zeMemGetIpcHandle_exit: { zeResult: ZE_RESULT_SUCCESS, pIpcHandle_val: { data: "\xde\x01\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\
\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00\x00" } }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe4000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe4000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe8000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffe8000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffec000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: 0xff03fffffffec000, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_SUCCESS, pBase_val: 0xff03fffffffe0000, pSize_val: 65536 }
lttng_ust_ze:zeMemGetAddressRange_entry: { hContext: 0x0000000006be2160, ptr: **0xff03ffffffff0000**, pBase: 0x00007ffdb07626e8, pSize: 0x00007ffdb07626e0 }
lttng_ust_ze:zeMemGetAddressRange_exit: { zeResult: ZE_RESULT_ERROR_UNKNOWN, pBase_val: 0x0000000000000000, pSize_val: 0 }

So it’s calling zeMemGetAddressRange 8 times (I launched with 4 MPI ranks, so 4 for recv and 4 for send), and the last time it calls it with an address, 0xff03ffffffff0000, that is 1 byte out of range of what is allocated. Hopefully this is useful for the debug!

[colleeneb]

A little more details. It looks like possibly it needs a guard on the pointer sent to zeMemGetAddressRange to sending a case where the pointer is outside of the allocation since a bock is 0.

To demonstrate, I printed out the value for mem_ptr that are sent to zeMemGetAddressRange in a loop here:

oneCCL/src/sched/entry/ze/ze_handle_exchange_entry.cpp

Line 118 in 303b41a

auto mem_ptr = in_buffers[buf_idx].first;

If we just look at rank 3 for a 4-rank run of the test from Riccardo:

What we expect for rank 3 is:

• Rank 3 sends 4096 floats to rank 0 and rank 2, and nothing to rank 1.
• Rank 3 receives 4096 floats from rank 0 and rank 2, and nothing from rank 1.

What we see in the output, noting that the GPU memory allocation starts at 0xff03fffffffe0000, for 65536 (0x10000) bytes (0xff03fffffffe0000 to 0xff03fffffffeffff):

 # memory location for send buffer, block 0 for sending to rank 0, it looks as expected, going to pt 0xff03fffffffe0000 at offset 0
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffe0000, idx: 0

 # memory location for send buffer, block 1 for sending to rank 1, it looks as expected, going to pt 0xff03fffffffe4000, which is 0x4000 (sizeof(float)*4096) bytes after 0xff03fffffffe0000 since we send that amount to rank 0.
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffe4000, idx: 1

 # memory location for send buffer, block 2 for sending to rank 2, it looks as expected, going to pt 0xff03fffffffe4000 again, since it is not sending to rank 1, so no additional offset.
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffe4000, idx: 2

 # memory location for send buffer, block 3 for sending to rank 3, it looks as expected, going to pt 0xff03fffffffe8000, which is 0x4000 (sizeof(float)*4096) bytes after the buffer for rank 2, since we send that amount to rank 2.
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffe8000, idx: 3

 # memory location for recv buffer, block 0 for receiving from rank 0, it looks as expected, going to pt 0xff03fffffffe8000 again, since rank 3 doesn’t receive anything from itself
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffe8000, idx: 4

 # memory location for recv buffer, block 1 for receiving from rank 1, it looks as expected, going to pt 0xff03fffffffec000, 0x4000 past the last buffer to give space for the buffer received from rank 0
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffec000, idx: 5

 # memory location for recv buffer, block 2 for receiving from rank 2, it looks as expected, going to pt 0xff03fffffffec000 again since doesn’t recive anything from rank 1
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03fffffffec000, idx: 6

 # ***memory location for recv buffer, block 3 for receiving from rank 3, it uses 0xff03fffffffec000 + 0x4000 since 0x4000 is the length of how much it should receive from rank 2, but since it receives nothing from rank 3, this memory address is 1 byte past the allocation space, and so it will crash***
2025:08:01-13:47:58:(175275) |CCL_DEBUG| ze_handle_exchange_entry.cpp:119 create_local_ipc_handles: mem_ptr: 0xff03ffffffff0000, idx: 7

So the last iteration, for rank 3's buffer for receiving from rank 3, it offsets by 0xff03fffffffec000 by 0x4000 since 0x4000 is the length of how much it should receive from rank 2, but since it receives nothing from rank 3, this memory address is 1 byte past the allocation space, and so it will crash. Possibly there needs to be a guard on the mem_ptr to avoid a case where the size of the bock to send/rec is 0 and the offset == total size or something like that.

For the question of “why does it only fail when >= 16384 bytes?”: It looks like by default LZ under the hood always allocates in 65536 byte chunks, even if the malloc_device is for less (at least this is what I see when tracing this test). so even with smaller sizes it's likely technically checking outside the amount allocated, but since LZ allocates more under the hood, it’s not until the test code uses a buffer size of 16384 that we are just at the 65536 size, and we see it.

Making the following change:

diff --git a/src/coll/algorithms/alltoallv.cpp b/src/coll/algorithms/alltoallv.cpp
index ebc247f..61ff993 100644
--- a/src/coll/algorithms/alltoallv.cpp
+++ b/src/coll/algorithms/alltoallv.cpp
@@ -416,9 +416,16 @@ ccl::status ccl_coll_build_topo_alltoallv(ccl_sched* main_sched,
         for (int idx = 0; idx < comm_size; idx++) {
             send_bufs[idx].set(coll_param.get_send_buf(), total_send_bytes, send_offsets[idx]);
             recv_bufs[idx].set(coll_param.get_recv_buf(), total_recv_bytes, recv_offsets[idx]);
+           if(recv_offsets[idx] == total_recv_bytes) {
+             recv_bufs[idx].set(nullptr, 0, 0);
+             LOG_DEBUG("   PROBLEM? ");
+           }
+           if(send_offsets[idx] == total_send_bytes) {
+             send_bufs[idx].set(nullptr, 0, 0);
+             LOG_DEBUG("   PROBLEM? ");
+           }
         }
     }

avoids the crash, and avoids calling zeMemGetAddressRange on the pointer that isn’t allocated (since if the offset is the same size as the number of total bytes, we know it's past the end of the allocation) but I’m not 100% sure it will be correct throughout the code.

[srirajpaul]

@colleeneb Thanks for the analysis. it is spot on. Will make the changes.

[jbrosenz]

Fix was merged and will be part of the oneCCL 2021.17 release / oneAPI 2025.3 release.

[colleeneb]

Amazing! Thanks so much!