Permalink
Browse files

Merge branch 'master' of github.com:jaredhoberock/thrust into normal_…

…distribution-fix-34
  • Loading branch information...
2 parents 76b1128 + 30c3573 commit bde962e8871545a1e8e51dc66e7cd421ed019a72 @jaredhoberock jaredhoberock committed Jun 12, 2012
@@ -117,6 +117,9 @@ bool CUDATestDriver::run_tests(const ArgumentSet &args, const ArgumentMap &kwarg
std::cout << "--verbose and --concise cannot be used together" << std::endl;
exit(EXIT_FAILURE);
}
+
+ // check error status before doing anything
+ if(check_cuda_error(concise)) return false;
bool result = true;
@@ -18,6 +18,7 @@
#include <thrust/detail/allocator/allocator_traits.h>
#include <thrust/detail/type_traits/pointer_traits.h>
#include <thrust/for_each.h>
+#include <memory>
namespace thrust
{
@@ -27,37 +28,63 @@ namespace allocator_traits_detail
{
-// when T has a trivial destructor and Allocator has no
-// destroy member function, destroying T has no effect
+// destroy_range has three cases:
+// if Allocator has an effectful member function destroy:
+// 1. destroy via the allocator
+// else
+// 2. if T has a non-trivial destructor, destroy the range without using the allocator
+// 3. if T has a trivial destructor, do a no-op
+
template<typename Allocator, typename T>
- struct has_no_effect_destroy
- : integral_constant<
- bool,
- has_trivial_destructor<T>::value && !has_member_destroy1<Allocator,T>::value
- >
+ struct has_effectful_member_destroy1
+ : has_member_destroy1<Allocator,T>
{};
-// we know that std::allocator::destroy's only effect is to
-// call T's destructor, so we needn't use it when destroying T
+// std::allocator::destroy's only effect is to invoke its argument's destructor
template<typename U, typename T>
- struct has_no_effect_destroy<std::allocator<U>, T>
- : has_trivial_destructor<T>
+ struct has_effectful_member_destroy1<std::allocator<U>, T>
+ : thrust::detail::false_type
{};
+// case 1: Allocator has an effectful 1-argument member function "destroy"
+template<typename Allocator, typename Pointer>
+ struct enable_if_destroy_range_case1
+ : thrust::detail::enable_if<
+ has_effectful_member_destroy1<
+ Allocator,
+ typename pointer_element<Pointer>::type
+ >::value
+ >
+{};
+// case 2: Allocator has no member function "destroy", but T has a non-trivial destructor
+template<typename Allocator, typename Pointer>
+ struct enable_if_destroy_range_case2
+ : thrust::detail::enable_if<
+ !has_effectful_member_destroy1<
+ Allocator,
+ typename pointer_element<Pointer>::type
+ >::value &&
+ !has_trivial_destructor<
+ typename pointer_element<Pointer>::type
+ >::value
+ >
+{};
+
+// case 3: Allocator has no member function "destroy", and T has a trivial destructor
+template<typename Allocator, typename Pointer>
+ struct enable_if_destroy_range_case3
+ : thrust::detail::enable_if<
+ !has_effectful_member_destroy1<
+ Allocator,
+ typename pointer_element<Pointer>::type
+ >::value &&
+ has_trivial_destructor<
+ typename pointer_element<Pointer>::type
+ >::value
+ >
+{};
-// destroy_range is a no op if element destruction has no effect
-template<typename Allocator, typename Pointer, typename Size>
- typename enable_if<
- has_no_effect_destroy<
- Allocator,
- typename pointer_element<Pointer>::type
- >::value
- >::type
- destroy_range(Allocator &, Pointer, Size)
-{
- // no op
-}
template<typename Allocator>
@@ -78,21 +105,44 @@ template<typename Allocator>
};
-// destroy_range is effectful if element destroy has effect
-// XXX need to check whether we really need to destroy via the allocator or not
+// destroy_range case 1: destroy via allocator
template<typename Allocator, typename Pointer, typename Size>
- typename disable_if<
- has_no_effect_destroy<
- Allocator,
- typename pointer_element<Pointer>::type
- >::value
- >::type
+ typename enable_if_destroy_range_case1<Allocator,Pointer>::type
destroy_range(Allocator &a, Pointer p, Size n)
{
thrust::for_each_n(p, n, destroy_via_allocator<Allocator>(a));
}
+// we must prepare for His coming
+struct gozer
+{
+ template<typename T>
+ inline __host__ __device__
+ void operator()(T &x)
+ {
+ x.~T();
+ }
+};
+
+// destroy_range case 2: destroy without the allocator
+template<typename Allocator, typename Pointer, typename Size>
+ typename enable_if_destroy_range_case2<Allocator,Pointer>::type
+ destroy_range(Allocator &, Pointer p, Size n)
+{
+ thrust::for_each_n(p, n, gozer());
+}
+
+
+// destroy_range case 3: no-op
+template<typename Allocator, typename Pointer, typename Size>
+ typename enable_if_destroy_range_case3<Allocator,Pointer>::type
+ destroy_range(Allocator &, Pointer, Size)
+{
+ // no op
+}
+
+
} // end allocator_traits_detail
@@ -15,10 +15,12 @@
*/
#include <thrust/detail/config.h>
+#include <thrust/detail/type_traits.h>
#include <thrust/detail/allocator/allocator_traits.h>
#include <thrust/detail/type_traits/pointer_traits.h>
#include <thrust/for_each.h>
#include <thrust/uninitialized_fill.h>
+#include <memory>
namespace thrust
{
@@ -27,6 +29,23 @@ namespace detail
namespace allocator_traits_detail
{
+// fill_construct_range has 2 cases:
+// if Allocator has an effectful member function construct:
+// 1. construct via the allocator
+// else
+// 2. construct via uninitialized_fill
+
+template<typename Allocator, typename T, typename Arg1>
+ struct has_effectful_member_construct2
+ : has_member_construct2<Allocator,T,Arg1>
+{};
+
+// std::allocator::construct's only effect is to invoke placement new
+template<typename U, typename T, typename Arg1>
+ struct has_effectful_member_construct2<std::allocator<U>,T,Arg1>
+ : thrust::detail::false_type
+{};
+
template<typename Allocator, typename Arg1>
struct construct2_via_allocator
@@ -47,19 +66,9 @@ template<typename Allocator, typename Arg1>
};
-template<typename Allocator, typename T, typename Arg1>
- struct needs_copy_construct_via_allocator
- : has_member_construct2<
- Allocator,
- T,
- Arg1
- >
-{};
-
-
template<typename Allocator, typename Pointer, typename Size, typename T>
typename enable_if<
- needs_copy_construct_via_allocator<
+ has_effectful_member_construct2<
Allocator,
typename pointer_element<Pointer>::type,
T
@@ -73,7 +82,7 @@ template<typename Allocator, typename Pointer, typename Size, typename T>
template<typename Allocator, typename Pointer, typename Size, typename T>
typename disable_if<
- needs_copy_construct_via_allocator<
+ has_effectful_member_construct2<
Allocator,
typename pointer_element<Pointer>::type,
T
@@ -66,6 +66,10 @@ struct function_attributes_t
size_t sharedSizeBytes;
};
+/*! Returns the current device ordinal.
+ */
+inline int current_device();
+
/*! Returns a copy of the device_properties_t structure
* that is associated with a given device.
*/
@@ -135,19 +135,24 @@ inline device_properties_t device_properties(int device_id)
return device_properties[device_id];
}
-inline device_properties_t device_properties(void)
+inline int current_device()
{
- int device_id = -1;
+ int result = -1;
- cudaError_t error = cudaGetDevice(&device_id);
+ cudaError_t error = cudaGetDevice(&result);
if(error)
throw thrust::system_error(error, thrust::cuda_category());
- if(device_id < 0)
+ if(result < 0)
throw thrust::system_error(cudaErrorNoDevice, thrust::cuda_category());
- return device_properties(device_id);
+ return result;
+}
+
+inline device_properties_t device_properties(void)
+{
+ return device_properties(current_device());
}
template <typename KernelFunction>
@@ -182,21 +182,32 @@ arch::function_attributes_t closure_attributes(void)
typedef closure_launcher<Closure> Launcher;
// cache the result of function_attributes(), because it is slow
- static bool result_exists = false;
- static arch::function_attributes_t result = {};
+ // only cache the first few devices
+ static const int max_num_devices = 16;
- if(!result_exists)
+ static bool attributes_exist[max_num_devices] = {0};
+ static arch::function_attributes_t function_attributes[max_num_devices] = {};
+
+ // XXX device_id ought to be an argument to this function
+ int device_id = arch::current_device();
+
+ if(device_id >= max_num_devices)
+ {
+ return arch::function_attributes(Launcher::get_launch_function());
+ }
+
+ if(!attributes_exist[device_id])
{
- result = arch::function_attributes(Launcher::get_launch_function());
+ function_attributes[device_id] = arch::function_attributes(Launcher::get_launch_function());
- // disallow the compiler to move the write to result_exists
- // before the initialization of result
+ // disallow the compiler to move the write to attributes_exist[device_id]
+ // before the initialization of function_attributes[device_id]
__thrust_compiler_fence();
- result_exists = true;
+ attributes_exist[device_id] = true;
}
- return result;
+ return function_attributes[device_id];
}
} // end namespace detail

0 comments on commit bde962e

Please sign in to comment.