Multi-GPU HIP distributed factorization fails on AMD MI300A with Multiple MPI ranks

[safa0rhan]

Hi,

I’m reporting a distributed multi-GPU failure in STRUMPACK on AMD/ROCm. I tried to reduce this carefully and rule out local setup mistakes before posting.

Summary

I can build STRUMPACK successfully with:

• MPI
• HIP
• OpenMP
• ParMETIS
• SLATE

and:

• 1 MPI rank / 1 GPU works
• 2 MPI ranks / 2 GPUs fail during distributed GPU factorization

I also tested the fix from issue #126 (8f095d2a), but it does not resolve the problem in my case.

The failure mode depends on OMP_NUM_THREADS:

• OMP_NUM_THREADS=8
  • fails in src/dense/HIPWrapper.cpp:351
  • hipBLAS assertion failed: 1
• OMP_NUM_THREADS=1
  • segfault in libblaspp.so / libslate.so
  • specifically in blas::device_free, called from:
    • slate::MatrixStorage<double>::allocateBatchArrays(...)
    • slate::impl::getrf_tntpiv(...)

So this appears to be a distributed ROCm multi-GPU failure somewhere in the STRUMPACK + SLATE/BLAS++ path.

Environment

Hardware:

• AMD Instinct MI300A
• gfx942

Software stack:

• GCC 14
• ROCm 7.2
• OpenMPI 5.0
• METIS 5.1
• ParMETIS 4.0
• AOCL 5.2
• ScaLAPACK
• HDF5 1.14.x

STRUMPACK:

• built from tag v8.0.0
• tested also after cherry-picking:
  • 8f095d2a Set device cuda/hip for each openmp thread. This isn't ideal as it assumes the threads are reused.

SLATE stack:

• locally built from v2025.05.28
• matching BLAS++ / LAPACK++

STRUMPACK configure highlights:

• STRUMPACK_USE_MPI=ON
• STRUMPACK_USE_HIP=ON
• STRUMPACK_USE_OPENMP=ON
• STRUMPACK_USE_CUDA=OFF
• TPL_ENABLE_SLATE=ON
• TPL_ENABLE_PARMETIS=ON
• HIP architecture set to gfx942

Note about ROCm 7.2 source compatibility

To build v8.0.0 against ROCm 7.2, I had to patch STRUMPACK source to replace the older hipBLAS complex typedef names:

• hipblasComplex -> hipFloatComplex
• hipblasDoubleComplex -> hipDoubleComplex

This was only a local source compatibility patch in STRUMPACK, not a modification of any vendor libraries.

What works

1 MPI rank / 1 GPU

The benchmark runs successfully with GPU enabled.

This validates:

• HIP build works
• runtime libraries are correct
• single-rank GPU factorization/solve works

MPI sanity checks

Plain MPI works:
mpirun -np 2 hostname

Standalone MPI_Init_thread(MPI_THREAD_MULTIPLE) test works on 2 ranks and returns:
provided=3

So the OpenMPI stack on this system can provide MPI_THREAD_MULTIPLE.

Reproducer

Run on one node with 2 MPI ranks and 2 GPUs visible, then:

mpirun -np 2 ./strumpack_benchmark <matrix_dir> --sp_enable_gpu --sp_reordering_method parmetis

The benchmark gets through:

• matrix read
• reorder
• symbolic factorization

and fails in numerical factorization.

Failure mode 1: OMP_NUM_THREADS=8

Run:
export OMP_NUM_THREADS=8
unset SLATE_GPU_AWARE_MPI
mpirun -np 2 ./strumpack_benchmark <matrix_dir> --sp_enable_gpu --sp_reordering_method parmetis

Failure:
hipBLAS assertion failed: 1 .../src/dense/HIPWrapper.cpp 351

The failing line is the hipblasDgemm(...) call in HIPWrapper.cpp.

Failure mode 2: OMP_NUM_THREADS=1

Run:
export OMP_NUM_THREADS=1
unset SLATE_GPU_AWARE_MPI
mpirun -np 2 ./strumpack_benchmark <matrix_dir> --sp_enable_gpu --sp_reordering_method parmetis

Failure:

• segfault in libblaspp.so
• blas::device_free(...)
• called from:
  • slate::MatrixStorage<double>::allocateBatchArrays(...)
  • slate::impl::getrf_tntpiv(...)

So with 1 OpenMP thread, the failure moves from hipblasDgemm to the SLATE / BLAS++ GPU memory-management path.

SLATE_GPU_AWARE_MPI test

I tested both:

• export SLATE_GPU_AWARE_MPI=0
• export SLATE_GPU_AWARE_MPI=1

