当内核为运行时，MVAPICH 在 CUDA 内存上死锁

Question

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

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

--启用-cuda --禁用-mcast

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

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

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

MPI_Isend/recv 在没有 CUDA 内核同时运行ning 时工作正常。但是在我的程序中，重要的是 MPI 在内核运行ning 时从 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 协议。 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 函数可能需要一些额外的代码更改。

当内核为运行时，MVAPICH 在 CUDA 内存上死锁

MVAPICH deadlocks on CUDA memory while kernel is running

c++

cuda

mpi

mpich

当内核为 运行 时，MVAPICH 在 CUDA 内存上死锁

MVAPICH deadlocks on CUDA memory while kernel is running

c++

cuda

mpi

mpich

当内核为运行时，MVAPICH 在 CUDA 内存上死锁