cudaEvent_t start, 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]);
MyKernel<<<100, 512, 0, stream[i]>>>(outputDev + i * size, inputDev + i * size, size);
cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventRecord的函数原型是 cudaError_t cudaEventRecord (cudaEvent_t event, cudaStream_t stream)
记录一个事件。如果stream 是非零的,当流中所有的操作完毕,事件被记录;否则,当CUDA context 中所有的操作完毕,事件被记录。
注意对事件的同步操作是在记录事件之后,顺序不能颠倒。
在本例中,cudaEventSynchronize(stop);这句可以用 cudaStreamSynchronize(0); 或者 cudaThreadSynchronize();来代替,能达到同样的效果。
代码来自官方文档。