內核運行時CUDA內存上的MVAPICH死鎖

Question

我試圖獲得一個與MVAPICH CUDA8一起工作的MPI-CUDA程序。我之前用openMPI成功運行過該程序，但是我想測試一下如果使用MVAPICH獲得更好的性能。不幸的是，如果一個CUDA內核在使用MVAPICH的同時運行，那麼程序會停留在MPI_Isend中。

我下載MVAPICH2-2.2並建立它與配置標誌

--enable-CUDA - 禁用MCAST

，使MPI CUDA的內存調用源。 mcast被禁用，因爲我無法在沒有標誌的情況下編譯它。當在同一時間沒有CUDA內核運行

export MV2_USE_CUDA=1 
export MV2_GPUDIRECT_GDRCOPY_LIB=/path/to/gdrcopy/ 
export MV2_USE_GPUDIRECT=1 

MPI_Isend/recv的做工精細：

我用下面的標誌運行應用程序之前。但是在我的程序中，MPI在內核運行時發送和接收來自GPU內存的數據並且接收GPU內存的數據是非常重要的。

我想出了兩種可能的原因。首先，由於某些原因，MVAPICH嘗試運行自己的CUDA內核以從GPU內存發送數據，並且該內核未被調度，因爲GPU已經被充分利用。第二種可能性：MVAPICH在某處使用cudaMemcpy（而不是異步版本），這會阻塞，直到內核完成執行。

有人可以確認我的一個假設嗎？ MVAPICH中有一個標誌解決了我不知道的這個問題嗎？

編輯：

這裏是 「simpel」 的代碼，說明我的問題。在使用openMPI執行代碼時，它會正確執行並終止。有了mvapich2，它會在標記爲MPI_Send函數時發生死鎖。

#include <stdio.h> 
#include <stdlib.h> 
#include <cuda.h> 
#include <mpi.h> 


__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(&argc, &argv); 
    int world_size, world_rank; 
    MPI_Comm_size(MPI_COMM_WORLD, &world_size); 
    MPI_Comm_rank(MPI_COMM_WORLD, &world_rank); 

    printf("Im rank %d\n", world_rank); 
    cudaSetDevice(world_rank); 

    double * dev_buffer; 
    cudaError_t err = cudaMalloc(&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(&stream); 
    cudaStreamCreate(&kernel_stream); 

    printf("[%d]launching kernel\n", world_rank); 
    kernel<<<208, 128, 0, kernel_stream>>>(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(); 
} 

Answer 1

我回到了這個問題，並使用gdb來調試代碼。

顯然，問題是在src/mpid/ch3/channels/mrail/src/gen2/ibv_send.c中實現的MVAPICH2的熱切協議。 eager協議使用不帶異步的cuda_memcpy，直到內核執行完成爲止。

通過將MV2_IBA_EAGER_THRESHOLD 1傳遞給mpirun，問題中發佈的程序正常運行。這可以防止MPI使用渴望的協議並使用rendez-vous協議。

修補MVAPICH2源代碼的確也解決了這個問題。我改變了同步cudaMemcpys到cudaMemcpyAsync中的文件

• SRC/MPID/CH3 /信道/ mrail/SRC /第二代/ ibv_send.c
• SRC/MPID/CH3 /信道/ mrail/SRC /第二代/ ibv_recv.c
• SRC/MPID/CH3/SRC/ch3u_request.c

在第三文件中的改變時，才需要對MPI_Isend/MPI_Irecv。其他MPI函數可能需要一些額外的代碼更改。