This did not resolve the OMP_NUM_THREADS=1 crash. The failure remained in the same libblaspp / libslate path.

Issue #126 relation

I checked issue #126 and confirmed that my original v8.0.0 tree did not include commit 8f095d2a.

I cherry-picked 8f095d2a, rebuilt, reinstalled, and reran the tests above.

Result:

• OMP_NUM_THREADS=8: still fails with hipBLAS assertion failed: 1 at HIPWrapper.cpp:351
• OMP_NUM_THREADS=1: still segfaults in libblaspp / libslate

So 8f095d2a does not fix this ROCm/MI300A case.

Why I think this is a real upstream distributed GPU issue

At this point I believe I have ruled out the common local mistakes:

• build is successful
• dependency chain is coherent
• runtime libs are correct
• single-rank GPU works
• MPI itself works
• issue Failure when using more than 1 GPU in STRUMPACK MPI #126 patch was applied and tested

The remaining failures are now inside:

• STRUMPACK HIP wrapper (hipblasDgemm)
• or SLATE/BLAS++ GPU memory handling (device_free, allocateBatchArrays, getrf_tntpiv)

Questions

1. Is distributed multi-GPU HIP on AMD/ROCm expected to work with:

   • STRUMPACK v8.0.0
   • SLATE v2025.05.28
   • ROCm 7.2
   • MI300A / gfx942
   • OpenMPI 5.0

2. Does the hipBLAS assertion failed: 1 at HIPWrapper.cpp:351 suggest a known ROCm multi-rank handle/context issue in STRUMPACK?

3. Does the OMP_NUM_THREADS=1 segfault in blas::device_free / slate::MatrixStorage::allocateBatchArrays suggest a BLAS++ / SLATE ROCm multi-GPU problem rather than STRUMPACK proper?

4. Is there a recommended STRUMPACK / SLATE / BLAS++ version combination for ROCm multi-GPU that is known to be more stable than:

   • STRUMPACK v8.0.0
   • SLATE v2025.05.28

[pghysels]

Question 1. Yes, it is expected to work, but it's not tested.

Is your fix for the complex types in ROCM 7.2 the same as this: cfba574 ?

Question 2. I have not seen this, but yes it suggests something wrong in the handle/context in STRUMPACK.

Question 3. That does seem to be a SLATE problem, unless it is related to Q2. You could try not using the tntpiv version of slate::getrf by commenting out this line:

STRUMPACK/src/sparse/fronts/FrontDenseMPI.cpp

Line 160 in 0437c39

slate_opts_.insert({slate::Option::MethodLU, slate::MethodLU::CALU});

Question 4. Unfortunately not. Both STRUMPACK and I think SLATE have not had much development recently.

[safa0rhan]

Thanks for the quick reply, I just tested both suggestions.

I applied cfba574f for the HIP complex types and rebuilt/reinstalled successfully. That fixes the ROCm 7.2 source compatibility issue, but it does not change the distributed multi-GPU runtime failures.

I also commented out this line in src/sparse/fronts/FrontDenseMPI.cpp:

slate_opts_.insert({slate::Option::MethodLU, slate::MethodLU::CALU});
Then I rebuilt/reinstalled and reran the 2-rank benchmark.

Result: no change in behavior.

With OMP_NUM_THREADS=1, it still segfaults in:

libblaspp.so: blas::device_free(...)
libslate.so: slate::MatrixStorage<double>::allocateBatchArrays(...)
libslate.so: slate::impl::getrf_tntpiv(...)
With OMP_NUM_THREADS=8, it still fails with:

hipBLAS assertion failed: 1
src/dense/HIPWrapper.cpp:351
One possibly important detail: even after commenting out MethodLU = CALU, the OMP_NUM_THREADS=1 backtrace still shows slate::impl::getrf_tntpiv(...).

So at least in this build/run, removing the explicit CALU option did not eliminate or visibly change the SLATE-side failure path.

[pghysels]

Try changing CALU to PartialPiv.
https://icl.bitbucket.io/slate/namespaceslate.html#aca16539a2a57b8f75e9f4cfbdee1f2c9

[safa0rhan]

Thank you so much for your help again. I reran the LU-method tests carefully and got the following.

1. MethodLU = CALU

slate_opts_.insert({slate::Option::MethodLU, slate::MethodLU::CALU});

OMP_NUM_THREADS=1
crashes in:
libblaspp.so: blas::device_free(...)
libslate.so: slate::MatrixStorage::allocateBatchArrays(...)
libslate.so: slate::impl::getrf_tntpiv(...)

OMP_NUM_THREADS=8
also crashes in the same SLATE/BLAS++ path
backtrace goes through:
slate::impl::getrf_tntpiv(...)

