Expose publicly get_temporary_buffer and return_temporary_buffer.

Test get_temporary_buffer & malloc

Fixes #160
Fixes #168
Fixes #1
Fixes #2

Author: jaredhoberock

testing/memory.cu

@@ -3,31 +3,37 @@
 #include <thrust/sort.h>
 #include <thrust/memory.h>
 #include <thrust/pair.h>
+#include <thrust/fill.h>
+#include <thrust/logical.h>
+
+
+struct my_system : thrust::device_system<my_system> {};
 
-struct my_tag : thrust::device_system<my_tag> {};
 
 template<typename T1, typename T2>
 bool is_same(const T1 &, const T2 &)
 {
   return false;
 }
 
+
 template<typename T>
 bool is_same(const T &, const T &)
 {
   return true;
 }
 
+
 void TestSelectSystemDifferentTypes()
 {
   using thrust::system::detail::generic::select_system;
 
-  // select_system(my_tag, device_system_tag) should return device_system_tag (the minimum tag)
-  bool is_device_system_tag = is_same(thrust::device_system_tag(), select_system(my_tag(), thrust::device_system_tag()));
+  // select_system(my_system, device_system_tag) should return device_system_tag (the minimum tag)
+  bool is_device_system_tag = is_same(thrust::device_system_tag(), select_system(my_system(), thrust::device_system_tag()));
   ASSERT_EQUAL(true, is_device_system_tag);
 
   // select_system(device_system_tag, my_tag) should return device_system_tag (the minimum tag)
-  is_device_system_tag = is_same(thrust::device_system_tag(), select_system(thrust::device_system_tag(), my_tag()));
+  is_device_system_tag = is_same(thrust::device_system_tag(), select_system(thrust::device_system_tag(), my_system()));
   ASSERT_EQUAL(true, is_device_system_tag);
 }
 DECLARE_UNITTEST(TestSelectSystemDifferentTypes);
@@ -45,42 +51,91 @@ void TestSelectSystemSameTypes()
   bool is_device_system_tag = is_same(thrust::device_system_tag(), select_system(thrust::device_system_tag(), thrust::device_system_tag()));
   ASSERT_EQUAL(true, is_device_system_tag);
 
-  // select_system(my_tag, my_tag) should return my_tag
-  bool is_my_tag = is_same(my_tag(), select_system(my_tag(), my_tag()));
-  ASSERT_EQUAL(true, is_my_tag);
+  // select_system(my_system, my_system) should return my_system
+  bool is_my_system = is_same(my_system(), select_system(my_system(), my_system()));
+  ASSERT_EQUAL(true, is_my_system);
 }
 DECLARE_UNITTEST(TestSelectSystemSameTypes);
 
 
-//template<typename T>
-//  thrust::pair<thrust::pointer<T,my_tag>, std::ptrdiff_t>
-//    get_temporary_buffer(my_tag, std::ptrdiff_t n)
-//{
-//  // communicate that my version of get_temporary_buffer
-//  // was correctly dispatched
-//  throw my_tag();
-//}
-//
+void TestGetTemporaryBuffer()
+{
+  const size_t n = 9001;
+
+  thrust::device_system_tag dev_tag;
+  typedef thrust::pointer<int, thrust::device_system_tag> pointer;
+  thrust::pair<pointer, std::ptrdiff_t> ptr_and_sz = thrust::get_temporary_buffer<int>(dev_tag, n);
+
+  ASSERT_EQUAL(ptr_and_sz.second, n);
+
+  const int ref_val = 13;
+  thrust::device_vector<int> ref(n, ref_val);
+
+  thrust::fill_n(ptr_and_sz.first, n, ref_val);
+
+  ASSERT_EQUAL(true, thrust::all_of(ptr_and_sz.first, ptr_and_sz.first + n, thrust::placeholders::_1 == ref_val));
+
+  thrust::return_temporary_buffer(dev_tag, ptr_and_sz.first);
+}
+DECLARE_UNITTEST(TestGetTemporaryBuffer);
+
+
+void TestMalloc()
+{
+  const size_t n = 9001;
+
+  thrust::device_system_tag dev_tag;
+  typedef thrust::pointer<int, thrust::device_system_tag> pointer;
+  pointer ptr = pointer(static_cast<int*>(thrust::malloc(dev_tag, sizeof(int) * n).get()));
+
+  const int ref_val = 13;
+  thrust::device_vector<int> ref(n, ref_val);
+
+  thrust::fill_n(ptr, n, ref_val);
+
+  ASSERT_EQUAL(true, thrust::all_of(ptr, ptr + n, thrust::placeholders::_1 == ref_val));
+
+  thrust::free(dev_tag, ptr);
+}
+DECLARE_UNITTEST(TestMalloc);
+
+
+static bool g_correctly_dispatched;
+
+
+template<typename T>
+  thrust::pair<thrust::pointer<T,my_system>, std::ptrdiff_t>
+    get_temporary_buffer(my_system sys, std::ptrdiff_t n)
+{
+  // communicate that my version of get_temporary_buffer
+  // was correctly dispatched
+  g_correctly_dispatched = true;
+
+  thrust::device_system_tag device_sys;
+  thrust::pair<thrust::pointer<T, thrust::device_system_tag>, std::ptrdiff_t> result = thrust::get_temporary_buffer<T>(device_sys, n);
+  return thrust::make_pair(thrust::pointer<T,my_system>(result.first.get()), result.second);
+}
+
+
 void TestGetTemporaryBufferDispatchImplicit()
 {
-  KNOWN_FAILURE;
-
-//  bool correctly_dispatched = false;
-//
-//  try
-//  {
-//    thrust::device_vector<int> vec(2);
-//
-//    // call something we know will invoke get_temporary_buffer
-//    thrust::sort(thrust::retag<my_tag>(vec.begin()),
-//                 thrust::retag<my_tag>(vec.end()));
-//  }
-//  catch(my_tag)
-//  {
-//    correctly_dispatched = true;
-//  }
-//
-//  ASSERT_EQUAL(true, correctly_dispatched);
+  if(is_same(thrust::device_system_tag(), thrust::system::cpp::tag()))
+  {
+    // XXX cpp uses the internal scalar backend, which currently elides user tags
+    KNOWN_FAILURE;
+  }
+  else
+  {
+    g_correctly_dispatched = false;
+
+    thrust::device_vector<int> vec(9001);
+
+    // call something we know will invoke get_temporary_buffer
+    my_system sys;
+    thrust::sort(sys, vec.begin(), vec.end());
+
+    ASSERT_EQUAL(true, g_correctly_dispatched);
+  }
 }
 DECLARE_UNITTEST(TestGetTemporaryBufferDispatchImplicit);
 

