Adding legacy is_hetero trait and supporting it more fully

[danhoeflinger]

Re-add legacy trait is_hetero, but supporting it more fully by adding an additional requirement for the is_hetero type to provide get_idx() which returns the buffer offset this iterator has to its base buffer.

Adding this function to sycl_iterator, incorporating it into __get_sycl_range implementation.
This allows us to support hetero types other than sycl_iterator if they support get_buffer() and get_idx().

[MikeDvorskiy]

Do I understand correctly that "soclomatic" uses such API?

[danhoeflinger]

SYCLomatic does not use sycl_iterator at all, but rather creates a type similar to sycl_iterator (device_pointer / device_iterator) which are similar. They already provide get_idx() within those types.

The idea is to require these is_hetero types to provide get_idx() and then use it within __get_sycl_range processing.

[MikeDvorskiy]

Don't quite understand. __get_sycl_range works fine without get_idx().

[MikeDvorskiy]

... Actually, the answer the next comment explain that. I see..

[MikeDvorskiy]

I think that new 4 lines are equivalent to the couple of line that we had before... No?

__process_input_iter<_LocalAccMode>(__first.base(), oneapi::dpl::end(__first.base().get_buffer()));

[danhoeflinger]

It is, but only if first.base() is a sycl_iterator. If first.base() is instead a dpct::device_pointer this fails because no overload exists with such combination of types for __process_input_iter (dpct::device_pointer,sycl_iterator).

[MikeDvorskiy]

I see..

[MikeDvorskiy]

I think it is equivalent to code that we had before.

const auto __offset = __first - oneapi::dpl::begin(__first.get_buffer());

[danhoeflinger]

Similar response to above, it is the same as long as __first is a sycl_iterator. If it is instead a dpct::device_pointer then this subtraction is undefined.

[MikeDvorskiy]

I see.
It looks we need to extend the contract for sycl_iterator type

[MikeDvorskiy]

Due to the necessity of these changes is not so obviously, it makes sense to add comments why we need some changes, otherwise, some new developer may re-write (revert )that code as redundant.

[adamfidel]

From what I understand, the issue is that we cannot easily get the index of the iterator because we are assuming that the base iterator type is a sycl_iterator. I believe we can solve this by instead of using oneapl::dpl::begin, explicitly constructing a begin iterator of the base iterator type directly from the buffer.

Suggested change

	auto res_src = __process_input_iter<_LocalAccMode>(oneapi::dpl::begin(__base_buffer) + __base_iter.get_idx(),
	using _BaseIterator = ::std::decay_t<decltype(__base_iter)>;
	auto __base_begin = _BaseIterator(__base_buffer);
	auto __idx = __first.base() - __base_begin;
	auto res_src = __process_input_iter<_LocalAccMode>(__base_begin + __idx,

In this way, we do not have to introduce a get_idx method to the sycl_iterator class.

One drawback to using the above proposed approach is that we rely on the base iterator to define a constructor that takes just a sycl::buffer. From what I can tell, this is true for both sycl_iterator and dpct::device_ptr.

[danhoeflinger]

If I understand your suggestion correctly, I think the problem here is that we would lose the information of the buffer offset if we use oneapi::dpl::begin.

[danhoeflinger]

Closing this in favor of #1383, I will add the requested comments there.