如何通过拆分数据，利用多线程多GPU高效计算大型矩阵向量乘法？

Question

我是一个刚接触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). 这里我假设所有设备的块数是相同的。我的问题如下

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

Answer 1

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

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

然而，正如评论中所讨论的那样，如果你使用流和异步API，那么完全可以写一个宽度第一的版本，它将从一个控制主机线程在多个设备上并发运行。