并行编程实战——CUDA编程的MPI和MPS

一、说明

大家通过操作系统的学习都知道进程是分配资源的最小单位，进程的一个特点是多个进程间可以共享一个CPU（时间片轮转等方式），这样可以发挥CPU最大的资源潜力。同样，GPU会不会有需要这种情况呢？答案是肯定的。
而随着软硬件技术的发展，多GPU和多CPU一样逐渐增长起来，即出现了分布式的GPU运算，这就需要一些类似于计算机分布式访问的管理控制的技术。用来管理多进程分配、调度等情况。特别是在某些场景下，单个进程对GPU占用较短而任务又较多时，这种任务管理调度的需求就变得更迫切了。

二、多进程服务

GPU的任务可以被CPU的调度触发，不过，CPU是可以并行，但默认的情况下，GPU则是分时执行（串行）。为了能够达到与CPU一样的应用，GPU也提供了多进程服务模式，即MPS模式。它可以让多个进程能够共享一个GPU，同时执行其内核。不过，目前来看，它只能在Linux应用。可以使用下面的命令来启用和停用MPS：

//start
nvidia-cuda-mps-control -d
//stop
echo quit | nvidia-cuda-mps-control

也可以通过设置CUDA_MPS_ACTIVE_THREAD_PERCENTAGE环境变量来限制每个MPS客户端的GPU线程占用比率。合理的设置此变量值可以达到类似CPU中核心与线程的比例最佳的目标方式。另外也可以设置GPU为进程独占模式，由MPS自主管理GPU资源。

三、消息传递接口

如果每个GPU都可以被进程单独的占满，正如CPU中一样，每个进程都能够将CPU资源占满，那么切换进程调度反而会降低效率。但实际上，往往由于硬件的强大，计算任务大多数并不会完全占据硬件的资源。而此时就是多进程服务MPS的用武之地。
MPI，消息传递接口是一个种并行计算接口，它用来对多进程的跨计算单元进行管理和控制。包括CPU核心、GPU及节点等。类似于一个多进程的普通程序，可以运行在不同的CPU核心上，然后在整体进行最终的结果处理。
GPU的MPI也是如此管理和调度进程，每个MPI代表着应用程序的一些部分，调用CUDA将MPI进程分配到多个核心上进行工作。也就是说，MPI可以映射到相关的GPU资源上，既可以一比一的进行，也可以比较理想的将多个MPI进程映射到同一个GPU资源上。
为了达到映射MPI的目的，需要安装使用OpenMPI，大家可以根据自己的实际环境来选择安装相关的版本。

四、MPI和MPS的应用

MPI和MPS本身可以理解为一种互相协调应用的关系。如果说MPS是CUDA自身支持的一种技术，而MPI则是一种外部引入的对MPS进行协作工作的技术（和CPU的核心调用有密切关系）。MPS作为一种对GPU资源应用的抽象，MPI是对MPS资源应用的抽象。大家可以简单的理解为，MPI是一种服务器而MPS是一种客户端。GPU则是MPI调度MPS下的最终的执行载体。
MPI处理应用层，负责跨节点、跨进程的并行工作并负责任务的拆分和数据的处理。而MPS则在系统层，用于更合理的分配和使用GPU资源的使用，提高其利用率和吞吐量。
每一个单独的MPI进程就是MPS负责的多进程任务中的一个任务。它们两个最典型的应用场景为：进程数量大于GPU数量大于节点数量，简单的归结到一点，就是任务多，资源少。如果每个任务以资源的渴求并不明显是，就更适合于使用MPS了。另外，在某些情况下，也可以将多GPU集群拆分成类似一上述的情况，达到相应的应用场景进行应用。
传统的MPS一个最大的缺点就是隔离性不强，不如虚拟机好样几乎可以完全隔离相关的资源和应用。所以在使用MPI和MPS是，要重视在异常情况下以其它进程的影响。另外，还需要特别说明的是，必须密切关注任务本身对GPU占用与切换进程的费效比，如果达到一定域值，那么MPS的优势就因为资源的抢占导致效率下降，丧失了优势。
在CUDA框架中，支持了CUDA-aware MPI来利用GPUDirect RDMA进行节点间的通信。由于目前没有分布式环境，暂先不以其进行阐述说明，后面有机会再相机展开。

五、例程

下面看一个例程：

#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime.h>
#include<sys/time.h>

doublegetTime(){
structtimevaltv;
    gettimeofday(&tv, NULL);
return tv.tv_sec * 1000.0 + tv.tv_usec / 1000.0;
}

__global__ voidvectorAdd(constfloat* A, constfloat* B, float* C, int N){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
float sum = A[idx] + B[idx];
for (int k = 0; k < 100; k++) {
            sum = sqrtf(sum * sum);
        }
        C[idx] = sum;
    }
}

intmain(int argc, char** argv){
int proID = 0;
if (argc > 1){
       proID = atoi(argv[1]);
    }

int N = 1024 * 1024 * 10;  
size_t size = N * sizeof(float);

float *hA = (float*)malloc(size);
float *hB = (float*)malloc(size);
float *hC = (float*)malloc(size);


for (int i = 0; i < N; i++) {
        hA[i] = rand() / (float)RAND_MAX;
        hB[i] = rand() / (float)RAND_MAX;
    }

float *dA, *dB, *dC;
    cudaMalloc(&dA, size);
    cudaMalloc(&dB, size);
    cudaMalloc(&dC, size);

double startTime = getTime();
    cudaFree(0);

double ctxTime = getTime() - startTime;

    cudaMemcpy(dA, hA, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB, size, cudaMemcpyHostToDevice);

int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;

double startKernelTime = getTime();

for (int it = 0; it < 10; it++) {
        vectorAdd<<<numBlocks, blockSize>>>(dA, dB, dC, N);
        cudaDeviceSynchronize();   
    }

double kernelTime = getTime() - startKernelTime;

// clear
    cudaFree(dA);
    cudaFree(dB);
    cudaFree(dC);
free(hA);
free(hB);
free(hC);

return0;
}

因为没有Linux上的环境，所以无法给出测试，代码仅供参考。

六、总结

每个技术都有自己应用的场景，优势往往就在于场景的适配性上。或者只要有一点偏差，可能这种应用的优势就会消失。所以，真正的设计和开发者一定要掌握具体的应用特点，不能盲目的乐观于其技术的长处而忽略了其的局限性。正如古人所说“橘生淮南则为橘，生淮北则为枳”。