Interactions with cray-mpich, GTL and memory buffer registering

[etiennemlb]

Using Cray MPICH and Cray NICs' RDMA capabilities implies registering a buffer. On the host its done either via hooks to malloc or pagefaults monitors. On GPU buffers IDK how it is handled, probably similarly. The fact that call cudaMalloc instead of hipmalloc may confuse the underlying Cray logic.

// $ module purge
// $ module load cpe/24.07
// $ module load craype-accel-amd-gfx90a craype-x86-trento
// $ module load PrgEnv-gnu
// $ module load gcc/13.2.0
// $ module load amd-mixed/6.4.3
// $ module load cmake/3.27.9
// $ module list
// $ source ...../scale-1.7.2-Linux/bin/scaleenv gfx90a
// $ nvcc -lineinfo -g -std=c++17 --offload-arch=gfx90a -O3 -I /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/include -o test_scale_mpi test_scale_mpi.cu -L/opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib -lmpi --gcc-install-dir=/opt/software/gcc/13.2.0/lib/gcc/x86_64-pc-linux-gnu/13.2.0 -Wl,-rpath,/opt/software/gcc/13.2.0/lib64 -L/opt/cray/pe/mpich/8.1.30/gtl/lib -lmpi_gtl_hsa
// $ MPICH_GPU_SUPPORT_ENABLED=1 srun --ntasks=2 test_scale_mpi
// process_vm_readv: Bad address
// Assertion failed in file ../src/mpid/ch4/shm/cray_common/cray_common_memops.c at line 461: 0
// process_vm_readv: Bad address
// Assertion failed in file ../src/mpid/ch4/shm/cray_common/cray_common_memops.c at line 461: 0
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(MPL_backtrace_show+0x26) [0x14ca0b3b700b]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x1c4a214) [0x14ca0acf3214]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x214449a) [0x14ca0b1ed49a]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x216066a) [0x14ca0b20966a]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x213cc67) [0x14ca0b1e5c67]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0xd99548) [0x14ca09e42548]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0xde4411) [0x14ca09e8d411]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(MPI_Waitall+0x3cc) [0x14ca09e8dc1c]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(MPL_backtrace_show+0x26) [0x1519dea4700b]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x1c4a214) [0x1519de383214]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x214449a) [0x1519de87d49a]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x216066a) [0x1519de89966a]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0x213cc67) [0x1519de875c67]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0xd99548) [0x1519dd4d2548]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(+0xde4411) [0x1519dd51d411]
// /opt/cray/pe/mpich/8.1.30/ofi/gnu/11.2/lib/libmpi_gnu_112.so.12(MPI_Waitall+0x3cc) [0x1519dd51dc1c]
// test_scale_mpi(+0x142e) [0x55bd5df4142e]
// /lib64/libc.so.6(+0x295d0) [0x1519dc0035d0]
// /lib64/libc.so.6(__libc_start_main+0x80) [0x1519dc003680]
// test_scale_mpi(+0x11a5) [0x55bd5df411a5]
// test_scale_mpi(+0x142e) [0x55c027a8442e]
// /lib64/libc.so.6(+0x295d0) [0x14ca089735d0]
// /lib64/libc.so.6(__libc_start_main+0x80) [0x14ca08973680]
// test_scale_mpi(+0x11a5) [0x55c027a841a5]
// MPICH ERROR [Rank 0] [job id 5249618.34] [Wed Jul 29 11:15:26 2026] [g1032] - Abort(1): Internal error
// MPICH ERROR [Rank 1] [job id 5249618.34] [Wed Jul 29 11:15:26 2026] [g1032] - Abort(1): Internal error

#include <cuda_runtime.h>
#include <mpi.h>

#include <cstdio>
#include <cstdlib>

#define CHECK_MPI(cmd)                      \
    do {                                    \
        int e = (cmd);                      \
        if(e != MPI_SUCCESS) {              \
            fprintf(stderr, "MPI error\n"); \
            MPI_Abort(MPI_COMM_WORLD, e);   \
        }                                   \
    } while(0)
#define CHECK_CUDA(cmd)                         \
    do {                                        \
        cudaError_t e = (cmd);                  \
        if(e != cudaSuccess) {                  \
            fprintf(stderr, "CUDA error: %s\n", \
                    cudaGetErrorString(e));     \
            exit(EXIT_FAILURE);                 \
        }                                       \
    } while(0)
