内核运行时 CUDA 内存上的 MVAPICH 死锁答案

【问题标题】：MVAPICH deadlocks on CUDA memory while kernel is running内核运行时 CUDA 内存上的 MVAPICH 死锁
【发布时间】：2017-07-14 04:39:12
【问题描述】：

我尝试让 MPI-CUDA 程序与 MVAPICH CUDA8 一起工作。我之前确实使用 openMPI 成功运行了该程序，但我想测试使用 MVAPICH 是否可以获得更好的性能。不幸的是，如果在使用 MVAPICH 时同时运行 CUDA 内核，程序会卡在 MPI_Isend 中。

我下载了 MVAPICH2-2.2 并使用配置标志从源代码构建它

--enable-cuda --disable-mcast

在 cuda 内存上启用 MPI 调用。 mcast 被禁用，因为没有标志我无法编译它。

我在运行应用程序之前使用了以下标志：

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

当同时没有 CUDA 内核运行时，MPI_Isend/recv 工作正常。但在我的程序中，重要的是 MPI 在内核运行时从 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();
}

【问题讨论】：

一个小的复制器会改善这个问题:)
添加了一个简单的代码示例来说明我的问题。
这看起来不应该工作。与 OpenMPI 一起工作的事实可能更像是一个意外。
你能详细说明为什么代码不应该工作吗？很抱歉代码不那么干净。正是我想出的短裤示例显示了问题。
您隐含地假设 MPI_Send 可以与正在运行的内核成功重叠，并且内核将同时保持内存一致性。但有很多原因可能导致情况并非如此。我不确定是否保证设备到设备的复制可以与正在运行的内核重叠，尤其是具有旨在完全占用相关设备的网格尺寸的内核。一般来说，“驻留内核”的整个想法（恕我直言）是一个破碎的想法。 linux上的内核启动延迟小于10us。如果这很关键，那么 MPI+GPus 是错误的解决方案

标签： c++ cuda mpi mpich

【解决方案1】：

我回到这个问题并使用gdb来调试代码。

显然，问题在于 src/mpid/ch3/channels/mrail/src/gen2/ibv_send.c 中实现的 MVAPICH2 的 Eager 协议。 Eager 协议使用不带异步的 cuda_memcpy，它会阻塞直到内核执行完成。

通过将 MV2_IBA_EAGER_THRESHOLD 1 传递给 mpirun，问题中发布的程序运行良好。这会阻止 MPI 使用 Eager 协议并改为使用 Rendez-vous 协议。

修补 MVAPICH2 源代码也可以解决问题。我将文件中的同步 cudaMemcpys 更改为 cudaMemcpyAsync

src/mpid/ch3/channels/mrail/src/gen2/ibv_send.c
src/mpid/ch3/channels/mrail/src/gen2/ibv_recv.c
src/mpid/ch3/src/ch3u_request.c

只有 MPI_Isend/MPI_Irecv 需要修改第三个文件。其他 MPI 函数可能需要一些额外的代码更改。

【讨论】：