First of all let's remember basic CUDA stream semantics:
- CUDA activity issued into the same stream will always execute in issue order.
- There is no defined relationship between the order of execution of CUDA activities issued into separate streams.
The CUDA default stream (assuming we have not overridden the default legacy behavior) has an additional characteristic of implicit synchronization, which roughly means that a CUDA operation issued into the default stream will not begin executing until all prior issued CUDA activity to that device has completed.
Therefore, if we issue 2 CUDA events (say, start and stop) into the legacy default stream, we can be confident that any and all CUDA activity issued between those two issue points will be timed (regardless of which stream they were issued into, or which host thread they were issued from). I would suggest for casual usage this is intuitive, and less likely to be misinterpreted. Furthermore, it should yield consistent timing behavior, run-to-run (assuming host thread behavior is the same, i.e. somehow synchronized).
OTOH, let's say we have a multi-streamed application. Let's assume that we are issuing kernels into 2 or more non-default streams:
Stream1: cudaEventRecord(start)|Kernel1|Kernel2|cudaEventRecord(stop)
Stream2: |Kernel3|
It does not really matter too much whether these were issued from the same host thread or from separate host threads. For example, let's say our single host thread activity looked like this (condensed):
cudaEventRecord(start, Stream1);
Kernel1<<<..., Stream1>>>(...);
Kernel2<<<..., Stream1>>>(...);
Kernel3<<<..., Stream2>>>(...);
cudaEventRecord(stop, Stream1);
What timing should we expect? Will Kernel3 be included in the elapsed time between start and stop?
In fact the answer is unknown, and could vary from run-to-run, and probably would depend on what else is happening on the device before and during the above activity.
For the above issue order, and assuming we have no other activity on the device, we can assume that immediately after the cudaEventRecord(start) operation, that the Kernel1 will launch and begin executing. Let's suppose it "fills the device" so that no other kernels can execute concurrently. Let's also assume that the duration of Kernel1 is much longer than the launch latency of Kernel2 and Kernel3. Therefore, while Kernel1 is executing, both Kernel2 and Kernel3 are queued for execution. At the completion of Kernel1, the device scheduler has the option of beginning either Kernel2 or Kernel3. If it chooses Kernel2 then at the completion of Kernel2 it can mark the stop event as completed, which will establish the time duration between start and stop as the duration of Kernel1 and Kernel2, approximately.
Device Execution: event(start)|Kernel1|Kernel2|event(stop)|Kernel3|
| Duration |
However, if the scheduler chooses to begin Kernel3 before Kernel2 (an entirely legal and valid choice based on the stream semantics) then the stop event cannot be marked as complete until Kernel2 finishes, which means the measured duration will now included the duration of Kernel1 plus Kernel2 plus Kernel3. There is nothing in the CUDA programming model to sort this out, which means the measured timing could alternate even run-to-run:
Device Execution: event(start)|Kernel1|Kernel3|Kernel2|event(stop)|
| Duration |
Furthermore, we could considerably alter the actual issue order, placing the issue/launch of Kernel3 before the first cudaEventRecord or after the last cudaEventRecord, and the above argument/variability still holds. This is where the meaning of the asynchronous nature of the cudaEventRecord call comes in. It does not block the CPU thread, but like a kernel launch it is asynchronous. Therefore all of the above activity can issue before any of it actually begins to execute on the device. Even if Kernel3 begins executing before the first cudaEventRecord, it will occupy the device for some time, delaying the beginning of execution of Kernel1, and therefore increasing the measured duration by some amount.
And if the Kernel3 is issued even after the last cudaEventRecord, because all these issue operations are asynchronous, the Kernel3 may still be queued up and ready to go when Kernel1 is complete, meaning the device scheduler can still make a choice about which to launch, making for possibly variable timing.
There are certainly other similar hazards that can be mapped out. This sort of possibility for variation in a multi-streamed scenario is what gives rise to the conservative advice to avoid trying to do cudaEvent based timing using events issued into the non-legacy-default stream.
Of course, if you for example use the visual profiler then there should be relatively little ambiguity about what was measured between two events (although it may still vary run-to-run). However, if you're going to use the visual profiler, you can read the duration directly off the timeline view, without needing an event elapsed time call.
Note that if you override the default stream legacy behavior, the default stream roughly becomes equivalent to an "ordinary" stream (especially for a single-threaded host application). In this case, we can't rely on the default stream semantics to sort this out. One possible option might be to precede any cudaEventRecord() call with a cudaDeviceSynchronize() call. I'm not suggesting this sorts out every possible scenario, but for single-device single host-thread applications, it should be equivalent to cudaEvent timing issued into default legacy stream.
Complex scenario timing might be best done using a profiler. Many folks also dispense entirely with cudaEvent based timing and revert to high-resolution host timing methodologies. In any event, the timing of a complex concurrent asynchronous system is non-trivial. The conservative advice intends to avoid some of these issues for casual use.
