generalize get_iterator_value() to respect execution policy

[3gx]

Unless it is a raw pointers, get_value_iterator [1] will not respect execution policy when dereferenced via *it. A naïve replacement of *it with

typename thrust::iterator_traits<Iterator>::value_type value;
thrust::copy(exec, it,it+1, &value); 
return value; 

makes the reproducer in NVIDIA/thrust#780 to die with

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  an illegal memory access was encountered
Aborted (core dumped)

[1] https://github.com/thrust/thrust/blob/master/thrust/detail/get_iterator_value.h#L29

[jaredhoberock]

I think incorporating the cross_system policy into the policy passed to thrust::copy might make your implementation work. The idea would be to use cpp::tag as one of the systems, and the user-provided execution policy as the other system. In this scenario, cpp::tag essentially communicates that &value lives on the stack.

[3gx]

cross_system is cuda specific, wouldn't it then be better to have a generic get_value_iterator implementation which dereferences iterator in an usual way *f, and cuda customisation which constructs cross_system and invoke thrust::copy with it? This will make get_value_iteratorbackend specific, which is not unreasonable.

[jaredhoberock]

We shouldn't make get_iterator_value() backend-specific. We already have thrust::copy(), which is backend-specific. We just need to figure out how to use it correctly for this problem. I'll keep thinking about it.

[3gx]

get_value is backend specific. Why can't same argument that made get_value backend specific be applied to get_iterator_vale ?

[jaredhoberock]

get_value() is an essential customization point to deal with remote memories. It has to exist in order for backends like CUDA to work at all. get_iterator_value() is just an application of thrust::copy().

I think we just need to figure out why the straightforward application of thrust::copy() makes your reproducer crash and solve that problem.

[jaredhoberock]

This implementation of get_iterator_value() makes the reproducer given in NVIDIA/thrust#780's work on my system:

template<typename DerivedPolicy, typename Iterator>
__host__ __device__
typename thrust::iterator_traits<Iterator>::value_type
get_iterator_value(thrust::execution_policy<DerivedPolicy>& exec, Iterator it)
{
  typename thrust::iterator_traits<Iterator>::value_type result;

  // treat &result as a cpp system iterator
  thrust::cpp::tag system2;

  thrust::copy(select_system(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), thrust::detail::derived_cast(thrust::detail::strip_const(system2))), it, it + 1, &result);

  return result;
} // get_iterator_value(exec,Iterator);

The idea is that select_system finds the cross_system thing and tells the CUDA backend where &result lives.

[3gx]

What is the purpose of thrust::detail::strip_const?

[jaredhoberock]

Huh, I also discovered the function detail::two_system_copy_n()

Probably you just want to use that directly instead of thrust::copy().