Refactoring of internal storage types (continued)

[akukanov]

This PR continues improving the internal SYCL backend infrastructure to deal with temporary buffers and device-computed results. It complements the storage types (__device/result/combined/_storage) with a "storage holder" to keep the allocated memory alive and pass it from a backend pattern to its caller.

Changes made:

• create a new utility header specially for storage types and related helpers, move respective types and functions there.
  • __result_and_scratch_storage and __future are not moved and will mostly remain intact until fully replaced.
• add a new __storage_holder class tempate to specify a set of storage buffers that a backend pattern may deal with.
  • scratch buffers are kept type-erased, only to be released at a later point
  • result buffers can be of several types that are defined by a parameter pack.
• add an implementation-detail test for the new template.
• use this new infrastructure to replace __result_and_scratch_storage in merge sort (__parallel_sort_impl)
• (unrelated but rather small) generalize utilities for optional result storage used in reduce-then-scan

How it is supposed to work:

• a backend pattern defines an alias to __storage_holder to specialize its return types and limit the number of scratch buffers.
• callers use that alias to create an instance of storage holder and pass it to the pattern by reference.
• the pattern creates and uses storage types as needed, and at the end deposits them to the holder via its __take method.
• a sycl::event that signals kernel completion remains (a part of) the returned value.
• the caller is responsible to ensure sufficient lifetime for the holder and its content
• the returned values can be extracted from the holder via its __copy_result method.
• the content of the holder can be extracted into a tuple for wrapping into __future.

[danhoeflinger]

If my understanding is correct, and we have device memory / buffer and an offset, this means we are holding alive both a result and a scratch, where the offset determines where we transition from one to the other.
Should we make this more clear here in the comments (and perhaps naming of the struct?).

[akukanov]

Changed my answer after re-reading the comment. The understanding is correct, but the goal of the struct is to hold the memory where the result resides. A part of that memory might have been used as a scratch, but at least now that does not matter anymore, as the memory will be released, not reused.

[danhoeflinger]

If the intent for these is to only be a "sink" whose main purpose is lifetime extension, we may want to consider making them more opaque (protected fields).

I think the one instance where we pull stuff out of this is to produce the final results from within __storage_holder. That storage holder could be friends with __result_keepalive to make more clear the intention here that this is a data sink, meant for lifetime extension and that extracting data from them is the exception.

It really is a preference thing, at the most a defense against future misuse, so perhaps just a small change to the comments would suffice without more complexity.

[akukanov]

Given that these are internal types designed specifically for the device data usage pipeline (allocate storage / use it on the device / transfer results to the host across code layers / keep alive until the kernel completes / extract results / deallocate) and all the same functions/classes that currently operate with these structs would all be declared friends, that sounds like just-in-case overengineering to me. Improving the comments is of course reasonable :)

[danhoeflinger]

Should this be a const member function of __result_keepalive?

Or if we don't want it as a member function, perhaps a static function to keep the implementation together?

[akukanov]

I prefer the keepalive and copyable state to be just simple structs without member functions, not encapsulating any functionality but instead fully exposing their state for any function to operate on.

[danhoeflinger]

I wont insist, but I think I would prefer a name like __raw_result_storage, and __raw_scratch_storage, instead of keepalive which, despite the comment above, may make it confusing in thinking these are RAII storage with ownership.

[akukanov]

Well, it's not a storage either, at least not in the sense of other "storage" types which provide access to data. Maybe __scratch/result_raw_state.

[danhoeflinger]

yeah or __scratch/result_raw_record?

[akukanov]

Renamed and updated comments.