CUDA程序不测量执行时间:cudaEventRecord

问题描述 投票:-1回答:1

我需要真正了解CUDA如何测量时间执行。

让我们关注这一部分(如果你想测试它,整个代码就在消息的末尾)。

// Launching Kernel and measuring its time
    cudaEventRecord(startComputation);
    MatProd << <BlockPerGrid, ThreadPerBlock >> >(C, A, B, dimAx, dimBx, dimCx, dimCy);
    cudaEventRecord(stopComputation);

    //cudaEventSynchronize(stopComputation); // this line must be HERE and it returns me a good computation time.

    cudaEventElapsedTime(&millisecondsPureComputation, startComputation, stopComputation);

    cudaDeviceSynchronize(); // putting this doesn't do the job
    //cudaEventSynchronize(stopComputation); // if I put it here instead it doesn't work.

    std::cout << "Computation time : " << millisecondsPureComputation << "ms" << std::endl;

我怎么理解这些事情。程序运行时,CPU多次调用内核。 cudaEventRecord,MatProd和cudaEventElapsedTime都在GPU上执行。

我的两个cudaEventRecord之间的时间是在我的cudaEventElapsedTime中计算的。

问题是:如果我的CPU在GPU计算方面太快,变量millisecondsPureComputaion将保持它的初始值:0。

因此,在显示计算之前,我必须向CPU说“等待GPU完成了cudaEventElapsedTime”。通过这种方式,变量millisecondsPureComputation将具有我们想要的值。

因此,放一个cudaDeviceSynchronise();就在cudaEventElapsedTime之后就足够了。

但实际上它在我这样做时不起作用,变量仍为0.唯一的非零数字的方法是在cudaEventElapsedTime之前放置cudaEvntSynchronize(stopComputation),我不明白为什么。

我的问题:

为什么我的方法放一个cudaDeviceSynchronise();不起作用?你能解释一下为什么要放一个cudaEventSynchronize(stopComputation);在cudaEventElapsedTime工作之前?它有什么特别之处呢?

#include <iostream>
#include <math.h>
#include <chrono>

__global__  void MatProd(float* C, float* A, float*B, int dimAx, int dimBx, int dimCx, int dimCy)
{
    int row = blockDim.y*blockIdx.y + threadIdx.y;
    int col = blockDim.x*blockIdx.x + threadIdx.x;

    double Result = 0;

    if (row <= dimCy - 1 && col <= dimCx - 1)
    {
        for (int k = 0; k < dimAx; k++)
        {
            Result += A[k + dimAx*row] * B[col + dimBx*k];
        }

        C[col + row*dimCx] = Result;
    }
}

int main(void)
{
    /* Initializing the inputs */
    // Matrix sizes
    int dimAx = 100;
    int dimAy = 100;
    int dimBx = 2;

    int dimBy = dimAx;
    int dimCx = dimBx;
    int dimCy = dimAy;

    // Matrix pointers
    float *A, *B, *C;

    // Variable to measure CUDA time execution.
    float millisecondsPureComputation = 0;
    cudaEvent_t startComputation, stopComputation;
    cudaEventCreate(&startComputation);
    cudaEventCreate(&stopComputation);

    // Memory allocation
    cudaMallocManaged(&A, dimAx*dimAy*sizeof(float));
    cudaMallocManaged(&B, dimBx*dimBy*sizeof(float));
    cudaMallocManaged(&C, dimCx*dimCy*sizeof(float));

    // Initializing matrices
    for (int i = 0; i < dimAy; i++)
    {
        for (int j = 0; j < dimAx; j++) 
        {
            A[j + dimAx*i] = j + 10 * i;
        }
    }
    for (int i = 0; i < dimBy; i++)
    {
        for (int j = 0; j < dimBx; j++)
        {
            B[j + dimBx*i] = (j + 1)*pow(i, 2);
        }
    }

    // Kernel properties

    int threadPerBlockx = 32;
    int threadPerBlocky = 32;
    int BlockPerGridx = 1 + (dimCx - 1) / threadPerBlockx;
    int BlockPerGridy = 1 + (dimCy - 1) / threadPerBlockx;


    dim3 BlockPerGrid(BlockPerGridx, BlockPerGridy, 1);
    dim3 ThreadPerBlock(threadPerBlockx, threadPerBlocky, 1);

    // Launching Kernel and measuring its time
    cudaEventRecord(startComputation);
    MatProd << <BlockPerGrid, ThreadPerBlock >> >(C, A, B, dimAx, dimBx, dimCx, dimCy);
    cudaEventRecord(stopComputation);

    //cudaEventSynchronize(stopComputation); // this line must be HERE and it returns me a good computation time.

    cudaEventElapsedTime(&millisecondsPureComputation, startComputation, stopComputation);

    cudaDeviceSynchronize(); // putting this doesn't do the job
    //cudaEventSynchronize(stopComputation); // if I put it here instead it doesn't work.

    std::cout << "Computation time : " << millisecondsPureComputation << "ms" << std::endl;

    cudaFree(A);
    cudaFree(B);
    cudaFree(C);

    return 0;
}

[编辑]我将代码更改为它,现在它可以工作,但我仍然不明白发生了什么..

cudaEventRecord(startComputation);
        MatProd << <BlockPerGrid, ThreadPerBlock >> >(C, A, B, dimAx, dimBx, dimCx, dimCy);

        //cudaDeviceSynchronize();
        cudaEventRecord(stopComputation);
        cudaDeviceSynchronize();
        cudaEventElapsedTime(&millisecondsPureComputation, startComputation, stopComputation);

