在 CUDA 中重叠内核执行和数据传输的最佳数据大小

Question

我有这个 cuda 内核，它可以对方阵的元素进行平方，它非常有趣。我想使用 3 个 cuda 流并将输入矩阵划分为多个块，以便我以循环方式使用流在给定块上执行 H2D MemcpyAsync、内核启动和 D2H MemcpyAsync。这是完整的源代码。

#include<iostream>
#include<vector>
#include<cuda.h>
#include<sys/time.h>
using namespace std;
__global__ void MatrixSquareKernel(int *inMatrix, int *outMatrix, size_t width, size_t     rowCount) {
    int myId = blockIdx.x * blockDim.x + threadIdx.x;
    size_t crntRow = 0;
    if(myId < width) {
            size_t mId;
            while(crntRow < rowCount) {
                    mId = myId * width + crntRow;enter code here
                    outMatrix[mId] = inMatrix[mId] * inMatrix[mId];
                    crntRow++;
            }
    }
 }
 int main() {

    size_t count = width * width;
    size_t size = count * sizeof(int);

    vector<cudaStream_t> streams(strCount);

    for(int i = 0; i < strCount; i++)
            cudaStreamCreate(&streams[i]);

    int *h_inMatrix, *h_outMatrix;
    int *d_inMatrix, *d_outMatrix;

    cudaHostAlloc((void **)&h_inMatrix, size, cudaHostAllocDefault);
    cudaHostAlloc((void **)&h_outMatrix, size, cudaHostAllocDefault);
    cudaMalloc((void **)&d_inMatrix, size);
    cudaMalloc((void **)&d_outMatrix, size);

    for(int i = 0; i = count; i++)
            h_inMatrix[i] = i;

    size_t optimalRows = 16;
    size_t iter = width/optimalRows + ((width % optimalRows == 0)? 0: 1);

    size_t chnkOffset, chnkSize, strId, sentRows;

    struct timeval start, stop;
    gettimeofday(&start, NULL);
    for(int i = 0; i < iter; i++){
            sentRows = i * optimalRows;
            chnkOffset = width * sentRows;
            chnkSize = width * optimalRows * sizeof(int);
            if(sentRows > width){
                    optimalRows -= sentRows - width; //Cutoff the extra rows in this chunk if it's larger than the remaining unsent rows
                    chnkSize = width * optimalRows * sizeof(int);
            }
            strId = i % strCount;
            cudaMemcpyAsync(d_inMatrix + chnkOffset, h_inMatrix + chnkOffset, chnkSize, cudaMemcpyHostToDevice, streams.at(strId));
            MatrixSquareKernel<<<1, width, 0, streams.at(strId)>>>(d_inMatrix + chnkOffset, d_outMatrix + chnkOffset, width, optimalRows);
            cudaMemcpyAsync(h_outMatrix + chnkOffset, d_outMatrix + chnkOffset, chnkSize, cudaMemcpyDeviceToHost, streams.at(strId));
    }
    cudaThreadSynchronize();
    gettimeofday(&stop, NULL);

    double elapsedTime = (stop.tv_sec - start.tv_sec) + (start.tv_usec - stop.tv_usec)/1e6;

    cout<<"Elapsed Time: "<<elapsedTime<<endl;


    for(int i = 0; i < strCount; i++)
            cudaStreamDestroy(streams[i]);

    cudaFreeHost(h_inMatrix);
    cudaFreeHost(h_outMatrix);
    cudaFree(d_inMatrix);
    cudaFree(d_outMatrix);

    return 0;
}

每个块都包含一定数量的行，因此变量optimalRows。现在，我给它分配了一个静态值。但我的目标是使用内核在一行矩阵上的完成时间和一行矩阵的传输时间来计算它的值。假设这个值为n。为了计算它，我正在为n求解方程T_tr(n * width * sizeof(int)) = n * T_k + T_k-overhead，其中T_tr(M)是M字节数据的传输时间，我可以通过考虑PCI/e总线的带宽来计算，@ 987654328@是矩阵单行平方的完成时间，T_k-overhead是内核启动的成本。为了测量T_k 和T_k-overhead 的值，我启动了两次内核，一次是仅对取T_k1 时间单位的矩阵的一行进行平方，另一种是将取为时间单位的矩阵的两行平方T_k2 时间单位。取差值将是每行矩阵的内核完成时间；因此， T_k = T_k2 - T_k1 和 T_k-overhead = 2*T_k1 - T_k2。我认为在给定这些参数的情况下为n 求解上述方程会给我一个大于1 的n 值，但它给我一个小于1 的值。

我错过了什么？我真的很感激你的想法。 谢谢

Answer 1

我想在给定这些参数的情况下为 n 求解上述方程 会给我一个大于 1 的 n 值，但它给了我 小于 1 的值。

您不是在最小化T_tr，您只是在寻找满足涉及T_tr 的条件的n。

小于 1 的值是有意义的。零值是一个明显的解决方案，并为您提供

T_tr(0) = T_k-overhead // always true

如果T_k-overhead = 2*T_k1 - T_k2 也是正确的，如果

N * size(T_k) == N * T_tr(T_k) // considering the problem perfectly linear

由于您的问题是线性的，因此当您的 GPU 利用率最高时条件为真。

这实际上是你应该首先做的：

1. 最大化 GPU 利用率
2. 重叠传输和执行

为了最大限度地提高利用率，您需要增加n，直到执行线性增加。您还需要通过改进内存模式来优化内核：

不是每个内核线程处理rowCount 并在内存访问方面取得进展，您应该将每个线程的单个矩阵元素与每个warp 的连续内存访问进行平方。 这也将简化内核，这通常也会增加 gpu 的使用（例如，每个 warp 使用更少的寄存器）

对于重叠执行和传输，您已经知道如何使用异步调用 + 流