我需要真正了解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);
在这里我的问题:
__host__
的bc)。在doc上他们说cudaEventRecord在事件中捕获了蒸汽的内容。我们称之为流的“内容”并不完全清楚。
但后来我不知道它是如何工作的。实际上,如果MatProd需要很长时间,CPU将在GPU完成其工作之前到达第二个cudaEventRecord。所以我应该得到一个错误的结果......?
我这样说是因为你解释我这些API函数是在主机上执行的。所以根据我的理解,它们将与内核并行启动。当我们在两个cudaEventRecord()之后同步时,我应该得到一个错误的结果......?
也许是因为我没有真正明白你在主机上执行的意思,但我把它理解为在CPU上启动的功能(因此,它不需要等待内核完成)。
在设备上执行的唯一事情是代码前面是__global__
或__device__
。其他所有内容,包括CUDA运行时API调用,以及实际的内核启动本身,都是主机代码。
由于尚未发生(第二)事件,因此您将获得零。
请阅读the documentation获取cudaEventElapsedTime
:
如果
cudaEventRecord()
已被调用两个事件,但其中一个或两个尚未完成(即,cudaEventQuery()
将至少在其中一个事件中返回cudaErrorNotReady
),则返回cudaErrorNotReady
。
这是你的情况下发生的事情,因为你没有做proper CUDA error checking你是盲目的。当两个事件都没有完成时(这意味着CUDA执行流没有到达两个事件),那么除了返回CUDA错误之外,cudaEventElapsedTime()
调用不执行任何操作。
如果在cudaDeviceSynchronize()
调用之前放置cudaEventSynchronize()
调用或适当的cudaEventElapsedTime()
调用,这将强制CPU线程在该点等待,直到事件完成。这将满足cudaEventElapsedTime()
呼叫的必要条件,并且您将获得经过时间的合理值。
添加进一步说明。让我们一步一步地考虑这个问题。
startComputation
事件“记录”到CUDA执行流中:cudaEventRecord(startComputation);
CUDA处理器(GPU)处于空闲状态。因此,在这一刻,特定的CUDA事件startComputation
被认为是“已记录”但不是“已完成”cudaEventRecord
调用之后的下一个项目,即内核启动:MatProd << <BlockPerGrid, ThreadPerBlock >> >(...)
。在此期间,CPU将内核启动作为要在CUDA执行流中处理的下一个项目。由于上面时间段1中的活动,CUDA处理器(GPU)有工作要做,所以它开始处理事件。该事件的处理将事件从“已记录”状态转换为“已完成”状态。cudaEventRecord(stopComputation);
就像在时间段1中一样,这将事件放入CUDA执行流中,以便在之后处理内核执行完成。因此,此新事件处于“已记录”状态,但不处于“已完成”状态。在此时间段3期间,GPU开始执行内核并且正忙于执行内核。cudaEventElapsedTime
)之间进行测量。为了进行此测量,两个事件必须处于“已完成”状态。在此时间段4期间,GPU仍在忙着处理内核,因此它没有前进来处理在时间段3中“已记录”但未“完成”的stopComputation
事件。因此,这两个事件中的第一个(startComputation
) )处于“已完成”状态,但两个事件中的第二个(stopComputation
)仍处于“已记录”状态。因此,cuda运行时API调用(如已指示)将返回错误,并且不会给出合理的度量。它要求两个事件在返回所请求的测量之前处于“已完成”状态。那么,在经过修改的代码中有什么不同,并且在经过时间请求之前包含同步函数?让我们在上面的时间段3结束后接受我们的时间线重播,因为到目前为止的一切都没有改变。但现在时间段4不同了:
cudaDeviceSynchronize()
)。在此时间段4期间,GPU仍在忙于处理内核。由于CUDA时间轴/流仍有工作要做,因此CPU线程在同步步骤中停止。它坐在那里等待。cudaDeviceSynchronize()
调用等待。cudaDeviceSynchronize()
调用等待。stopComputation
。此事件的处理将stopComputation
的状态从“RECORDED”转换为“COMPLETED”。由于GPU在时间段7期间仍在执行某些操作,因此CPU线程在cudaDeviceSynchronize()
调用处等待。cudaDeviceSynchronize()
调用时等待,因此它将继续执行CPU线程中的下一个项目,即经过时间测量的请求。作为先前活动的结果,两个事件(startComputation
和stopComputation
)都处于“COMPLETED”状态,因此事件经过时间测量请求是合法的,并且该调用将返回合理的测量(并且没有错误)。