MPI 直接传递 GPU buffer 数据的原理——调试 libmpi.so from MPI with-cuda

原创已于 2024-11-25 19:35:29 修改 · 577 阅读

8 ·

本内容遵循CC 4.0 BY-SA版权协议

标签

#HPC #LLM #GPU #CUDA #MPI

于 2024-11-24 18:23:15 首次发布

由于新的机器中遇到了新的问题，故搭建环境部分再写一遍

1, 搭建环境

1.1 环境描述

（1) ubuntu 22.04

(2) 中安装了 cuda-12.4 和 rocm-6.0.2

(3) Intel CPU 和两张 RTX 2080TI + NVLink，不带MIxxx 卡

1.2 搭建 openmpi debug 环境

下载代码：

git clone https://github.com/open-mpi/ompi.git
cd ompi
git checkout v5.0.6
git submodule update --recursive --init

配置：

mkdir build/

cd build/

在上述环境中，需要按找如下做配置

$ ../configure --prefix=/home/hipper/ex_openmpi_withcuda/tmp4_withcuda_dbg506/localmpi --enable-debug --with-cuda=/usr/local/cuda --with-cuda-libdir=/usr/local/cuda/lib64/stubs

而如下配置

$ ../configure --prefix=/home/hipper/ex_openmpi_withcuda/tmp4_withcuda_dbg506/localmpi --enable-debug --with-cuda=/usr/local/cuda --with-cuda-libdir=/usr/local/cuda/lib64/stubs --without-rocm

仅多了 --without-rocm，则会导致出现错误，原因待查：

编译安装：

$ make -j
$ make install

验证安装成功：

设置环境变量：

hipper@hipper-G21:~/ex_openmpi_withcuda/tmp4_withcuda_dbg506/localmpi/bin$ export PATH=$PWD:$PATH

hipper@hipper-G21:~/ex_openmpi_withcuda/tmp4_withcuda_dbg506/localmpi/lib$ export LD_LIBRARY_PATH=$PWD

$ export PATH=$PWD:$PATH

$ export LD_LIBRARY_PATH=$PWD

2. 示例程序编译运行

源码：

twoGPU_snd_rec_dev_buf.c

// oneGPU_snd_rec_dev_buf.c
#include <stdio.h>
#include <string.h>
#include <mpi.h>
#include <cuda_runtime.h>
 
int main(int argc, char *argv[])
{
    char message[20];
    int myrank, tag=99;
    MPI_Status status;
 
    /* Initialize the MPI library */
    MPI_Init(&argc, &argv);
    /* Determine unique id of the calling process of all processes participating
       in this MPI program. This id is usually called MPI rank. */
    MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
        int size = 60;
    if (myrank == 0) {
        char s_buf_h[60]="Hello cuda mpi in same one GPU";
        char* s_buf_d = NULL;
        cudaMalloc((void**)&s_buf_d, 60);
        cudaMemcpy(s_buf_d, s_buf_h, 60, cudaMemcpyHostToDevice);
 
        //MPI rank 0
        MPI_Send(s_buf_d,size,MPI_CHAR,1,100,MPI_COMM_WORLD);
    } else {
        char* r_buf_d = NULL;
        char r_buf_h[60];
        cudaMalloc((void**)&r_buf_d, 60);
        //MPI rank n-1
        MPI_Recv(r_buf_d,size,MPI_CHAR,0,100,MPI_COMM_WORLD, &status);
 
        cudaMemcpy(r_buf_h, r_buf_d, size, cudaMemcpyDeviceToHost);
        printf("received %s\n", r_buf_h);
    }
 
    /* Finalize the MPI library to free resources acquired by it. */
    MPI_Finalize();
    return 0;
}

编译 debug 示例程序：

Makefile

EXE := twoGPU_snd_rec_dev_buf
all: $(EXE)
 
%: %.c
        mpicc -g $< -o $@ $(INC) $(LD_FLAGS)
 
INC      := -I/usr/local/cuda/include
LD_FLAGS := -L/usr/local/cuda/lib64 -lcudart
 
.PHONY: clean
clean:
        -rm -rf $(EXE)

编译运行：

$ make
$ mpiexec  -np 2 ./twoGPU_snd_rec_dev_buf

成功执行：

3. 调试程序 with xterm

3.0 安装启动 tightVNC

（Virtual Network Computing）

3.0.1 安装 tightVNC 和 xterm

sudo apt install xfce4 xfce4-goodies tightvncserver
sudo apt-get install  xfonts-base
sudo apt install xterm

酌情执行：

sudo apt-get install xfonts-100dpi

sudo apt-get install xfonts-75dpi

3.0.2 设置vnc密码

$ vncpasswd

键入两次密码；

删除密码：

$ rm ~/.vnc/passwd

3.0.3 启动 tightVNC

