カーネルの実行中にCUDAメモリにMVAPICHデッドロックが発生する

Question

MVAPICH CUDA8で動作するMPI-CUDAプログラムを取得しようとしています。前にopenMPIでプログラムを正常に実行しましたが、MVAPICHでより良いパフォーマンスが得られるかどうかをテストしたいと思います。残念なことに、MVAPICHを使用しているときにCUDAカーネルが同時に実行されている場合、プログラムはMPI_Isendに固執します。

私はMVAPICH2-2.2をダウンロードし、MPIは、CUDAメモリに呼び出しを可能にするために、構成フラグ

--enable-CUDA --disable-MCAST

とソースから構築されました。 mcastはフラグなしでコンパイルできなかったため無効になっています。

私は、アプリケーションを実行する前に、以下のフラグを使用：

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

MPI_Isend/recvの作業罰金を同時に何のCUDAカーネルが実行されていないとき。しかし私のプログラムでは、カーネルが動いている間にMPIがGPUメモリとデータを送受信することが重要です。

私は、そのような2つの理由が考えられました。まず、何らかの理由でMVAPICHが自分のCUDAカーネルをGPUメモリから送信しようとします。このカーネルはGPUがすでに完全に利用されているためスケジュールされません。 2番目の可能性：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が熱心なプロトコルを使用するのを防ぎ、代わりにランデブー・プロトコルを使用します。

MVAPICH2ソースコードにパッチを当てることで問題が解決します。私は、ファイルにcudaMemcpyAsyncに同期cudaMemcpysを変更

• のsrc/MPID/CH3 /チャンネル/ mrail/SRC/GEN2/ibv_send.c
• のsrc/MPID/CH3 /チャンネル/ mrail/SRC/GEN2/ibv_recv.c
• のsrc/MPID/CH3/SRC/ch3u_request.c
3番目のファイルの変更がMPI_Isend/MPI_Irecvのためにのみ必要とされている

。他のMPI関数では、コードをいくつか変更する必要があります。