[FEA]: Expose thrust::detail::contiguous_iterator_raw_pointer_cast

[bernhardmgruber]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

Thrust

Is your feature request related to a problem? Please describe.

I have found thrust::detail::contiguous_iterator_raw_pointer_cast quite useful to unwrap iterators to their underlying pointers. I use this to unwrap a zip iterator and check whether I can use a kernel operating on contiguous memory or not:

    cuda::std::apply([](auto... it) {
            if constexpr((thrust::is_contiguous_iterator<decltype(it)>::value && ...))
                fast_path(thrust::detail::contiguous_iterator_raw_pointer_cast(it)...);
            else
                generic(it...);
        }, zip_iterator.get_iterator_tuple());

The function contiguous_iterator_raw_pointer_cast seems to do exactly what I need, but it is declared in a detail namespace.

Describe the solution you'd like

Move contiguous_iterator_raw_pointer_cast out of the detail namespace and make it part of thrust's public API.

Describe alternatives you've considered

I can always circumvent this by using thrust::raw_pointer_cast(&*it) instead, but it would be nice if I could use the already implemented function.

Additional context

No response

[bernhardmgruber]

CC @ahendriksen

[miscco]

I am not too happy about exposing this. It is a highly dangerous facility that should be used internally.

Is there any reason you cannot use / improve an existing thrust algorithm?

[bernhardmgruber]

I am not too happy about exposing this. It is a highly dangerous facility that should be used internally.

I think it's less dangerous than me calling thrust::raw_pointer_cast(&*it) on an unknown iterator type, potentially forgetting to guard it with a if constexpr (thrust::is_contiguous_iterator_v<decltype(it)>).

Is there any reason you cannot use / improve an existing thrust algorithm?

@ahendriksen and I would like to upstream some improvements at some point, but for now work externally. Independently, other uses may also want to use that functionality to define their own algorithms.

[bernhardmgruber]

@alliepiper mentioned during our VC yesterday that she is fine with exposing the functionality since it's there for a while and seems to have stabilized.

[alliepiper]

It is a highly dangerous facility that should be used internally.

Some more context:

Thrust provides a trait to identify contiguous iterators. It is defined for all contiguous thrust iterators and some common STL vector implementations, etc, and users are allowed to opt-in through some THRUST_PROCLAIM_... macros to declare that their iterator types are contiguous. Doing so allows them to enable both performance and compile-time optimizations by allowing these iterators to be converted to raw pointers.

The utilities exposed by this PR are safest solutions I can think of to solve the problem of converting generic iterators to raw pointers. If an input pointer type is not explicitly marked as contiguous, it is safely and by default returned unmodified. Only known contiguous iterators are decayed to a raw pointer.