device_vector using usm device memory and becoming allocator aware

[danhoeflinger]

Update: There are a pair of fundamental issues discovered in the comments below. I think that using USM device memory may not be possible to achieve this functionality in this way. Read the comments for more details.

---

dpct::device_vector currently uses sycl::usm::alloc::shared memory to provide users with the ability to have seamless usage on host and device.

This PR changes the semantics of dpct::device_vector to instead use sycl::usm::alloc::device memory, while preserving the ability to seamlessly use the dpct::device_vector from the host.

This PR also upgrades dpct::device_vector to be close to satisfying AllocatorAwareContainer. However, it diverges from that requirement in the following ways:

1. construct() calls do not use ::std::allocator_traits<allocator_type>::construct(alloc, pointer, args), but rather dpct::device_allocator_traits<allocator_type>::construct(pointer,args) which uses a static construct() call from the supplied allocator, or provides a default if one is not defined.
2. destroy() calls do not use ::std::allocator_traits<allocator_type>::destroy(alloc, pointer), but rather the static functiondpct::device_allocator_traits<allocator_type>::destroy(pointer) which uses a static destroy() call from the supplied allocator, or provides a default if one is not defined.

These changes are done to avoid passing the allocator object into the sycl kernel, as construction and destruction must occur on the device in a kernel. The selection of the call can happen at compile time and must not involve the individual instance of the allocator.

To avoid unnecessary overhead, for dpct::device_vector a kernel is only launched to destroy data if a custom static destroy() function is provided in the allocator. Prior to this PR, no construction or destruction was done, and we do not want to add extra kernel launches to "destroy" data when this will generally result in a no-op in effect.

Changes to the testing can be found here:
danhoeflinger/SYCLomatic-test#1

note: this PR is meant to be a draft on my fork for discussion prior to introducing it to SYCLomatic itself.

[danhoeflinger]

The shift to use usm device memory rather than usm shared memory for device_vector will change the performance of users of this type. However, when SYCLomatic is run, it injects generated headers into the user's repository, rather than relying upon the headers like a library. This means that only if users re-migrate their code will they experience the difference in performance.

It is difficult to say if the performance changes will be a net benefit or detriment because it depends on usage pattern.
The usm shared memory implementation lets the runtime decide where to hold the memory based on the usage. The new implementation always keeps it on the device, and transfers data to and/or from the device each time it is used on the host.

With this in mind, would it be good to provide the legacy shared usm implementation of dpct::device_vector under another name?. If so, do we provide it as it was prior to this PR, or work to add AllocatorAwareContainer requirements to it as well?

[danhoeflinger]

I've realized I'm missing a big piece here, which is proper handling of device_pointer, device_iterator and a solution does not seem easy.

Currently, those merely use value_type & as their reference rather than device_reference<value_type>, and making that switch does not play well with the details of oneDPL. Below is an example of a functor used in fill which assumes a raw reference, and will not function with a device_reference<value_type>:

template <typename _SourceT>
struct fill_functor
{
    _SourceT __value;
    template <typename _TargetT>
    void
    operator()(_TargetT& __target) const
    {
        __target = __value;
    }
};

This also means that the tests that are working are not really using the code inside the __SYCL_DEVICE_ONLY__ macros to access data within kernels, but rather just depending on device_iterator using raw value_type & when dereferenced, as we are passing device_iterator to kernels themselves.

[danhoeflinger]

All of these operators which use both __assign_from() and __get_value() are using two sycl kernels, when this should be able to be accomplished with a single custom kernel.

I have an implementation which attempts to accomplish this on a branch. It uses the preprocessor macro __SYCL_DEVICE_ONLY__ to choose to either submit a custom single_task, or just run simple commands if already on the device.

However, it seems that with this structure based on __SYCL_DEVICE_ONLY__, the compiler does not realize it needs to compile this kernel and it results in a runtime exception, failing to find the custom kernel. My hypothesis is that the kernel submission must be present in an active code path when __SYCL_DEVICE_ONLY__ is defined for the compiler to recognize it. This is a somewhat fundamental issue with this strategy.

[danhoeflinger]

Actually, there is a bug here which is allowing the __assign_from() and __get_value() to work, in the current state of the code. __SYCL_DEVICE_ONLY should be __SYCL_DEVICE_ONLY__, and this is allowing the kernels to be compiled. Fixing that bug breaks the code generally, due to the runtime failure described above.

Using __SYCL_DEVICE_ONLY__ in this fashion does not seem to actually be viable.

[danhoeflinger]

If users really wanted explicitly device USM memory, we could provide explicitly host_accessible_device_vector, which would function seamlessly on the host with device data.
From this container, you could have begin(), end() functions which would provide device accessible iterators. It may also be possible to provide explicitly host_begin() and host_end() to provide host accessible iterators.

However, the current shared USM implementation provides the seamless experience of host and device accessible memory currently, without requiring the programmer to specify where the data will be used. I don't think that is likely to be a winning trade-off.

The allocator aware functionality can be broken off and added separately.