__global__ void init_kernel(float* buf, float value, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if(i < n)
        buf[i] = value;
}
int main(int argc, char** argv) {
    CHECK_MPI(MPI_Init(&argc, &argv));
    int rank, size;
    CHECK_MPI(MPI_Comm_rank(MPI_COMM_WORLD, &rank));
    CHECK_MPI(MPI_Comm_size(MPI_COMM_WORLD, &size));
    const int N = 1 << 20;
    float*    d_send;
    float*    d_recv;
    CHECK_CUDA(cudaMalloc(&d_send, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_recv, N * sizeof(float)));
    init_kernel<<<(N + 255) / 256, 256>>>(d_send, (float)rank, N);
    CHECK_CUDA(cudaGetLastError());
    CHECK_CUDA(cudaDeviceSynchronize());
    int         dst = (rank + 1) % size;
    int         src = (rank - 1 + size) % size;
    MPI_Request reqs[2];
    CHECK_MPI(MPI_Irecv(d_recv, N, MPI_FLOAT, src, 123, MPI_COMM_WORLD, &reqs[0]));
    CHECK_MPI(MPI_Isend(d_send, N, MPI_FLOAT, dst, 123, MPI_COMM_WORLD, &reqs[1]));
    CHECK_MPI(MPI_Waitall(2, reqs, MPI_STATUSES_IGNORE));
    CHECK_CUDA(cudaDeviceSynchronize());
    CHECK_CUDA(cudaFree(d_send));
    CHECK_CUDA(cudaFree(d_recv));
    CHECK_MPI(MPI_Finalize());
    return 0;
}

[etiennemlb]

I guess this makes SCALE ... not scale ;)., Nah seriously, while libfabric or cray's gpu translat layer can do shady things, this is common, what is the plan to support these cray mpich ? Large AMD supercomputers/ai factory where scale is being considered are based on Cray or Bull so we are going to see the issue.

[etiennemlb]

• When I launch the program fora single rank I get a segfault, mpich uses hsa_amd_pointer_info to get a hsa hsa_amd_pointer_info_s, the function returns success with a ptr type of 0 (HSA_EXT_POINTER_TYPE_UNKNOWN). cray mpich treats it as host mem and does a memcpy which segfaults.

[ChrisKitching]

Thanks for the report! Very much agree that this stuff is extremely important to get working.

This is not the first "Fun adventure with Cray" we have had, but it is the first we've had with Cray MPI! :D

what is the plan to support these cray mpich ?

In general, the answer is the same as for any other CUDA library. Cray mpich isn't magic, it's "just" a library that uses the CUDA Runtime/Driver API to do stuff, so rebuilding it against the SCALE runtime should produce the desired effect. There's examples in this repo of doing that with openmpi. Like with any CUDA library, this boils down to "rebuild it like normal, but with SCALE instead of NVIDIA CUDA".

The mental model to use is to think of CUDA as a programming language (like C or Rust), not an NVIDIA-specific platform. There isn't any conceptual difference between the CUDA API for doing IPC buffer transfer or the HIP/hsa one (they even have almost exactly the same names), so why does MPI need a special codepath for AMD mode? Recompiling the CUDA codepaths with SCALE should produce equivalent results, and if it doesn't then we've implemented some of the CUDA runtime APIS wrong.

For closed-source programs with no device code that use the CUDA C APIs and have dynamically linked it, having them load our CUDA libraries instead of NVIDIA's should work. We have rather strong ABI compatibility with NVIDIA, given how many programs access the CUDA APIs via dlopen - sometimes via Python.

A common mistake is to believe that the rocm/hip backend of MPI should be used. That won't work: SCALE is not built on the HIP/rocm runtime, but does currently contain a few (violently mangled) bits ripped from rocm that might make it look like we are to someone who goes hunting through our files. With a few exceptions, SCALE is complete reimplementation of everything from the kernel driver up, so interoperability with the rocm/hip runtime is not expected (and not possible, since the HIP APIs do different things from the corresponding CUDA ones).

Interoperability/ABI compatibility with HIP is not possible, since the goals of the two projects are different: SCALE is aiming to provide a perfectly CUDA-equivalent experience, including all the strange edgecase behavoiurs that matter to real code. HIP is aiming to provide a GPU programming experience on AMD, intentionally deviating from NVIDIA CUDA behaviour in many cases. Clearly just wiring up corresponding functions to each would end badly.

