CU-LAUN0002 · default_stream_kernel_launch

Detect one kernel launch that names no stream and therefore takes the default one.

This is a deterministic rule for cuda. Read its implementation.

Definition

Read the execution configuration a launch states and report one whose fourth argument is absent or is the null stream, in a translation unit that takes part in stream work at all. A launch with no stream runs on the legacy default stream, which synchronizes implicitly against every other stream on the device, so one such launch in the middle of an overlapped pipeline drains the overlap the rest of the code was built to get. The cost does not show up in the launch itself. It shows up as the copies and kernels around it losing their concurrency, which is why reading it here is worth more than profiling for it later.

Evidence

Each finding names the kernel, the grid and block it launches with, and the function that launches it. The Boolean result reports whether this launch takes the default stream where another stream exists to be drained.

Exceptions

A translation unit that neither creates a stream nor is handed one has nothing to serialize against, so its launches are not reported. Being handed a stream counts as much as creating one, since a function that receives a cudaStream_t and launches without it drains exactly the overlap its caller set up. Setup or teardown outside the hot path costs nothing extra either. A project compiled with per-thread default streams gets a different default whose behavior is the opposite, so it turns this rule off rather than threading a stream it does not need.

Examples

Bad

cudaStreamCreate(&stream);
scale<<<grid, block>>>(data);

Good

cudaStreamCreate(&stream);
scale<<<grid, block, 0, stream>>>(data);

A file that names no stream anywhere returns `False` for `scale<<<grid, block>>>(data)`,
because nothing it can reach was overlapping in the first place.

References

• Cites “CUDA C++ Programming Guide”, streams and the default stream. Open reference
• Cites “The NVIDIA Technical Blog”, GPU pro tip, CUDA 7 streams simplify concurrency. Open reference
• Cites “CUDA C++ Best Practices Guide”, concurrent kernel execution. Open reference