在这里我的问题:

  • 所以,cudaEventRecord(),cudaEventElapsedTime()在我的情况下在主机上执行,如果我理解得很好(文档中的__host__的bc)。

在doc上他们说cudaEventRecord在事件中捕获了蒸汽的内容。我们称之为流的“内容”并不完全清楚。

但后来我不知道它是如何工作的。实际上,如果MatProd需要很长时间,CPU将在GPU完成其工作之前到达第二个cudaEventRecord。所以我应该得到一个错误的结果......?

我这样说是因为你解释我这些API函数是在主机上执行的。所以根据我的理解,它们将与内核并行启动。当我们在两个cudaEventRecord()之后同步时,我应该得到一个错误的结果......?

也许是因为我没有真正明白你在主机上执行的意思,但我把它理解为在CPU上启动的功能(因此,它不需要等待内核完成)。

c++ cuda
1个回答
1
投票

在设备上执行的唯一事情是代码前面是__global____device__。其他所有内容,包括CUDA运行时API调用,以及实际的内核启动本身,都是主机代码。

由于尚未发生(第二)事件,因此您将获得零。

请阅读the documentation获取cudaEventElapsedTime

如果cudaEventRecord()已被调用两个事件,但其中一个或两个尚未完成(即,cudaEventQuery()将至少在其中一个事件中返回cudaErrorNotReady),则返回cudaErrorNotReady

这是你的情况下发生的事情,因为你没有做proper CUDA error checking你是盲目的。当两个事件都没有完成时(这意味着CUDA执行流没有到达两个事件),那么除了返回CUDA错误之外,cudaEventElapsedTime()调用不执行任何操作。

如果在cudaDeviceSynchronize()调用之前放置cudaEventSynchronize()调用或适当的cudaEventElapsedTime()调用,这将强制CPU线程在该点等待,直到事件完成。这将满足cudaEventElapsedTime()呼叫的必要条件,并且您将获得经过时间的合理值。

添加进一步说明。让我们一步一步地考虑这个问题。

  1. 在时间段1中,由于此调用,CPU代码将startComputation事件“记录”到CUDA执行流中:cudaEventRecord(startComputation); CUDA处理器(GPU)处于空闲状态。因此,在这一刻,特定的CUDA事件startComputation被认为是“已记录”但不是“已完成”
  2. 在时间段2中,CPU线程向前移动到上一个cudaEventRecord调用之后的下一个项目,即内核启动:MatProd << <BlockPerGrid, ThreadPerBlock >> >(...)。在此期间,CPU将内核启动作为要在CUDA执行流中处理的下一个项目。由于上面时间段1中的活动,CUDA处理器(GPU)有工作要做,所以它开始处理事件。该事件的处理将事件从“已记录”状态转换为“已完成”状态。
  3. 在时间段3中,CPU线程向前移动到上一次内核启动之后的下一个项目,这是另一个事件记录调用:cudaEventRecord(stopComputation);就像在时间段1中一样,这将事件放入CUDA执行流中,以便在之后处理内核执行完成。因此,此新事件处于“已记录”状态,但不处于“已完成”状态。在此时间段3期间,GPU开始执行内核并且正忙于执行内核。
  4. 在时间段4中,CPU线程前进到上一个事件记录调用之后的下一个项目,这是对运行时API的请求,以在两个事件(cudaEventElapsedTime)之间进行测量。为了进行此测量,两个事件必须处于“已完成”状态。在此时间段4期间,GPU仍在忙着处理内核,因此它没有前进来处理在时间段3中“已记录”但未“完成”的stopComputation事件。因此,这两个事件中的第一个(startComputation) )处于“已完成”状态,但两个事件中的第二个(stopComputation)仍处于“已记录”状态。因此,cuda运行时API调用(如已指示)将返回错误,并且不会给出合理的度量。它要求两个事件在返回所请求的测量之前处于“已完成”状态。

那么,在经过修改的代码中有什么不同,并且在经过时间请求之前包含同步函数?让我们在上面的时间段3结束后接受我们的时间线重播,因为到目前为止的一切都没有改变。但现在时间段4不同了:

  1. 在时间段4中,CPU线程向前移动以处理CUDA事件记录调用之后的下一个项目,但该指令是同步指令(cudaDeviceSynchronize())。在此时间段4期间,GPU仍在忙于处理内核。由于CUDA时间轴/流仍有工作要做,因此CPU线程在同步步骤中停止。它坐在那里等待。
  2. 在时间段5中,GPU仍然忙于处理内核。 CPU线程卡在cudaDeviceSynchronize()调用等待。
  3. 在时间段6中,GPU仍然忙于处理内核。 CPU线程卡在cudaDeviceSynchronize()调用等待。
  4. 在时间段7中,GPU完成对内核的处理,并继续前进到CUDA流中记录的下一个工作,即cuda事件stopComputation。此事件的处理将stopComputation的状态从“RECORDED”转换为“COMPLETED”。由于GPU在时间段7期间仍在执行某些操作,因此CPU线程在cudaDeviceSynchronize()调用处等待。
  5. 在时间段8中,GPU已经完成了对其发出的所有工作的处理并返回到空闲状态。因此,CPU不再需要在cudaDeviceSynchronize()调用时等待,因此它将继续执行CPU线程中的下一个项目,即经过时间测量的请求。作为先前活动的结果,两个事件(startComputationstopComputation)都处于“COMPLETED”状态,因此事件经过时间测量请求是合法的,并且该调用将返回合理的测量(并且没有错误)。
© www.soinside.com 2019 - 2024. All rights reserved.