When I launch the program fora single rank I get a segfault, mpich uses hsa_amd_pointer_info to get a hsa hsa_amd_pointer_info_s, the function returns success with a ptr type of 0 (HSA_EXT_POINTER_TYPE_UNKNOWN). cray mpich treats it as host mem and does a memcpy which segfaults.

Calling directly into hsa/rocm while working with SCALE is definitely not going to work.

SCALE uses the following components of rocm in the following ways:

• hsakmt (a thin C wrapper around the kernel module) is the primary way we interact with the GPU.
• hsa (a higher-level, quite awful library on top of hsakmt) is used for some initialisation stuff (where it has some useful firmware bug mitigations), but is out of the loop for most other things. This is to a significant extent technical debt: we built SCALE on top of hsa in the very early days, and have been steadily fleeing it ever since (mostly for performance reasons).
• We wrap recompiled/patched versions of the various rocm math libraries (rocblas, rocfft, etc.) to provide the NVIDIA equivalents. We're in the process of moving away from this in favour of recompiling these libraries using our compiler/runtime, since the existing wrapper approach requires the SCALE and HIP runtimes to coexist, despite being mutually incompatible. This causes a variety of terrible problems, though is likely not the cause of your current problem.

Notably, we do not use the HSA memory allocator: we've had to write our own memory allocator to work around a driver/firmware bug in which the GPU could end up with page tables that are out-of-sync with the kernel (causing the GPU and CPU to disagree about whether or not a buffer is allocated). This also ends up with us being able to do powerful optimisations as a result of choosing a more convenient address mapping scheme. hsa_amd_pointer_info will give the wrong answer when given a pointer allocated by SCALE, since it will be expecting the hsa/HIP address mapping scheme and/or inspecting hsa's own memory allocator's datastructures, which we don't use. This isn't really a concern, since interopability with hsa is not expected to work.

---

I hope this rather long post is informative. Concrete next steps could include:

• Try recompiling the cray mpi, if the source is available.
• Try using the CUDA version of Cray MPI, but make it load our CUDA .so files instead of NVIDIA's.
• Try OpenMPI

... I'm going to go add a documentation page about all this now, since I think you're the third person to ask approximately this in the past few days days ;).

[etiennemlb]

OK this great,ill try to see if linking the cuda gtl works.

…

