3.2.6.6 事件
通过在应用的任意点上异步地记载事件和查询事件是否真正被记载,运行时提供了精密地监测设备运行进度的方式,事件同时也提供精确计时功能。当事件记载点前面,事件指定的流中的所有任务或者命令全部完成时,事件被记载。只有记载点之前所有的流中的任务/命令都已完成,0号流的事件才会记载。
下面的代码创建了两个事件:
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventCreate(&start);
cudaEventCreate(&stop);
按下面的方式,这些事件可用于计算上节代码的运行时间:
cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>> (outputDev + i * size, inputDev + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
cudaEventRecord(stop, 0);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>> (outputDev + i * size, inputDev + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
cudaEventRecord(stop, 0);
可用这种方式销毁事件
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaEventDestroy(stop);
3.2.6.7 同步调用
直到设备真正完成任务,同步函数调用的控制权才会返回给主机线程。在主机线程执行任何其它CUDA调用前,通过调用cudaSetDeviceFlags()并传入指定标签(参见参考手册)可以指定主机线程的让步,阻塞,或自旋状态。