We just noticed that the MPI_Irecv and MPI_Send commands in our reproducer have
typos. We fixed them but we are still getting the same results.
Bruce
From: 'Palmer, Bruce J' via Open MPI users <[email protected]>
Date: Tuesday, June 2, 2026 at 8:51 AM
To: [email protected] <[email protected]>
Cc: Panyala, Ajay <[email protected]>
Subject: Re: [EXTERNAL] [OMPI users] GPU-aware MPI
Any additional thoughts on this?
Bruce
From: 'Palmer, Bruce J' via Open MPI users <[email protected]>
Date: Thursday, May 28, 2026 at 9:00 AM
To: [email protected] <[email protected]>
Cc: Panyala, Ajay <[email protected]>
Subject: Re: [EXTERNAL] [OMPI users] GPU-aware MPI
Hi Edgar,
I should have posted more information about my reproducer. I’ve been looking at
running using 4 processors on 2 SMP nodes with each node having 1 GPU. Ranks 0
and 2 allocate memory on the GPU. Rank 0 sends a message to rank 3 and rand 2
sends a message to rank 1. Rank 3 opens the allocation created by rank 2 using
a cudaIpcOpenMemHandle call and uses the pointer returned by
cudaIpcOpenMemHandle in an MPI_Irecv call that expects a message from rank 0.
Similarly, rank 1 uses a pointer from cudaIpcOpenMemHandle to a GPU allocation
created by rank 0 in an MPI_Irecv call that expects a message from rank 2. This
reproduces the behavior of a onesided put call in Global Arrays using the
progress ranks runtime.
The essential feature is that a pointer to GPU memory that was allocated by a
different rank is being used in an MPI_Send/Recv call and that pointer is
obtained via a cudaIpcOpenMemHandle call.
Bruce
From: [email protected] <[email protected]> on behalf of Edgar
Gabriel <[email protected]>
Date: Wednesday, May 27, 2026 at 2:10 PM
To: [email protected] <[email protected]>
Cc: Panyala, Ajay <[email protected]>
Subject: RE: [EXTERNAL] [OMPI users] GPU-aware MPI
You don't often get email from [email protected]. Learn why this is
important<https://aka.ms/LearnAboutSenderIdentification>
Check twice before you click! This email originated from outside PNNL.
If I understand correctly, the issue that is being asked is: Process 0 has lets
say a buffer on GPU 0. That buffer has been imported to PE 1 using the GPU IPC
mechanism and is now mapped using a different virtual address into the address
space of PE 1. Can PE 1 use that new virtual address in a communication
operation with Open MPI? Is my understanding correct?
I think the challenge might be that if PE 1 uses that virtual address in a
Send/Recv operation, the internal protocols will try to (potentially) open an
IPC handle for that buffer as well, and I am not sure that PE1 can do that,
since the owner of that buffer is PE 0.
@bosilca ?
Thanks
Edgar
From: 'Pritchard Jr., Howard' via Open MPI users <[email protected]>
Sent: Wednesday, May 27, 2026 4:01 PM
To: [email protected]
Cc: Panyala, Ajay <[email protected]>
Subject: Re: [EXTERNAL] [OMPI users] GPU-aware MPI
Hello Bruce,
I think a little more info is needed. Could you post the output you get from
running
ompi_info
?
double check that the ompi_info you are running is in the same folder as the
mpicc you’re using.
thanks,
Howard
From: "'Palmer, Bruce J' via Open MPI users"
<[email protected]<mailto:[email protected]>>
Reply-To: "[email protected]<mailto:[email protected]>"
<[email protected]<mailto:[email protected]>>
Date: Wednesday, May 27, 2026 at 10:44 AM
To: Open MPI Users <[email protected]<mailto:[email protected]>>
Cc: "Panyala, Ajay" <[email protected]<mailto:[email protected]>>
Subject: [EXTERNAL] [OMPI users] GPU-aware MPI
Hi,
I’m trying to modify the Global Arrays library so that it supports global
arrays hosted on GPU memory. I have a version of the progress ranks runtime
that works by copying data to a buffer on the host before sending it to another
process located on a different SMP node but I’d like to eliminate the host
memory copies by using GPU-aware MPI. I’ve implemented this in the code but it
seems to be failing because I can’t use a pointer to GPU memory in an MPI send
or receive call that was created via a cudaIpcOpenMemHandle call.
Are pointers to GPU memory created from IpcMemHandles supposed to work with
GPU-aware MPI? This would be critical for our progress ranks runtime since it
would be effectively replacing the POSIX-shared memory strategy that we use for
handling messaging related to global arrays hosted on regular host memory.
I’ve included a small test code using just Cuda and MPI that reproduces the
strategy we want to use inside Global Arrays.
Bruce
To unsubscribe from this group and stop receiving emails from it, send an email
to
[email protected]<mailto:[email protected]>.
To unsubscribe from this group and stop receiving emails from it, send an email
to
[email protected]<mailto:[email protected]>.
To unsubscribe from this group and stop receiving emails from it, send an email
to
[email protected]<mailto:[email protected]>.
To unsubscribe from this group and stop receiving emails from it, send an email
to
[email protected]<mailto:[email protected]>.
To unsubscribe from this group and stop receiving emails from it, send an email
to
[email protected]<mailto:[email protected]>.
To unsubscribe from this group and stop receiving emails from it, send an email
to [email protected].
#include "mpi.h"
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#define NLEN 1024*1024
#define MPI_TAG 12383
int main(int argc, char **argv)
{
int me, nprocs, iproc;
int *sbuf, *rbuf, *gpuBuf, *devbuf;
void *vbuf;
int nghbr, sndr;
int *hostIDs;
int i;
int lowest;
int ngpu, totgpu, myGPU;
cudaIpcMemHandle_t *handles;
cudaIpcMemHandle_t handle;
MPI_Request req;
MPI_Status status;
int my_master;
int *masters;
int t_ok, ok;
/* initialize MPI */
MPI_Init(&argc, &argv);
MPI_Comm_rank(MPI_COMM_WORLD, &me);
MPI_Comm_size(MPI_COMM_WORLD, &nprocs);
/* Find host IDs for all processors */
hostIDs = (int*)malloc(sizeof(int)*nprocs);
sbuf = (int*)malloc(sizeof(int)*nprocs);
for (i=0; i<nprocs; i++) sbuf[i] = 0;
sbuf[me] = gethostid();
MPI_Allreduce(sbuf,hostIDs,nprocs,MPI_INT,MPI_SUM,MPI_COMM_WORLD);
/* Find the lowest rank with same host ID */
lowest = nprocs;
for (i=0; i<nprocs; i++) {
if (hostIDs[i] == hostIDs[me] && i<lowest) lowest = i;
}
/* find the highest rank with the same host ID */
for (i=me; i<nprocs; i++) {
if (hostIDs[i] == hostIDs[me]) my_master = i;
else break;
}
/* create global list of masters for each process */
masters = (int*)malloc(sizeof(int)*nprocs);
for (i=0; i<nprocs; i++) sbuf[i] = 0;
sbuf[me] = my_master;
MPI_Allreduce(sbuf,masters,nprocs,MPI_INT,MPI_SUM,MPI_COMM_WORLD);
free(sbuf);
/* find total number of GPUs visible on each node. Check to see if this
* is greater than or equal to number of processors on each node minus one */
cudaGetDeviceCount(&ngpu);
t_ok = 0;
if (ngpu >= my_master-lowest) t_ok = 1;
MPI_Allreduce(&t_ok,&ok,1,MPI_INT,MPI_PROD,MPI_COMM_WORLD);
if (!ok) {
if (me == 0) printf("Not enough GPUs per node\n");
MPI_Abort(MPI_COMM_WORLD,1);
}
/* if not a master, allocate memory on GPU */
if (me != my_master) {
myGPU = me-lowest;
cudaSetDevice(myGPU);
cudaMalloc(&vbuf, sizeof(int)*NLEN);
gpuBuf = (int*)vbuf;
} else {
myGPU = -1;
}
handles = (cudaIpcMemHandle_t*)malloc(sizeof(cudaIpcMemHandle_t)*nprocs);
cudaIpcGetMemHandle(&handle,gpuBuf);
MPI_Allgather(&handle,sizeof(cudaIpcMemHandle_t),MPI_BYTE,handles,
sizeof(cudaIpcMemHandle_t),MPI_BYTE,MPI_COMM_WORLD);
/* initialize send buf */
sbuf = (int*)malloc(sizeof(int)*NLEN);
if (me != my_master-1 && me != my_master) {
nghbr = (me+1)%nprocs;
for (i=0; i<NLEN; i++) {
sbuf[i] = ((me+1)%nprocs)*NLEN+i;
}
} else if (me != my_master) {
nghbr = (me+2)%nprocs;
for (i=0; i<NLEN; i++) {
sbuf[i] = ((me+2)%nprocs)*NLEN+i;
}
} else {
nghbr = -1;
}
printf("p[%d] nghbr: %d\n",me,nghbr);
/* post GPU-aware Irecv, if necessary. */
if (me == my_master) {
int id;
if (lowest > 0) {
sndr = lowest-2;
} else {
sndr = nprocs-2;
}
id = 0;
cudaSetDevice(id);
printf("p[%d] opening handle on %d, receiving from %d\n",me,lowest,sndr);
fflush(stdout);
cudaIpcOpenMemHandle(&vbuf,handles[lowest],cudaIpcMemLazyEnablePeerAccess);
devbuf = (int*)vbuf;
cudaDeviceSynchronize();
MPI_Irecv(devbuf,NLEN,MPI_INT,sndr,MPI_TAG,MPI_COMM_WORLD,&req);
}
/* send data to neighbor */
if (nghbr != -1 && hostIDs[me] == hostIDs[nghbr]) {
/* Data is on the same SMP node so use a memcpy */
int id = nghbr-lowest;
cudaSetDevice(id);
cudaIpcOpenMemHandle(&vbuf,handles[nghbr],cudaIpcMemLazyEnablePeerAccess);
devbuf = (int*)vbuf;
cudaDeviceSynchronize();
cudaMemcpy(devbuf,sbuf,sizeof(int)*NLEN,cudaMemcpyHostToDevice);
cudaIpcCloseMemHandle(vbuf);
cudaDeviceSynchronize();
} else if (nghbr != -1) {
printf("p[%d] posting message to %d\n",me,masters[nghbr]);
fflush(stdout);
/* Data is on different SMP node so use MPI */
MPI_Send(sbuf,NLEN,MPI_INT,masters[nghbr],MPI_TAG,MPI_COMM_WORLD);
}
printf("p[%d] Posting wait\n",me);
fflush(stdout);
/* wait on Irecv, if necessary */
if (me == my_master) {
int id = 0;
MPI_Wait(&req, &status);
cudaIpcCloseMemHandle(devbuf);
cudaDeviceSynchronize();
}
MPI_Barrier(MPI_COMM_WORLD);
printf("p[%d] Completed wait\n",me);
fflush(stdout);
/* check for correctness */
t_ok = 1;
if (me != my_master) {
/* copy data from GPU to host */
int id = me-lowest;
cudaSetDevice(id);
cudaMemcpy(sbuf,gpuBuf,sizeof(int)*NLEN,cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
for (i=0; i<NLEN; i++) {
if (sbuf[i] != me*NLEN+i) t_ok = 0;
}
}
MPI_Allreduce(&t_ok,&ok,1,MPI_INT,MPI_PROD,MPI_COMM_WORLD);
if (!ok) {
if (me == 0) printf("Error moving data\n");
MPI_Abort(MPI_COMM_WORLD,1);
}
if (ok && me == 0) {
printf("Completed test\n");
fflush(stdout);
}
MPI_Finalize();
return 0;
}
