Remove the use of CUDA API wrappers

[makortel]

This issue is to track progress for removing the use of CUDA API wrappers. The library turned out to not be that useful (see some discussion in #279 (comment))).

We are currently using the following components (extracted with git grep, may be incomplete)

• cuda::device_t::create_stream() done partly in Replace use of API wrapper stream and event with plain CUDA, part 1 #389, completed in Replace use of API wrapper stream and event with plain CUDA, part 2 #404
• cuda::stream_t<> done partly in Replace use of API wrapper stream and event with plain CUDA, part 1 #389, completed in Replace use of API wrapper stream and event with plain CUDA, part 2 #404 (except for HeterogeneousEDProducer)
• cuda::stream::wrap() done in Replace use of API wrapper stream and event with plain CUDA, part 1 #389
• cuda::stream_t<>::enqueue::callback() done in Replace use of API wrapper stream and event with plain CUDA, part 1 #389 (except for HeterogeneousEDProducer)
• cuda::event_t done partly in Replace use of API wrapper stream and event with plain CUDA, part 1 #389, completed in Replace use of API wrapper stream and event with plain CUDA, part 2 #404
• cuda::event::create() done partly in Replace use of API wrapper stream and event with plain CUDA, part 1 #389, completed in Replace use of API wrapper stream and event with plain CUDA, part 2 #404
• cuda::memory::copy() done in Replace CUDA API wrapper memory operations with native CUDA calls #395
• cuda::memory::async::copy() done in Replace CUDA API wrapper memory operations with native CUDA calls #395
• cuda::memory::device::zero() done in Replace CUDA API wrapper memory operations with native CUDA calls #395
• cuda::memory::device::async::set() done in Replace CUDA API wrapper memory operations with native CUDA calls #395
• cuda::memory::device::make_unique() done in Replace use of CUDA API wrapper unique_ptrs with CUDAUtilities unique_ptrs #396
• cuda::memory::host::make_unique() done in Replace use of CUDA API wrapper unique_ptrs with CUDAUtilities unique_ptrs #396
• cuda::device::count() done in Replace remaining cuda::device operations with native CUDA calls. #408
• cuda::device::get() done in Replace remaining cuda::device operations with native CUDA calls. #408
• cuda::device::set() done in Replace remaining cuda::device operations with native CUDA calls. #408
• cuda::device::current::get() done in Replace remaining cuda::device operations with native CUDA calls. #408
• cuda::device::current::scoped_override_t<> done in Replace use of cuda::device::current::scoped_override_t with plain CUDA #407
• cuda::launch() done in Implement library-only wrappers for launching CUDA kernels #390
• cuda::throw_if_error() done in Replace cuda::throw_if_error() with cudaCheck() #414
• cuda::is_success() done in Replace use of API wrapper stream and event with plain CUDA, part 1 #389

[makortel]

@fwyzard About the cuda::throw_if_error(), given that we are mostly using cudaCheck(), how about changing cudaCheck() such that there would be a compile-time flag whether to abort() or throw an exception? (such that in CMSSW master it would throw an exception, patatrack and private forks could still use the abort())

[makortel]

I'll start with removing the streams and events.

[fwyzard]

OK.
I'll have a look at providing an equivalent of cuda::launch().

[fwyzard]

cuda::launch() does not support functions objects or lambdas - just plain functions, right ?

[makortel]

cuda::launch() does not support functions objects or lambdas - just plain functions, right ?

I believe it is just a wrapper for the kernel launch <<<...>>>, so the passed function has to be __global__ function (or whatever CUDA kernel launch supports).

[makortel]

First part of streams and events is done in #389.

[fwyzard]

@makortel I have prepared an alternative for cuda::launch().

Do you prefer to keep the same syntax, e.g.

launch(kernel, {gridDim, blockDim, sharedMem = 0, stream = nullptr}, args...);

or something more like what Cupla uses, e.g.

launch(kernel)(gridDim, blockDim, sharedMem = 0, stream = nullptr)(args...);

?

[makortel]

@fwyzard I don't have a clear preference (I would have probably gone with the bare kernel launch syntax, i.e. <<<...>>>, but am really fine with ~any syntax since the use is not that widespread).

[fwyzard]

OK, I'll stick to the current syntax then, since it is simpler to implement.

The advantage over the kernel<<<...>>>(...) syntax is that launch(kernel, {...}, ...) can be used also in the host compiler.

[fwyzard]

I have asked @waredjeb to look into the cuda::memory operations, replacing

• cuda::memory::device::make_unique() with cudautils::make_device_unique()
• cuda::memory[::async]::copy() with cudaMemcpy[Async]()
• cuda::memory[::async]::zero() and ::set() with cudaMemset[Async]()

[makortel]

• cuda::memory::device::make_unique() with cudautils::make_device_unique()

The cudautils::make_device_unique() is not a direct replacement as it (currently) requires the CUDA stream (although I'm thinking to add a variant that caches the allocation but is not "tied" to a CUDA stream).

IIRC all calls to cuda::memory::device::make_unique() are in unit tests. I was thinking that maybe the API wrappers could be acceptable there. On the other hand, using the CUDAStreamCache to get the stream for the allocator is not too hard either.

[fwyzard]

Would it work to explicitly pass stream 0 (actually nullptr) to use the default stream ?

[makortel]

Would it work to explicitly pass stream 0 (actually nullptr) to use the default stream ?

Probably, since all *Async API calls accept that (right?).

[fwyzard]

I think they do.

[fwyzard]

@fwyzard About the cuda::throw_if_error(), given that we are mostly using cudaCheck(), how about changing cudaCheck() such that there would be a compile-time flag whether to abort() or throw an exception? (such that in CMSSW master it would throw an exception, patatrack and private forks could still use the abort())

Sorry @makortel looks like I never answered you: given that an uncaught exception results anyway in an abort() can we just make the replacement for good ?

[makortel]

@fwyzard

given that an uncaught exception results anyway in an abort() can we just make the replacement for good ?

I didn't follow if you meant "cudaCheck() should throw exception" or "cudaCheck() should abort"?

(also, exceptions from places where we call CUDA APIs should get caught by the framework)

[fwyzard]

I meant, cudaCheck() should just throw exceptions.

exceptions from places where we call CUDA APIs should get caught by the framework

Right... but they would still result in a stack trace and in the job ending, shouldn't they ?

Edit: see #398 .

[makortel]

#404 completes the work for streams and events.

[makortel]

I can take care of cuda::device::current::scoped_override_t<> after #404 gets merged.

[waredjeb]

I can work on the remaining cuda::devicecalls!

[fwyzard]

The last references to the CUDA API Wrappers were removed via #417 .