thrust/detail/allocator/temporary_allocator.inl

@@ -16,43 +16,21 @@
 
 #include <thrust/detail/config.h>
 #include <thrust/detail/allocator/temporary_allocator.h>
-#include <thrust/system/detail/generic/select_system.h>
-#include <thrust/system/detail/generic/memory.h>
+#include <thrust/detail/temporary_buffer.h>
 #include <thrust/system/detail/bad_alloc.h>
-#include <thrust/pair.h>
-#include <thrust/detail/raw_pointer_cast.h>
 
 namespace thrust
 {
 namespace detail
 {
-namespace temporary_allocator_detail
-{
-
-template<typename T, typename System, typename Pair>
-  thrust::pair<thrust::pointer<T,System>, typename thrust::pointer<T,System>::difference_type>
-    down_cast_pair(Pair p)
-{
-  // XXX should use a hypothetical thrust::static_pointer_cast here
-  thrust::pointer<T,System> ptr = thrust::pointer<T,System>(static_cast<T*>(thrust::raw_pointer_cast(p.first)));
-
-  typedef thrust::pair<thrust::pointer<T,System>, typename thrust::pointer<T,System>::difference_type> result_type;
-  return result_type(ptr, p.second);
-} // end down_cast_pair()
-
-
-} // end temporary_allocator_detail
 
 
 template<typename T, typename System>
   typename temporary_allocator<T,System>::pointer
     temporary_allocator<T,System>
       ::allocate(typename temporary_allocator<T,System>::size_type cnt)
 {
-  using thrust::system::detail::generic::get_temporary_buffer;
-
-  // XXX use thrust::get_temporary_buffer here should we add it
-  pointer_and_size result = temporary_allocator_detail::down_cast_pair<T,System>(get_temporary_buffer<T>(m_system.derived(), cnt));
+  pointer_and_size result = thrust::get_temporary_buffer<T>(m_system, cnt);
 
   // handle failure
   if(result.second < cnt)
@@ -71,12 +49,10 @@ template<typename T, typename System>
   void temporary_allocator<T,System>
     ::deallocate(typename temporary_allocator<T,System>::pointer p, typename temporary_allocator<T,System>::size_type n)
 {
-  using thrust::system::detail::generic::return_temporary_buffer;
-
-  // XXX use thrust::return_temporary_buffer here should we add it
-  return return_temporary_buffer(m_system.derived(), p);
+  return thrust::return_temporary_buffer(m_system, p);
 } // end temporary_allocator
 
+
 } // end detail
 } // end thrust
 

thrust/detail/malloc_and_free.h

@@ -27,7 +27,7 @@ namespace thrust
 {
 
 template<typename System>
-pointer<void,System> malloc(thrust::dispatchable<System> &system, std::size_t n)
+pointer<void,System> malloc(thrust::detail::dispatchable_base<System> &system, std::size_t n)
 {
   using thrust::system::detail::generic::malloc;
 
@@ -55,7 +55,7 @@ void free(int *volatile ptr)
 #endif // THRUST_DEVICE_COMPILER
 
 template<typename System, typename Pointer>
-void free(thrust::dispatchable<System> &system, Pointer ptr)
+void free(thrust::detail::dispatchable_base<System> &system, Pointer ptr)
 {
   using thrust::system::detail::generic::free;
 

thrust/detail/temporary_buffer.h

@@ -0,0 +1,70 @@
+/*
+ *  Copyright 2008-2012 NVIDIA Corporation
+ *
+ *  Licensed under the Apache License, Version 2.0 (the "License");
+ *  you may not use this file except in compliance with the License.
+ *  You may obtain a copy of the License at
+ *
+ *      http://www.apache.org/licenses/LICENSE-2.0
+ *
+ *  Unless required by applicable law or agreed to in writing, software
+ *  distributed under the License is distributed on an "AS IS" BASIS,
+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ *  See the License for the specific language governing permissions and
+ *  limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/dispatchable.h>
+#include <thrust/pair.h>
+#include <thrust/detail/pointer.h>
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/system/detail/generic/memory.h>
+
+namespace thrust
+{
+namespace detail
+{
+namespace get_temporary_buffer_detail
+{
+
+
+template<typename T, typename System, typename Pair>
+  thrust::pair<thrust::pointer<T,System>, typename thrust::pointer<T,System>::difference_type>
+    down_cast_pair(Pair p)
+{
+  // XXX should use a hypothetical thrust::static_pointer_cast here
+  thrust::pointer<T,System> ptr = thrust::pointer<T,System>(static_cast<T*>(thrust::raw_pointer_cast(p.first)));
+
+  typedef thrust::pair<thrust::pointer<T,System>, typename thrust::pointer<T,System>::difference_type> result_type;
+  return result_type(ptr, p.second);
+} // end down_cast_pair()
+
+
+} // end get_temporary_buffer_detail
+} // end detail
+
+
+template<typename T, typename System>
+  thrust::pair<thrust::pointer<T,System>, typename thrust::pointer<T,System>::difference_type>
+    get_temporary_buffer(thrust::detail::dispatchable_base<System> &system, typename thrust::pointer<T,System>::difference_type n)
+{
+  using thrust::system::detail::generic::get_temporary_buffer;
+
+  return thrust::detail::get_temporary_buffer_detail::down_cast_pair<T,System>(get_temporary_buffer<T>(system.derived(), n));
+} // end get_temporary_buffer()
+
+
+template<typename System, typename Pointer>
+  void return_temporary_buffer(thrust::detail::dispatchable_base<System> &system, Pointer p)
+{
+  using thrust::system::detail::generic::return_temporary_buffer;
+
+  return return_temporary_buffer(system.derived(), p);
+} // end return_temporary_buffer()
+
+
+} // end thrust
+

thrust/memory.h

@@ -26,6 +26,7 @@
 #include <thrust/detail/raw_pointer_cast.h>
 #include <thrust/detail/raw_reference_cast.h>
 #include <thrust/detail/malloc_and_free.h>
+#include <thrust/detail/temporary_buffer.h>
 
 namespace thrust
 {

thrust/system/cpp/detail/malloc_and_free.h

@@ -34,14 +34,14 @@ namespace detail
 // note that malloc returns a raw pointer to avoid
 // depending on the heavyweight thrust/system/cpp/memory.h header
 template<typename System>
-inline void *malloc(thrust::system::cpp::detail::dispatchable<System> &, std::size_t n)
+  void *malloc(dispatchable<System> &, std::size_t n)
 {
   return std::malloc(n);
 } // end malloc()
 
 
 template<typename System, typename Pointer>
-inline void free(thrust::system::cpp::detail::dispatchable<System> &, Pointer ptr)
+  void free(dispatchable<System> &, Pointer ptr)
 {
   std::free(thrust::raw_pointer_cast(ptr));
 } // end free()

thrust/system/cuda/detail/malloc_and_free.h

@@ -36,7 +36,8 @@ namespace detail
 
 // note that malloc returns a raw pointer to avoid
 // depending on the heavyweight thrust/system/cuda/memory.h header
-inline void *malloc(tag, std::size_t n)
+template<typename System>
+  void *malloc(dispatchable<System> &, std::size_t n)
 {
   void *result = 0;
 
@@ -51,8 +52,8 @@ inline void *malloc(tag, std::size_t n)
 } // end malloc()
 
 
-template<typename Pointer>
-inline void free(tag, Pointer ptr)
+template<typename System, typename Pointer>
+  void free(dispatchable<System> &, Pointer ptr)
 {
   cudaError_t error = cudaFree(thrust::raw_pointer_cast(ptr));
 

thrust/system/cuda/detail/memory.inl

@@ -71,12 +71,14 @@ void swap(reference<T> a, reference<T> b)
 
 pointer<void> malloc(std::size_t n)
 {
-  return pointer<void>(thrust::system::cuda::detail::malloc(tag(), n));
+  tag cuda_tag;
+  return pointer<void>(thrust::system::cuda::detail::malloc(cuda_tag, n));
 } // end malloc()
 
 void free(pointer<void> ptr)
 {
-  return thrust::system::cuda::detail::free(tag(), ptr.get());
+  tag cuda_tag;
+  return thrust::system::cuda::detail::free(cuda_tag, ptr.get());
 } // end free()
 
 } // end cuda