Breaking change from Thrust 2 (CCCL 1) to Thrust 3 (CCCL 2) needs specialization of pointer_element

[correaa]

When upgrading to CUDA 13.2, I got a breaking change for my library.
After some trial and error I worked around this breaking change by adding an specialization:

#if THRUST_VERSION >= 300200  // CCCL 2
namespace thrust::detail {

template<typename T, ::boost::multi::dimensionality_type D, typename ElementPtr, class Layout, bool IsConst>
struct pointer_element<::boost::multi::detail::subarray_ptr<T, D, ElementPtr, Layout, IsConst>> {
	using type = std::conditional_t<D == 1,
		std::conditional_t<IsConst,
			std::add_const_t<typename pointer_element<ElementPtr>::type>,
			typename pointer_element<ElementPtr>::type>,
		void>;
};

}  // end namespace thrust::detail
#endif

(What is after type = is not very important, it seems that I could have put just void.)

The point is that, this workaround seems fragile, because I am specializing a detail class.
Do you know of a way to achieve the same effect without messing with a detail class.

I don't mind specializations, I have a few here: https://github.com/correaa/boost-multi/blob/develop/include/boost/multi/adaptors/thrust.hpp (including the one mentioned above) and here: https://github.com/correaa/boost-multi/blob/develop/include/boost/multi/adaptors/thrust/fix_pointer_traits.hpp

although I would prefer other mechanisms, like adding extra traits to my own classes (iterators), or some Thrust-blessed mechanism.

For reference these are test passing for CUDA 13.2 https://gitlab.com/correaa/boost-multi/-/jobs/13881966958

[miscco]

I dont see the error message, but have you tried specializing cuda::std::pointer_traits for your contiguous iterator?

The issue with pointer_traits and readable_traits or incrementable_traits is that we cannot sniff out whether the user has specialized them, because the standard does not require any information there from the STL

So we are kind of relying on cuda::std::pointer_traits

[correaa]

yeah, the linked log was to show the fixed state (and hash details, etc).

The breaking change is shown in this earlier log (before the detail::pointer_element hack): https://gitlab.com/correaa/boost-multi/-/jobs/13878030269#L1830

[correaa]

Another (more direct) way in which pointer_element asks to be specialized: https://gitlab.com/correaa/boost-multi/-/jobs/13906779267

(it seems that defining the pointer typedef inside my iterator doesn't help avoiding pointer_element to need a specialization.

[miscco]

We are trying to replace thrust::detail::pointer_traits with cuda::std::pointer_traits but its a bit gnarly because the former offers some nonstandard extensions especially around void pointer

[correaa]

So, that means that in the future (now?) I can specialize cuda::std::pointer_traits instead that this ::detail class.

Please, if you can, confirm, if currently there is no better workaround than specializing this ::detail class.

Having said that, I still prefer Thrust to deduce stuff from the iterator itself. Even if that intrusively means adding non-standard traits to my iterators. I think this it more practical than defining specializations.

Case in point, I still think it is a good idea if the iterators themselves can carry the Thrust system information: #5284

[miscco]

To be clear, it should not be necessary to specialize cuda::std::pointer_traits because that should already pull in the right properties in a standard conforming way.

Thrust was using some hacks around void pointers and the raw_pointer_cast machinery

[correaa]

To be clear, it should not be necessary to specialize cuda::std::pointer_traits because that should already pull in the right properties in a standard conforming way.

From internal (to the iterator) traits? Which one?

As far as I see my iterators are generally quite standard-conforming

I wonder what typedef should I add to my iterator class to make it standard-conforming.

Thrust was using some hacks around void pointers and the raw_pointer_cast machinery

Ok