On 31 July 2026 21:29:23 UTC, Chris Kitching ***@***.***> wrote: ChrisKitching left a comment (spectral-compute/scale-validation#66) Thanks for the report! Very much agree that this stuff is extremely important to get working. This is not the first "Fun adventure with Cray" we have had, but it is the first we've had with Cray MPI! :D > what is the plan to support these cray mpich ? In general, the answer is the same as for any other CUDA library. Cray mpich isn't magic, it's "just" a library that uses the CUDA Runtime/Driver API to do stuff, so rebuilding it against the SCALE runtime should produce the desired effect. There's examples in this repo of doing that with openmpi. Like with any CUDA library, this boils down to "rebuild it like normal, but with SCALE instead of NVIDIA CUDA". The mental model to use is to think of CUDA as a programming language (like C or Rust), not an NVIDIA-specific platform. There isn't any conceptual difference between the CUDA API for doing IPC buffer transfer or the HIP/hsa one (they even have almost exactly the same names), so why does MPI need a special codepath for AMD mode? Recompiling the CUDA codepaths with SCALE should produce equivalent results, and if it doesn't then we've implemented some of the CUDA runtime APIS wrong. For closed-source programs with no device code that use the CUDA C APIs and have dynamically linked it, having them load our CUDA libraries instead of NVIDIA's should work. We have rather strong ABI compatibility with NVIDIA, given how many programs access the CUDA APIs via dlopen - sometimes via Python. A common mistake is to believe that the rocm/hip backend of MPI should be used. That won't work: SCALE is not built on the HIP/rocm runtime, but does currently contain a few (violently mangled) bits ripped from rocm that might make it look like we are to someone who goes hunting through our files. With a few exceptions, SCALE is complete reimplementation of everything from the kernel driver up, so interoperability with the rocm/hip runtime is not expected (and not _possible_, since the HIP APIs do different things from the corresponding CUDA ones). Interoperability/ABI compatibility with HIP is not possible, since the goals of the two projects are different: SCALE is aiming to provide a perfectly CUDA-equivalent experience, including all the strange edgecase behavoiurs that matter to real code. HIP is aiming to provide a GPU programming experience on AMD, intentionally deviating from NVIDIA CUDA behaviour in many cases. Clearly just wiring up corresponding functions to each would end badly. > When I launch the program fora single rank I get a segfault, mpich uses hsa_amd_pointer_info to get a hsa hsa_amd_pointer_info_s, the function returns success with a ptr type of 0 (HSA_EXT_POINTER_TYPE_UNKNOWN). cray mpich treats it as host mem and does a memcpy which segfaults. Calling directly into hsa/rocm while working with SCALE is definitely not going to work. SCALE uses the following components of rocm in the following ways: - hsakmt (a thin C wrapper around the kernel module) is the primary way we interact with the GPU. - hsa (a higher-level, quite awful library on top of hsakmt) is used for some initialisation stuff (where it has some useful firmware bug mitigations), but is out of the loop for most other things. This is to a significant extent technical debt: we built SCALE on top of hsa in the very early days, and have been steadily fleeing it ever since (mostly for performance reasons). - We wrap recompiled/patched versions of the various rocm math libraries (rocblas, rocfft, etc.) to provide the NVIDIA equivalents. We're in the process of moving away from this in favour of recompiling these libraries using our compiler/runtime, since the existing wrapper approach requires the SCALE and HIP runtimes to coexist, despite being mutually incompatible. This causes a variety of terrible problems, though is likely not the cause of your current problem. Notably, we do not use the HSA memory allocator: we've had to write our own memory allocator to work around a driver/firmware bug in which the GPU could end up with page tables that are out-of-sync with the kernel (causing the GPU and CPU to disagree about whether or not a buffer is allocated). This also ends up with us being able to do powerful optimisations as a result of choosing a more convenient address mapping scheme. `hsa_amd_pointer_info` will give the wrong answer when given a pointer allocated by SCALE, since it will be expecting the hsa/HIP address mapping scheme and/or inspecting hsa's own memory allocator's datastructures, which we don't use. This isn't really a concern, since interopability with hsa is not expected to work. I hope this rather long post is informative. Concrete next steps could include: - Try recompiling the cray mpi, if the source is available. - Try using the CUDA version of Cray MPI, but make it load our CUDA `.so` files instead of NVIDIA's. - Try OpenMPI ... I'm going to go add a documentation page about all this now, since I think you're the third person to ask approximately this in the past few days days ;).

[etiennemlb]

Hey,
so we dont have access to the code, the cray-mpich and its addons (gtl)) for GPUs are closed source. cray-mpich works better than openmpi on many cray machines (for now at least).
Looking at the symbols and gdb, the gtl is also used for mpi reductions and in that case cuda launches kernels, so unless scales compiles ptx to cdnaX it wont work. A rebuild from Cray/HPE is needed.

Anyway, trying anyway to see what happens:

/usr/bin/ld: warning: libcudart.so.12, needed by /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so, not found (try using -rpath or -rpath-link)
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaUnregisterFatBinary@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaDeviceGetPCIBusId@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cuCtxEnablePeerAccess'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaFree@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaInitModule@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaRegisterFatBinaryEnd@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaStreamCreateWithFlags@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cuStreamWriteValue64_v2'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cuStreamWaitValue32_v2'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaDeviceCanAccessPeer@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaFreeHost@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaRegisterFatBinary@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaGetDeviceCount@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaRegisterFunction@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaPopCallConfiguration@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaMalloc@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaGetDevice@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaMemcpyAsync@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaHostUnregister@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaEventCreateWithFlags@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `__cudaPushCallConfiguration@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaStreamCreate@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaGetErrorString@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaHostRegister@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaStreamDestroy@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaEventQuery@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaLaunchKernel@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaStreamSynchronize@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaEventRecord@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaMallocHost@libcudart.so.12'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cuStreamWaitValue64_v2'
/usr/bin/ld: /opt/cray/pe/mpich/8.1.30/gtl/lib/libmpi_gtl_cuda.so: undefined reference to `cudaEventDestroy@libcudart.so.12'
nvcc: error: linker command failed with exit code 1 (use -v to see invocation)

Indeed, Scale provides libcudart.so.1, patchelf'ing the gtl wont help, the symbols are versioned.

Cheers