DeviceScan::ExclusiveSum blocks execution thread about cub HOT 4 CLOSED

Example
https://github.com/sh1ng/thrust_async/blob/master/example_device_scan.cu

that code compiled with

nvcc -gencode arch=compute_61,code=sm_61 --default-stream=per-thread example_device_scan.cu -o a -I../cub/ -O3

from cub.

There are no blocking API calls in device-scan. How is your Visual Profiler indicating that it blocks the CPU thread? (Be sure to check the "profile execution on CPU" checkbox when configuring the session's profiling options.)

When I compile that program using your string, run it through the Visual profiler, I see the following timeline, which shows kernel execution overlapping with the CPU's subsequent call to cudaStreamSynchronize() at the end of the program.

Let me know if I'm missing anything,

Duane

from cub.

I've seen the similar picture. But why kernels don't overlap(or overlap just a little)?

from cub.

They do overlap eachother by a little in this pic, and certainly they overlap the host.

TBH, overlapping is bad in this case. Each scan is large enough to fill the GPU by itself -- if the scheduler overlapped them more, then overall performance would go down because they would be contending for hardware resources (like L2 cache thrashing).

from cub.