title: MVAPICH deadlocks on CUDA memory while kernel is running
body: <p>I try to get a MPI-CUDA program working with MVAPICH CUDA8. I did run the program successfully with openMPI before but I want to test if I get better performance with MVAPICH. Unfortunately the program gets stuck in MPI_Isend if a CUDA kernel is running at the same time when using MVAPICH.</p>
  
  <p>I downloaded MVAPICH2-2.2 and built it from the source with the configuration flags </p>
  
  <p>--enable-cuda --disable-mcast</p>
  
  <p>to enable MPI calls on cuda memory. mcast was disabled because I could not compile it without the flag.</p>
  
  <p>I used the following flags before running the application:</p>
  
  <pre><code>export MV2_USE_CUDA=1
  export MV2_GPUDIRECT_GDRCOPY_LIB=/path/to/gdrcopy/
  export MV2_USE_GPUDIRECT=1
  </code></pre>
  
  <p>MPI_Isend/recv work fine when at the same time no CUDA kernel is running. But in my program it is important that MPI sends and receives data from and to GPU memory, while a kernel is running. </p>
  
  <p>I came up with two possible reasons for that behavior. First, MVAPICH tries to run his own CUDA kernel for some reason to send data from GPU memory and this kernel does not get scheduled because the GPU is already fully utilized. Second possibility: MVAPICH uses cudaMemcpy somewhere (not the async version), which blocks until the kernel finishes execution.</p>
  
  <p>Could someone confirm one of my assumptions? And is there a flag in MVAPICH that solves this problem that I am not aware of?</p>
  
  <p><strong>EDIT:</strong></p>
  
  <p>Here a "simpel" code that illustrates my problem. When executing the code with openMPI, it executes and terminates correctly. With mvapich2 it deadlocks at the marked MPI_Send function.</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;cuda.h&gt;
  #include &lt;mpi.h&gt;
  
  
  __global__ void kernel(double * buffer, int rank)
  {
      volatile double *buf = buffer;
      if(rank == 0){
          while(buf[0] != 3){}
      } else {
          while(buf[0] != 2){}
      }
  }
  
  
  int main(int argc, char **argv)
  {
      double host_buffer[1];
      MPI_Init(&amp;argc, &amp;argv);
      int world_size, world_rank;
      MPI_Comm_size(MPI_COMM_WORLD, &amp;world_size);
      MPI_Comm_rank(MPI_COMM_WORLD, &amp;world_rank);
  
      printf("Im rank %d\n", world_rank);
      cudaSetDevice(world_rank);
  
      double * dev_buffer;
      cudaError_t err = cudaMalloc(&amp;dev_buffer, sizeof(double));
      if(world_rank == 0){
          host_buffer[0] = 1;
          cudaError_t err = cudaMemcpy(dev_buffer, host_buffer, sizeof(double), cudaMemcpyHostToDevice);
          MPI_Send(dev_buffer, 1, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD);
          printf("[%d]First send does not deadlock\n", world_rank);
      }else {
          MPI_Recv(dev_buffer, 1, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
          printf("[%d]Received first message\n", world_rank);
      }
  
      cudaStream_t stream, kernel_stream;
      cudaStreamCreate(&amp;stream);
      cudaStreamCreate(&amp;kernel_stream);
  
      printf("[%d]launching kernel\n", world_rank);
      kernel&lt;&lt;&lt;208, 128, 0, kernel_stream&gt;&gt;&gt;(dev_buffer, world_rank);
  
      if(world_rank == 0){
          //rank 0
          host_buffer[0] = 2;
          cudaMemcpyAsync(
              dev_buffer, host_buffer, sizeof(double),
              cudaMemcpyHostToDevice,
              stream
          );
          cudaStreamSynchronize(stream);
  
          printf("[%d]Send message\n", world_rank);
          MPI_Send(dev_buffer, 1, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD); //mvapich2 deadlocks here
          printf("[%d]Message sent\n", world_rank);
  
          printf("[%d]Receive message\n", world_rank);
          MPI_Recv(dev_buffer, 1, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
          printf("[%d]Message received\n", world_rank);
  
          cudaStreamSynchronize(kernel_stream);
          printf("[%d]kernel finished\n", world_rank);
  
      } else {
          //rank 1
          printf("[%d]Receive message\n", world_rank);
          MPI_Recv(dev_buffer, 1, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
          printf("[%d]Message received\n", world_rank);
  
          cudaStreamSynchronize(kernel_stream);
          printf("[%d]kernel finished\n", world_rank);
  
          host_buffer[0] = 3;
          cudaMemcpyAsync(
              dev_buffer, host_buffer, sizeof(double),
              cudaMemcpyHostToDevice,
              stream
          );
          cudaStreamSynchronize(stream);
  
          printf("[%d]Send message\n", world_rank);
          MPI_Send(dev_buffer, 1, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD);
          printf("[%d]Message sent\n", world_rank);
  
      }
      printf("[%d]Stopped execution\n", world_rank);
      MPI_Finalize();
  }
  </code></pre>