测量内核在使用流时所花费的总时间答案

【问题标题】：Measuring the total time taken by the kernel when using streams测量内核在使用流时所花费的总时间
【发布时间】：2014-03-29 17:30:11
【问题描述】：

我希望分析在内核上花费的总时间，运行多次，并且想知道这段代码是否会给我在流式内核上的总花费，或者返回的时间是否需要乘以启动次数.

cudaEvent_t start, stop;    
cudaEventCreate(&start);
cudaEventCreate(&stop);


for(x=0; x<SIZE; x+=N*2){

     gpuErrchk(cudaMemcpyAsync(data_d0, data_h+x, N*sizeof(char), cudaMemcpyHostToDevice, stream0));
     gpuErrchk(cudaMemcpyAsync(data_d1, data_h+x+N, N*sizeof(char), cudaMemcpyHostToDevice, stream1));


     gpuErrchk(cudaMemcpyAsync(array_d0, array_h, wrap->size*sizeof(node_r), cudaMemcpyHostToDevice, stream0));
     gpuErrchk(cudaMemcpyAsync(array_d1, array_h, wrap->size*sizeof(node_r), cudaMemcpyHostToDevice, stream1));

     cudaEventRecord(start, 0);
        GPU<<<N/512,512,0,stream0>>>(array_d0, data_d0, out_d0 );
        GPU<<<N/512,512,0,stream1>>>(array_d1, data_d1, out_d1);
     cudaEventRecord(stop, 0);

     gpuErrchk(cudaMemcpyAsync(out_h+x, out_d0 , N * sizeof(int), cudaMemcpyDeviceToHost, stream0));
     gpuErrchk(cudaMemcpyAsync(out_h+x+N, out_d1 ,N *  sizeof(int), cudaMemcpyDeviceToHost, stream1));

} 

float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
printf("Time %f ms\n", elapsedTime);

【问题讨论】：

标签： c cuda

【解决方案1】：

在 CUDA 中添加了基于事件的时序，以实现片上执行的细粒度时序（例如，即使只有一个内核调用被事件启动/停止调用括起来，您也应该获得准确的时间）。但是流和乱序执行给cudaEventRecord()记录的“时间戳”的含义带来了歧义。 cudaEventRecord() 接受一个流参数，据我所知，它尊重该流参数；但是流的执行可能会受到其他流的影响，例如如果他们正在争夺某些资源。

因此最好在 NULL 流上调用 cudaEventRecord() 进行序列化。

有趣的是，英特尔与 RDTSC 有着相似的历史，他们在同一产品中引入了超标量执行和时间戳记录。（对于 NVIDIA，它是 CUDA 1.1；对于 Intel，它是 Pentium。）同样，Intel 必须修改他们对依赖 RDTSC 作为序列化指令的开发人员的指导，告诉他们显式序列化以获得有意义的时序结果。

Why isn't RDTSC a serializing instruction?

【讨论】：

【解决方案2】：

它不会捕获所有循环过程中内核的总执行时间。

来自documentation：

如果之前在事件中调用了 cudaEventRecord()，则此调用将覆盖事件中的任何现有状态。任何检查事件状态的后续调用将只检查最近一次调用 cudaEventRecord() 的完成情况。

如果您认为每次循环的执行时间大致相同，那么您可以将结果乘以循环次数。

请注意，在调用cudaEventElapsedTime() 之前，您应该在stop 事件上发出cudaEventSynchronize() 调用

【讨论】：