thrust: CUDA specific zero init hack

Improve performance when pre-initialized memory is not required.
This is not in main branch because it is not portablel to rocThrust,
and may depend on CUDA/thrust version (but so far it works with
CUDA 10 and 11).

include/thrust/detail/allocator/default_construct_range.h

@@ -0,0 +1,37 @@
+/*
+ *  Copyright 2008-2013 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>
+
+namespace thrust
+{
+namespace detail
+{
+
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+inline void default_construct_range(Allocator &a, Pointer p, Size n);
+
+
+} // end detail
+} // end thrust
+
+#include <thrust/detail/allocator/default_construct_range.inl>
+
+

include/thrust/detail/allocator/default_construct_range.inl

@@ -0,0 +1,111 @@
+/*
+ *  Copyright 2008-2013 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.
+ */
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/allocator/allocator_traits.h>
+#include <thrust/detail/type_traits.h>
+#include <thrust/detail/type_traits/pointer_traits.h>
+#include <thrust/for_each.h>
+#include <thrust/uninitialized_fill.h>
+
+namespace thrust
+{
+namespace detail
+{
+namespace allocator_traits_detail
+{
+
+
+template<typename Allocator>
+  struct construct1_via_allocator
+{
+  Allocator &a;
+
+  __host__ __device__
+  construct1_via_allocator(Allocator &a)
+    : a(a)
+  {}
+
+  template<typename T>
+  inline __host__ __device__
+  void operator()(T &x)
+  {
+    allocator_traits<Allocator>::construct(a, &x);
+  }
+};
+
+
+// we need to construct T via the allocator if...
+template<typename Allocator, typename T>
+  struct needs_default_construct_via_allocator
+    : thrust::detail::or_<
+        has_member_construct1<Allocator,T>,               // if the Allocator does something interesting
+        thrust::detail::not_<has_trivial_constructor<T> > // or if T's default constructor does something interesting
+      >
+{};
+
+
+// we know that std::allocator::construct's only effect is to call T's 
+// default constructor, so we needn't use it for default construction
+// unless T's constructor does something interesting
+template<typename U, typename T>
+  struct needs_default_construct_via_allocator<std::allocator<U>, T>
+    : thrust::detail::not_<has_trivial_constructor<T> >
+{};
+
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  typename enable_if<
+    needs_default_construct_via_allocator<
+      Allocator,
+      typename pointer_element<Pointer>::type
+    >::value
+  >::type
+    default_construct_range(Allocator &a, Pointer p, Size n)
+{
+//  thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, construct1_via_allocator<Allocator>(a));
+}
+
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  typename disable_if<
+    needs_default_construct_via_allocator<
+      Allocator,
+      typename pointer_element<Pointer>::type
+    >::value
+  >::type
+    default_construct_range(Allocator &a, Pointer p, Size n)
+{
+//  thrust::uninitialized_fill_n(allocator_system<Allocator>::get(a), p, n, typename pointer_element<Pointer>::type());
+}
+
+
+} // end allocator_traits_detail
+
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  void default_construct_range(Allocator &a, Pointer p, Size n)
+{
+  return allocator_traits_detail::default_construct_range(a,p,n);
+}
+
+
+} // end detail
+} // end thrust
+

include/thrust/detail/allocator/destroy_range.h

@@ -0,0 +1,34 @@
+/*
+ *  Copyright 2008-2013 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>
+
+namespace thrust
+{
+namespace detail
+{
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  inline void destroy_range(Allocator &a, Pointer p, Size n);
+
+} // end detail
+} // end thrust
+
+#include <thrust/detail/allocator/destroy_range.inl>
+

include/thrust/detail/allocator/destroy_range.inl

@@ -0,0 +1,164 @@
+/*
+ *  Copyright 2008-2013 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.
+ */
+
+#include <thrust/detail/allocator/destroy_range.h>
+#include <thrust/detail/allocator/allocator_traits.h>
+#include <thrust/detail/type_traits/pointer_traits.h>
+#include <thrust/for_each.h>
+#include <memory>
+
+namespace thrust
+{
+namespace detail
+{
+namespace allocator_traits_detail
+{
+
+
+// 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_effectful_member_destroy
+    : has_member_destroy<Allocator,T>
+{};
+
+// std::allocator::destroy's only effect is to invoke its argument's destructor
+template<typename U, typename T>
+  struct has_effectful_member_destroy<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_destroy<
+          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_destroy<
+          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_destroy<
+          Allocator,
+          typename pointer_element<Pointer>::type
+        >::value &&
+        has_trivial_destructor<
+          typename pointer_element<Pointer>::type
+        >::value
+      >
+{};
+
+
+
+template<typename Allocator>
+  struct destroy_via_allocator
+{
+  Allocator &a;
+
+  __host__ __device__
+  destroy_via_allocator(Allocator &a)
+    : a(a)
+  {}
+
+  template<typename T>
+  inline __host__ __device__
+  void operator()(T &x)
+  {
+    allocator_traits<Allocator>::destroy(a, &x);
+  }
+};
+
+
+// destroy_range case 1: destroy via allocator
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  typename enable_if_destroy_range_case1<Allocator,Pointer>::type
+    destroy_range(Allocator &a, Pointer p, Size n)
+{
+  //thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, destroy_via_allocator<Allocator>(a));
+}
+
+
+// we must prepare for His coming
+struct gozer
+{
+  __thrust_exec_check_disable__
+  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>
+__host__ __device__
+  typename enable_if_destroy_range_case2<Allocator,Pointer>::type
+    destroy_range(Allocator &a, Pointer p, Size n)
+{
+  //thrust::for_each_n(allocator_system<Allocator>::get(a), p, n, gozer());
+}
+
+
+// destroy_range case 3: no-op
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  typename enable_if_destroy_range_case3<Allocator,Pointer>::type
+    destroy_range(Allocator &, Pointer, Size)
+{
+  // no op
+}
+
+
+} // end allocator_traits_detail
+
+
+template<typename Allocator, typename Pointer, typename Size>
+__host__ __device__
+  void destroy_range(Allocator &a, Pointer p, Size n)
+{
+  return allocator_traits_detail::destroy_range(a,p,n);
+}
+
+
+} // end detail
+} // end thrust
+