Multiple cuda IPC mappings without closing previous handles causing `st.volatile.global`,`ld.volatile.global` peer-memory accesses to misbehave

Accelerated Computing CUDA CUDA Programming and Performance
1

I’ve noticed that if I open two identical CUDA IPC mappings back-to-back—each time allocating a fresh buffer but never closing the old handle—then subsequent access to the peer memory region through ld.volatile.global/st.volatile.global for certain ranks does not seem to have any effect.

Below is a pseudo code to illustrate the issue which I am facing:

struct Signal {...}; // some structure for barrier design
struct RankSignal { Signal* signals[8]}; 
// Called once before any kernel launches
static
  void initialize_multi_gpu_barrier(RankSignals* sg__, Signal** self_sg__) {
    int nRanks, myRank;
    MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
    MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));
    CUDACHECK(cudaSetDevice(myRank));
    
    cudaIpcMemHandle_t self_data_handle;
    cudaIpcMemHandle_t data_handles[8];
    Signal* buffer;
    CUDACHECK(cudaMalloc(&buffer, sizeof(Signal)));
    CUDACHECK(cudaMemset(buffer, 0, sizeof(Signal)));
    CUDACHECK(cudaIpcGetMemHandle(&self_data_handle, buffer));
    MPICHECK(MPI_Allgather(/*sendbuf=*/&self_data_handle, 
                  /*sendcount=*/sizeof(cudaIpcMemHandle_t),
                  /*send type=*/MPI_BYTE, 
                  /*recvbuf=*/data_handles, 
                  /*recvcount=*/sizeof(cudaIpcMemHandle_t),
                  /*recv type=*/MPI_BYTE, 
                  MPI_COMM_WORLD));
    Signal* ipc_ptrs[8];
    for (int i = 0; i < nRanks; i++) {
      if (i == myRank)
        ipc_ptrs[i] = buffer;
      else
        CUDACHECK(cudaIpcOpenMemHandle((void**)&ipc_ptrs[i], data_handles[i],
                                      cudaIpcMemLazyEnablePeerAccess));
        sg__->signals[i] = ipc_ptrs[i];
    }
    (*self_sg__) = ipc_ptrs[myRank];
}

void finalize_multi_gpu_barrier(RankSignals* sg,
                                  Signal*      self_sg,
                                  int          myRank,
                                  int          nRanks) {
    cudaDeviceSynchronize();
    for (int r = 0; r < nRanks; ++r) {
      if (r == myRank) continue;
      CUDACHECK(cudaIpcCloseMemHandle(sg->signals[r]));
      sg->signals[r] = nullptr;
    }
    CUDACHECK(cudaFree(self_sg));
    sg->signals[myRank] = nullptr;
}

__global__ void some_kernel(...){
      ... some work (which does not write to global memory)
   //write to Signal object across peers through st.volatile.global.u32

  // read from Signal object, on the current GPU through 
  // ld.volatile.global.u32 --- waiting for peer GPUs to write to it -- in  spin loop // <--- line 1----+
     ... some work (which does not write to global memory)                                              |
}                                                                                                       |
                                                                                                        |
// Later, before launching a compute kernel:                                                            |
void run_work() {                                                                                       |
     // First set of work                                                                               |
     initialize_multi_gpu_barrier(...);                                                                 |
     launch some_kernel(...);                                                                           |
     // finalize_multi_gpu_barrier(...); // <---- line 2                                                |
                                                                                                        |
     // second set of work                                                                              |
     initialize_multi_gpu_barrier(...);                                                                 |
     launch some_kernel(...);  // ----------------------------------------------------------------------+	

}

I am using onlyld.volatile.global and st.volatile.global to access peer memory. (I am not sure if that is the root cause of the problem). Based on the documentation of volatile:

The semantics of volatile operations are equivalent to a relaxed memory operation with system-scope.

Since, I am not concerned about work writing to global memory, I did not use any fence (__threadfence_system) or acquire.sys/release.sys. And the approach just works fine with the line withfinalize_mulit_gpu_barrier() included (line-2).

When I allocate shared buffers back to back without a deallocation (line -2) in between, the second kernel_launch has some visibility issue of certain writes made by peer GPUs to the shared buffer.

Say Rank 0’s GPU keeps spinning on a zero valued shared location, to which say Rank 2’s GPU surely made a write (the write of Rank 2 is not seen by Rank 0 and it keeps on seeing zero). [I am not sure about the semantics of the level at which the flushes happen with just st.volatile/ld.volatile. Is some stale value being read by GPU in Rank 0 from L1 cache? From volatile semantics, I don’t expect that to be the case. I am not sure of the flushing semantics of acquire.sys/release.sys or __threadfence_system() either, but using them does not seem to have any effect here either.]

Can opening a second IPC mapping (with a new cudaMalloc) without ever closing the first one lead to strange behavior of writes when I use the peer pointer? In other words, do I have to call cudaIpcCloseMemHandle() (and free the original buffer) before creating a new IPC mapping, even if the handles refer to completely separate allocations? Any confirmation or pointer to the right best practice would help!

For context, I am using only one barrier barrier_at_start with some modifications from the vLLM source code.