Backport changes for gcc-4.2.1 compatibility

Fix typo in CHANGELOG
NVIDIA · Jul 18, 2012 · 343b34d · 343b34d
1 parent 9065ccb
commit 343b34d
Show file tree

Hide file tree

Showing 3 changed files with 227 additions and 131 deletions.
diff --git a/CHANGELOG b/CHANGELOG
@@ -1,5 +1,5 @@
 #######################################
-#           Thrust v1.5.2             #
+#           Thrust v1.5.3             #
 #######################################
 
 Summary

diff --git a/thrust/detail/backend/cuda/detail/alignment.h b/thrust/detail/backend/cuda/detail/alignment.h
@@ -0,0 +1,223 @@
+/*
+ *  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
+
+namespace thrust
+{
+namespace detail
+{
+namespace backend
+{
+namespace cuda
+{
+namespace detail
+{
+namespace alignment_of_detail
+{
+
+
+template<typename T> class alignment_of_impl;
+
+template<typename T, std::size_t size_diff>
+  struct helper
+{
+  static const std::size_t value = size_diff;
+};
+
+template<typename T>
+  class helper<T,0>
+{
+  public:
+    static const std::size_t value = alignment_of_impl<T>::value;
+};
+
+template<typename T>
+  class alignment_of_impl
+{
+  private:
+    struct big { T x; char c; };
+
+  public:
+    static const std::size_t value = helper<big, sizeof(big) - sizeof(T)>::value;
+};
+
+
+} // end alignment_of_detail
+
+
+template<typename T>
+  struct alignment_of
+    : alignment_of_detail::alignment_of_impl<T>
+{};
+
+
+template<std::size_t Align> struct aligned_type;
+
+// __align__ is CUDA-specific, so guard it
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+
+// implementing aligned_type portably is tricky:
+
+#  if THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC
+// implement aligned_type with specialization because MSVC
+// requires literals as arguments to declspec(align(n))
+template<> struct aligned_type<1>
+{
+  struct __align__(1) type { };
+};
+
+template<> struct aligned_type<2>
+{
+  struct __align__(2) type { };
+};
+
+template<> struct aligned_type<4>
+{
+  struct __align__(4) type { };
+};
+
+template<> struct aligned_type<8>
+{
+  struct __align__(8) type { };
+};
+
+template<> struct aligned_type<16>
+{
+  struct __align__(16) type { };
+};
+
+template<> struct aligned_type<32>
+{
+  struct __align__(32) type { };
+};
+
+template<> struct aligned_type<64>
+{
+  struct __align__(64) type { };
+};
+
+template<> struct aligned_type<128>
+{
+  struct __align__(128) type { };
+};
+
+template<> struct aligned_type<256>
+{
+  struct __align__(256) type { };
+};
+
+template<> struct aligned_type<512>
+{
+  struct __align__(512) type { };
+};
+
+template<> struct aligned_type<1024>
+{
+  struct __align__(1024) type { };
+};
+
+template<> struct aligned_type<2048>
+{
+  struct __align__(2048) type { };
+};
+
+template<> struct aligned_type<4096>
+{
+  struct __align__(4096) type { };
+};
+
+template<> struct aligned_type<8192>
+{
+  struct __align__(8192) type { };
+};
+#  elif (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && (THRUST_GCC_VERSION < 40300)
+// implement aligned_type with specialization because gcc 4.2
+// requires literals as arguments to __attribute__(aligned(n))
+template<> struct aligned_type<1>
+{
+  struct __align__(1) type { };
+};
+
+template<> struct aligned_type<2>
+{
+  struct __align__(2) type { };
+};
+
+template<> struct aligned_type<4>
+{
+  struct __align__(4) type { };
+};
+
+template<> struct aligned_type<8>
+{
+  struct __align__(8) type { };
+};
+
+template<> struct aligned_type<16>
+{
+  struct __align__(16) type { };
+};
+
+template<> struct aligned_type<32>
+{
+  struct __align__(32) type { };
+};
+
+template<> struct aligned_type<64>
+{
+  struct __align__(64) type { };
+};
+
+template<> struct aligned_type<128>
+{
+  struct __align__(128) type { };
+};
+
+#  else
+// assume the compiler allows template parameters as
+// arguments to __align__ 
+template<std::size_t Align> struct aligned_type
+{
+  struct __align__(Align) type { };
+};
+#  endif // THRUST_HOST_COMPILER
+#else
+template<std::size_t Align> struct aligned_type
+{
+  struct type { };
+};
+#endif // THRUST_DEVICE_COMPILER
+
+
+template<std::size_t Len, std::size_t Align>
+  struct aligned_storage
+{
+  union type
+  {
+    unsigned char data[Len];
+
+    typename aligned_type<Align>::type align;
+  };
+};
+
+
+} // end detail
+} // end cuda
+} // end backend
+} // end detail
+} // end thrust
+
diff --git a/thrust/detail/backend/cuda/detail/uninitialized.h b/thrust/detail/backend/cuda/detail/uninitialized.h
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <thrust/detail/config.h>
+#include <thrust/detail/backend/cuda/detail/alignment.h>
 #include <cstddef>
 #include <new>
 
@@ -30,143 +31,15 @@ namespace cuda
 {
 namespace detail
 {
-namespace uninitialized_detail
-{
-
-
-template<typename T> struct alignment_of_impl;
-
-template<typename T, std::size_t size_diff>
-  struct helper
-{
-  static const std::size_t value = size_diff;
-};
-
-template<typename T>
-  struct helper<T,0>
-{
-  static const std::size_t value = alignment_of_impl<T>::value;
-};
-
-template<typename T>
-  struct alignment_of_impl
-{
-  struct big { T x; char c; };
-
-  static const std::size_t value = helper<big, sizeof(big) - sizeof(T)>::value;
-};
-
-
-template<typename T>
-  struct alignment_of
-    : uninitialized_detail::alignment_of_impl<T>
-{};
-
-
-template<std::size_t Align> struct aligned_type;
-
-// __align__ is CUDA-specific, so guard it
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-// implement aligned_type with specialization because MSVC
-// requires literals as arguments to declspec(align(n))
-template<> struct aligned_type<1>
-{
-  struct __align__(1) type { };
-};
-
-template<> struct aligned_type<2>
-{
-  struct __align__(2) type { };
-};
-
-template<> struct aligned_type<4>
-{
-  struct __align__(4) type { };
-};
-
-template<> struct aligned_type<8>
-{
-  struct __align__(8) type { };
-};
-
-template<> struct aligned_type<16>
-{
-  struct __align__(16) type { };
-};
-
-template<> struct aligned_type<32>
-{
-  struct __align__(32) type { };
-};
-
-template<> struct aligned_type<64>
-{
-  struct __align__(64) type { };
-};
-
-template<> struct aligned_type<128>
-{
-  struct __align__(128) type { };
-};
-
-template<> struct aligned_type<256>
-{
-  struct __align__(256) type { };
-};
-
-template<> struct aligned_type<512>
-{
-  struct __align__(512) type { };
-};
-
-template<> struct aligned_type<1024>
-{
-  struct __align__(1024) type { };
-};
-
-template<> struct aligned_type<2048>
-{
-  struct __align__(2048) type { };
-};
-
-template<> struct aligned_type<4096>
-{
-  struct __align__(4096) type { };
-};
-
-template<> struct aligned_type<8192>
-{
-  struct __align__(8192) type { };
-};
-#else
-template<std::size_t Align> struct aligned_type
-{
-  struct type { };
-};
-#endif // THRUST_DEVICE_COMPILER
-
-template<std::size_t Len, std::size_t Align>
-  struct aligned_storage
-{
-  union type
-  {
-    unsigned char data[Len];
-
-    typename aligned_type<Align>::type align;
-  };
-};
-
-
-} // end uninitialized_detail
 
 
 template<typename T>
   class uninitialized
 {
   private:
-    typename uninitialized_detail::aligned_storage<
+    typename aligned_storage<
       sizeof(T),
-      uninitialized_detail::alignment_of<T>::value
+      alignment_of<T>::value
     >::type storage;
 
     __device__ __forceinline__ const T* ptr() const