2. No explicit MethodLU

// no slate::Option::MethodLU insertion

OMP_NUM_THREADS=1
crashes in:
libblaspp.so: blas::device_free(...)
libslate.so: slate::MatrixStorage::allocateBatchArrays(...)
libslate.so: slate::impl::getrf(...)

OMP_NUM_THREADS=8
also crashes in the same SLATE/BLAS++ path
backtrace goes through:
slate::impl::getrf(...)

3. MethodLU = PartialPiv

slate_opts_.insert({slate::Option::MethodLU, slate::MethodLU::PartialPiv});

OMP_NUM_THREADS=1
crashes in:
libblaspp.so: blas::device_free(...)
libslate.so: slate::MatrixStorage::allocateBatchArrays(...)
slate::impl::getrf(...)

OMP_NUM_THREADS=8
also crashes in the same SLATE/BLAS++ path
backtrace goes through:
slate::impl::getrf(...)

Interpretation

So the MethodLU choice is definitely taking effect:

• CALU leads to slate::impl::getrf_tntpiv(...)
• PartialPiv and no explicit method lead to slate::impl::getrf(...)

But all three cases still fail in the same broader place:

• blas::device_free(...)
• slate::MatrixStorage<double>::allocateBatchArrays(...)

So at least in this setup, changing the LU method changes the SLATE factorization path, but does not avoid the crash.

[safa0rhan]

Update: reduced reproducer using STRUMPACK's own test

I have done additional diagnostics and found a smaller reproducer that does not involve my benchmark code or HDF5 reader.

The failure can be reproduced with STRUMPACK's own test:

export OMP_NUM_THREADS=1
export SLATE_GPU_AWARE_MPI=0

srun -n 2 \
  --gpus-per-task=1 \
  --gpu-bind=map_gpu:0,1 \
  /strumpack_amd_build/src/STRUMPACK/build/test/test_structure_reuse_mpi \
  /strumpack_amd_build/src/STRUMPACK/examples/sparse/data/pde900.mtx \
  --sp_enable_gpu \
  --sp_matching 0 \
  --sp_reordering_method parmetis

This fails on one node with 2 MPI ranks and 2 MI300A GPUs.

The matrix is tiny:

reading 900 by 900 matrix with 4,380 nnz
estimated memory usage (exact solver) = 0.171296 MB
So this does not appear to be memory pressure or a matrix-size issue.

The failure is:

Segmentation fault
blas::device_free(...)
slate::MatrixStorage<double>::allocateBatchArrays(...)
slate::impl::getrf_tntpiv<Target::Devices, double>(...)
This is the same failure path I observed in my benchmark.

I also verified Slurm GPU binding. With:

--gpus-per-task=1 --gpu-bind=map_gpu:0,1
each MPI rank sees:

ROCR_VISIBLE_DEVICES=0
but rocm-smi --showuniqueid reports different physical GPU unique IDs for the two ranks. So Slurm is correctly assigning one physical GPU per rank and remapping each assigned GPU to local ordinal 0.

This means the failure is independent of:

my JOREK matrix
my HDF5 reader
my distributed CSR construction
residual computation
large matrix memory pressure
The failing path appears to be:

STRUMPACK distributed multifrontal factorization
  -> FrontDenseMPI
  -> SLATE distributed dense frontal LU
  -> slate::getrf_tntpiv(Target::Devices)
  -> SLATE / BLAS++ device workspace management
This also looks very similar to issue #126, except that my case is HIP/ROCm on MI300A rather than CUDA. In both cases:

1 MPI rank / 1 GPU works
multiple MPI ranks / GPUs fail
test_structure_reuse_mpi fails
the failure happens during multifrontal numerical factorization
the matrix is small enough that memory pressure is unlikely
One additional note: I found and fixed a bug in my benchmark where STRUMPACK options were not being explicitly parsed with set_from_options(...). However, this does not explain the failure reported here, because the reduced reproducer above uses STRUMPACK's own test_structure_reuse_mpi executable and still fails.

My current question is therefore more specific:

Is the FrontDenseMPI path using SLATE Target::Devices expected to work for one MPI rank per GPU on multiple GPUs?

In particular, is the following path currently tested/supported on CUDA or HIP?

FrontDenseMPI
  -> slate::Matrix::fromScaLAPACK(...)
  -> slate::getrf(..., Target::Devices)
  -> getrf_tntpiv
  -> MatrixStorage::allocateBatchArrays
If this is expected to work, I can provide the complete backtrace and build configuration. If it is not currently supported or known to be fragile, it would be useful to document that distributed STRUMPACK GPU factorization with more than one MPI rank/GPU is not currently reliable on this stack.