Scrub temporary_array for concrete spaces

[GoogleCodeExporter]

temporary_array's tag type should be inferred from the function's iterator 
argument.  We should typically not use a concrete backend tag

Original issue reported on code.google.com by jaredhoberock on 7 Nov 2011 at 10:59

[GoogleCodeExporter]

Here's a general recipe for doing this sort of thing until we have c++11 
decltype:

#include <thrust/system/detail/generic/select_system.h>
#include <thrust/iterator/iterator_traits.h>

namespace some_system
{
namespace detail
{
namespace algo_detail
{

// the algorithm implementation takes a templated Tag to allow us to get
// the specific type of the tag used during ADL dispatch
template<typename Tag, typename Iterator>
void algo(Tag, Iterator first, Iterator last)
{
  temporary_array<int, Tag> scratch_space;

  ...
} // end algo

} // end algo_detail

// the algorithm entry point takes a specific system tag
// to recover the original user tag, dispatch again using select_system
template<typename Iterator> void algo(system::tag, Iterator first, Iterator 
last)
{
  using thrust::system::detail::generic::select_system;

  typedef typename thrust::iterator_space<Iterator>::type space;

  return algo_detail::algo(select_system(space()), first, last));
} // end algo

} // end detail
} // end some_system

When we have c++11 decltype, we can do this:

#include <thrust/system/detail/generic/select_system.h>
#include <thrust/iterator/iterator_traits.h>

namespace some_system
{
namespace detail
{

template<typename Iterator> void algo(system::tag, Iterator first, Iterator 
last)
{
  using thrust::system::detail::generic::select_system;
  typedef typename thrust::iterator_space<Iterator>::type space;

  temporary_array<int, decltype(select_system(space()))> scratch_space;

  ...
}

}
}

Or use some equivalent shorthand.

Original comment by jaredhoberock on 4 Feb 2012 at 1:14

[GoogleCodeExporter]

cuda::detail::copy_device_to_device uses temporary_array<..., cpp::tag> to 
explicitly allocate scratch on the host.  We want to keep this.

cuda::detail::launch_closure uses temporary_array<..., cuda::tag> to marshal 
kernel parameters when they get big. It's not clear how to get the original 
entry point's iterator tags into launch_closure, because it tends to get called 
by non-entry points. Punt on this one for now.

Original comment by jaredhoberock on 6 Feb 2012 at 11:12

[GoogleCodeExporter]

This issue was closed by revision b9f9a4a80256.

Original comment by jaredhoberock on 6 Feb 2012 at 11:57

• Changed state: Fixed