Busy-wait inside cudaDeviceSynchronize while calling Thrust algorithms

[andrewmcdonald-oxb]

Hello. Not sure if this is the right place to ask a question. But I've inherited an app to profile on x64 Linux, which makes use of Thrust (specifically exclusive_scan_n in this case). The profile shows it spends a lot of CPU time busy-waiting inside cudaDeviceSynchronize.

Is there some way of telling Thrust to sleep the CPU thread instead while waiting for the result from the GPU? I've read the docs I could find but couldn't spot anything about making this choice. Does Thrust just use the default Cuda stream? Perhaps there's a way to set a flag globally on that? I'm familiar with GPU programming but haven't much experience with Cuda yet, so I might be missing something.

[alliepiper]

The wait behavior is controlled by cudaSetDeviceFlags (docs). The cudaDeviceScheduleYield should do what you're looking for.

Thrust uses the default stream by default, but this can be changed by passing thrust::device.on(stream) as the first argument to the algorithm, where stream is a cudaStream_t.

If you want more control over synchronization, you may also want to look at cub::DeviceScan::ExclusiveScan (docs). While Thrust algorithms are inherently synchronous, the CUB algorithms provide the user with more flexibility in this regard.

[andrewmcdonald-oxb]

Thanks Allison, that's helpful! Setting the flag reduces the CPU load for our TensorRT calls as well due to their internals.

I'm opting for the BlockingSync flag at the moment; does the Yield one re-yield if the scheduler picks the thread back up again and the operation still hasn't finished? I guess we're getting away from Thrust now though.

[alliepiper]

I would assume so, but I'm not sure. You might want to ask about that one on the CUDA dev forums.