如何通过拆分数据和高效使用多线程多 GPU 来计算大型矩阵向量乘法？答案

【问题标题】：How to compute large matrix-vector multiplication by splitting data and using multithreading multiple GPU efficiently?如何通过拆分数据和高效使用多线程多 GPU 来计算大型矩阵向量乘法？
【发布时间】：2020-10-02 00:38:57
【问题描述】：

我是 GPU 计算的新手。我想计算XX^TY 的结果，其中 X 是一个大矩阵（N 乘 P），Y 是一个向量（N 乘 1）。现在我想将矩阵 X 分成几个块，并首先计算每个设备的本地结果。最后，我使用全环归约得到总和结果。我的演示代码如下

typedef unsigned long long uint64;

struct DevicePro {
    // pointer of device memory
    float* deviceData;
    float* deviceVec;
    float* deviceRes;

    // cuda context
    cudaStream_t stream;

    int gpuID;
    uint64 blockSize; // block size of each matrix-vec computation
    uint64 dataSplitIndex; // data start index for each device
    uint64 resultLength;
}


// allocate memory on each device
    for (int i = 0; i < numDevice; ++i) {
        int currDevice = deviceProList[i];
        cudaSetDevice(currDevice.gpuID);

        gpuErrchk(cudaStreamCreate(&currDevice.stream));
        gpuErrchk(cudaMalloc(reinterpret_cast<void **>(&currDevice.deviceData), rows * currDevice.blockSize * sizeof(float)));
        gpuErrchk(cudaMalloc(reinterpret_cast<void **>(&currDevice.deviceVec), rows * sizeof(float)));
        gpuErrchk(cudaMalloc(reinterpret_cast<void **>(&currDevice.deviceRes), rows * sizeof(float)));
    }

    // start computation
    // first we need to know how many blocks we have
    uint64 numBlock = static_cast<uint64>(dataSplitSize / deviceProList[0].blockSize) + 1;
    for (int i = 0; i < numDevice; ++i) {
        for (uint64 block = 0; block < numBlock; ++block) {
            DevicePro currDevice = deviceProList[i];
            //transfer data and launch kernel
        }
    }


    // free cuda memory
    for (int i = 0; i < numDevice; ++i) {
        int currDevice = deviceProList[i];
        cudaSetDevice(currDevice.gpuID);

        gpuErrchk(cudaFree(currDevice.deviceData));
        gpuErrchk(cudaFree(currDevice.deviceRes));
        gpuErrchk(cudaFree(currDevice.deviceVec));
    }

即使我拆分它们，我也无法将整个数据放入 GPU 内存中。这就是为什么我需要这个循环for (uint64 block = 0; block < numBlock; ++block)。在这里，我假设所有设备的块数相同。我的问题如下

我可以使用多线程在每个设备上启动数据传输和内核计算吗？否则，设备 2 必须等到设备 1 完成其所有调用。

【问题讨论】：

我不相信你最后的评论是正确的——如果你固定内存，使用所有 API 的流和异步版本并将循环重写为呼吸优先而不是深度优先，你可以使用一个线程执行此操作，并且在所有设备上仍然存在重叠
@talonmies 如果每个设备的块号不同怎么办？
您在问题中说您假设块数是相同的。我的发言正是基于此。如果不是相同的宽度，首先仍然可以工作。想象一下，预先计算工作拆分并将其推送到每个设备的工作堆栈上。代替固定的跳闸循环，有一个堆栈弹出。如果一个设备在另一个之前达到空，则空设备将被跳过。当你到达一个设备时，你又回到了上面的深度优先操作
@talonmies 你能给我看一个演示吗？更重要的是，如果我使用广度优先的方式来做到这一点，那么对于最后一个设备，它必须等到所有以前的异步 API 完成。对吗？
对不起，没有。我不是写你是框架原型来解释我认为我在上面的 cmets 中已经充分描述的东西

标签： multithreading cuda gpu

【解决方案1】：

我可以使用多线程在每个设备上启动数据传输和内核计算吗？

是的。运行时 API 是线程安全的，如果设计正确（甚至可能是 OpenMP 或其他主机多线程编译器驱动的抽象），这当然可以通过每个 GPU 的主机线程来完成。

但是，正如 cmets 中所讨论的，如果您使用流和异步 API，则完全有可能编写一个广度优先版本，该版本可以从单个控制主机线程在多个设备上同时运行。

【讨论】：