Can I control synchronization?

[rgov]

I have a host function which performs several device-based thrust::transform calls in succession. I ran a profile of my code and was surprised that it spends half its time in synchronization. Can I get it to synchronize only after the last transformation?

[rgov]

My solution so far is to patch synchronize_stream() in thrust/system/cuda/detail/par.h to remove the call to cudaDeviceSynchronize() and then call it on my own. This is definitely fraught with peril.

[mjhcamp]

I ran into the same issue with thrust::fill. Here is how I solved it:

struct async_device_execution_policy : thrust::device_execution_policy<async_device_execution_policy> {};

__host__ __device__
cudaError_t synchronize_stream(async_device_execution_policy&)
{
	return cudaSuccess;
}

And then use the following lines to launch the kernel:

async_device_execution_policy policy;
thrust::fill(policy, ptr, ptr + size, false);

Or in your case use the thrust::transform overload that requires an execution policy parameter.

[alliepiper]

#1568 adds a thrust::cuda::par_nosync execution policy that skips all optional synchronizations.