2

我有这个 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。为了计算它，我正在求解方程T_tr(n * width * sizeof(int)) = n * T_k + T_k-overhead，n其中T_tr(M)是数据字节的传输时间M，我可以通过考虑 PCI/e 总线的带宽来计算，T_k是对矩阵的单行求平方的完成时间，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会给我一个n大于的值，1但它给我的值小于1.

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

4

1 回答 1

1

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

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

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

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

也是T_k-overhead = 2*T_k1 - T_k2正确的iff

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

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

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

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

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

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

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

于 2013-08-24T22:59:46.900 回答