$ tightvncserver :1

关闭刚才启动的 tightVNC

$  tightvncserver -kill :1

重新启动并指定分辨率：

$ tightvncserver :1 -geometry 1920x1080

3.0.4 配置防火墙

让防火墙打开端口 5901：

$ sudo ufw allow 5901

防火墙命令解释：

ufw: 是缩写 of "Uncomplicated Firewall"，是 Ubuntu 中一个比较用户友好的防火墙管理工具；
allow: 指示防火墙允许通过指定的端口；
5901: 这是要开放的端口号。在这个密令中中，5901 通常用于 VNC服务。

3.1 登陆 VNC

在服务器上启动 tightVNC后，现在在 MacOS 登陆 VNC

MacOS 登陆Linux 远程桌面：
Finder cmd+K
vnc://ip地址:5901
输入vncserver的密码

例如：

vnc://hanmeimei@192.168.1.102:5901

输入刚才设置的密码后登陆：

3.2 开始调试

打开 terminal，进入服务器中 mpi项目的目录，检查环境变量：

符合期待，可以执行app：

执行看看：

使用 xterm gdb 来调试 app：

$ mpiexec -np 2  xterm -e gdb  ./two_dev

会产生两个 xterm 虚拟终端，因为 -np 指定了 2 个rank。

分别输入参数，如果有参数的话；

输入 start 启动程序，

并轮换输入 n +回车

3.3 追踪 MPI_Send(dev_buffer) 的实现

rank 0 中，

MPI_Send 定义

打开源代码：

继续跟踪，试图找到与cuda 有关的内容：

b cuMemcpyAsync

b cuMemcpy

接下来需要找时间分析这个调用堆栈

而 rank 1 中，

MPI_Recv() 的调用堆栈：

综合上边两个调用堆栈的信息，我们发现如下图的等值关系，MPI_Send的 dest，便是MPI_Recv 的 src地址。这两者处在不同的rank中，那么它们是不是同一个块显存空间呢？

那么，接下来我们调研 MPI_Send 走到 cuMemcpyAsync(dest, src,...)中，dest的来源，declaration。

通过跟踪代码可以发现，openmpi在实现这个MPI_Send 时，先将 DevBuffer的数据 cp到 HostMem，这块 HostMem是作为 rank 之间 sharedMem而存在。

这涉及到如下结构体：

/*
 * Shared Memory (SM) component instance.
 */
mca_btl_sm_component_t mca_btl_sm_component = {
    .super =
        {
            /* First, the mca_base_component_t struct containing meta information
               about the component itself */
            .btl_version =
                {
                    MCA_BTL_DEFAULT_VERSION("sm"),
                    .mca_open_component = mca_btl_sm_component_open,
                    .mca_close_component = mca_btl_sm_component_close,
                    .mca_register_component_params = mca_btl_sm_component_register,
                },
            .btl_data =
                {/* The component is checkpoint ready */
                 .param_field = MCA_BASE_METADATA_PARAM_CHECKPOINT},

            .btl_init = mca_btl_sm_component_init,
            .btl_progress = mca_btl_sm_component_progress,
        } /* end super */
};

在函数int mca_btl_sm_sendi(。。。）的如下代码段，获得共享内存的指针：


    /* allocate a fragment, giving up if we can't get one */
    frag = (mca_btl_sm_frag_t *) mca_btl_sm_alloc(btl, endpoint, order, length,
                                                  flags | MCA_BTL_DES_FLAGS_BTL_OWNERSHIP);
    if (OPAL_UNLIKELY(NULL == frag)) {
        if (descriptor) {
            *descriptor = NULL;
        }

        return OPAL_ERR_OUT_OF_RESOURCE;
    }

    /* fill in fragment fields */
    frag->hdr->len = length;
    frag->hdr->tag = tag;

    /* write the match header (with MPI comm/tag/etc. info) */
    memcpy(frag->segments[0].seg_addr.pval, header, header_size);

    /* write the message data if there is any */
    /* we can't use single-copy semantics here since as caller will consider the send
       complete when we return */
    if (payload_size) {
        uint32_t iov_count = 1;
        struct iovec iov;

        /* pack the data into the supplied buffer */
        iov.iov_base = (IOVBASE_TYPE *) ((uintptr_t) frag->segments[0].seg_addr.pval + header_size);
        iov.iov_len = length = payload_size;

        (void) opal_convertor_pack(convertor, &iov, &iov_count, &length);

其中共享内存地址存储在 iov.iov_base = 之中。

并通过 (void) opal_convertor_pack(convertor, &iov, &iov_count, &length); 调用传递给 cuMemcpyAsync(dest, src, ...);

作为 dest 指针。

显然这个效率并不高。

作为一个印证，现在用nvprof跟踪mpi app 的执行：