Add libcu++ dependency; initial round of NV_IF_TARGET ports.

.gitmodules

@@ -1,3 +1,6 @@
 [submodule "cub"]
 	path = dependencies/cub
 	url = ../cub.git
+[submodule "libcudacxx"]
+	path = dependencies/libcudacxx
+	url = ../libcudacxx.git

cmake/ThrustInstallRules.cmake

@@ -24,27 +24,35 @@ configure_file("${Thrust_SOURCE_DIR}/thrust/cmake/thrust-header-search.cmake.in"
 install(FILES "${Thrust_BINARY_DIR}/thrust/cmake/thrust-header-search.cmake"
   DESTINATION "${install_location}")
 
-# Depending on how Thrust is configured, CUB's CMake scripts may or may not be
-# included, so maintain a set of CUB install rules in both projects. By default
-# CUB headers are installed alongside Thrust -- this may be disabled by turning
-# off THRUST_INSTALL_CUB_HEADERS.
-option(THRUST_INSTALL_CUB_HEADERS "Include cub headers when installing." ON)
+# Depending on how Thrust is configured, libcudacxx and CUB's CMake scripts may
+# or may not be include()'d, so force include their install rules when requested.
+# By default, these projects are installed alongside Thrust. This is controlled by
+# THRUST_INSTALL_CUB_HEADERS and THRUST_INSTALL_LIBCUDACXX_HEADERS.
+option(THRUST_INSTALL_CUB_HEADERS "Include CUB headers when installing." ON)
 if (THRUST_INSTALL_CUB_HEADERS)
-  install(DIRECTORY "${Thrust_SOURCE_DIR}/dependencies/cub/cub"
-    DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}"
-    FILES_MATCHING
-      PATTERN "*.cuh"
-  )
+  # Use a function to limit scope of the CUB_*_DIR vars:
+  function(_thrust_install_cub_headers)
+    # Fake these for the logic in CUBInstallRules.cmake:
+    set(CUB_SOURCE_DIR "${Thrust_SOURCE_DIR}/dependencies/cub/")
+    set(CUB_BINARY_DIR "${Thrust_BINARY_DIR}/cub-config/")
+    set(CUB_ENABLE_INSTALL_RULES ON)
+    set(CUB_IN_THRUST OFF)
+    include("${Thrust_SOURCE_DIR}/dependencies/cub/cmake/CubInstallRules.cmake")
+  endfunction()
 
-  # Need to configure a file to store THRUST_INSTALL_HEADER_INFIX
-  install(DIRECTORY "${Thrust_SOURCE_DIR}/dependencies/cub/cub/cmake/"
-    DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/cub"
-    PATTERN cub-header-search EXCLUDE
-  )
-  set(install_location "${CMAKE_INSTALL_LIBDIR}/cmake/cub")
-  configure_file("${Thrust_SOURCE_DIR}/dependencies/cub/cub/cmake/cub-header-search.cmake.in"
-    "${Thrust_BINARY_DIR}/dependencies/cub/cub/cmake/cub-header-search.cmake"
-    @ONLY)
-  install(FILES "${Thrust_BINARY_DIR}/dependencies/cub/cub/cmake/cub-header-search.cmake"
-    DESTINATION "${install_location}")
+  _thrust_install_cub_headers()
+endif()
+
+option(THRUST_INSTALL_LIBCUDACXX_HEADERS "Include libcudacxx headers when installing." ON)
+if (THRUST_INSTALL_LIBCUDACXX_HEADERS)
+  # Use a function to limit scope of the libcudacxx_*_DIR vars:
+  function(_thrust_install_libcudacxx_headers)
+    # Fake these for the logic in libcudacxxInstallRules.cmake:
+    set(libcudacxx_SOURCE_DIR "${Thrust_SOURCE_DIR}/dependencies/libcudacxx/")
+    set(libcudacxx_BINARY_DIR "${Thrust_BINARY_DIR}/libcudacxx-config/")
+    set(libcudacxx_ENABLE_INSTALL_RULES ON)
+    include("${Thrust_SOURCE_DIR}/dependencies/libcudacxx/cmake/libcudacxxInstallRules.cmake")
+  endfunction()
+
+  _thrust_install_libcudacxx_headers()
 endif()

testing/allocator.cu

@@ -2,6 +2,9 @@
 #include <thrust/detail/config.h>
 #include <thrust/device_malloc_allocator.h>
 #include <thrust/system/cpp/vector.h>
+
+#include <nv/target>
+
 #include <memory>
 
 template <typename T>
@@ -60,9 +63,12 @@ DECLARE_VARIABLE_UNITTEST(TestAllocatorCustomCopyConstruct);
 template <typename T>
 struct my_allocator_with_custom_destroy
 {
-  typedef T         value_type;
-  typedef T &       reference;
-  typedef const T & const_reference;
+  // This is only used with thrust::cpp::vector:
+  using system_type = thrust::cpp::tag;
+
+  using value_type = T;
+  using reference = T &;
+  using const_reference = const T &;
 
   static bool g_state;
 
@@ -80,9 +86,7 @@ struct my_allocator_with_custom_destroy
   __host__ __device__
   void destroy(T *)
   {
-#if !__CUDA_ARCH__
-    g_state = true;
-#endif
+    NV_IF_TARGET(NV_IS_HOST, (g_state = true;));
   }
 
   value_type *allocate(std::ptrdiff_t n)
@@ -119,12 +123,14 @@ bool my_allocator_with_custom_destroy<T>::g_state = false;
 template <typename T>
 void TestAllocatorCustomDestroy(size_t n)
 {
+  my_allocator_with_custom_destroy<T>::g_state = false;
+
   {
     thrust::cpp::vector<T, my_allocator_with_custom_destroy<T> > vec(n);
   } // destroy everything
 
-  if (0 < n)
-    ASSERT_EQUAL(true, my_allocator_with_custom_destroy<T>::g_state);
+  // state should only be true when there are values to destroy:
+  ASSERT_EQUAL(n > 0, my_allocator_with_custom_destroy<T>::g_state);
 }
 DECLARE_VARIABLE_UNITTEST(TestAllocatorCustomDestroy);
 
@@ -203,7 +209,6 @@ void TestAllocatorTraitsRebind()
 }
 DECLARE_UNITTEST(TestAllocatorTraitsRebind);
 
-#if THRUST_CPP_DIALECT >= 2011
 void TestAllocatorTraitsRebindCpp11()
 {
   ASSERT_EQUAL(
@@ -251,5 +256,3 @@ void TestAllocatorTraitsRebindCpp11()
   );
 }
 DECLARE_UNITTEST(TestAllocatorTraitsRebindCpp11);
-#endif // C++11
-

testing/cuda/pair_sort.cu

@@ -4,16 +4,11 @@
 #include <thrust/execution_policy.h>
 
 
-template<typename ExecutionPolicy, typename Iterator1, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator>
 __global__
-void stable_sort_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 is_supported)
+void stable_sort_kernel(ExecutionPolicy exec, Iterator first, Iterator last)
 {
-#if (__CUDA_ARCH__ >= 200)
-  *is_supported = true;
   thrust::stable_sort(exec, first, last);
-#else
-  *is_supported = false;
-#endif
 }
 
 
@@ -43,19 +38,14 @@ void TestPairStableSortDevice(ExecutionPolicy exec)
 
   thrust::device_vector<P> d_pairs = h_pairs;
 
-  thrust::device_vector<bool> is_supported(1);
-
-  stable_sort_kernel<<<1,1>>>(exec, d_pairs.begin(), d_pairs.end(), is_supported.begin());
+  stable_sort_kernel<<<1,1>>>(exec, d_pairs.begin(), d_pairs.end());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
-  if(is_supported[0])
-  {
-    // sort on the host
-    thrust::stable_sort(h_pairs.begin(), h_pairs.end());
+  // sort on the host
+  thrust::stable_sort(h_pairs.begin(), h_pairs.end());
 
-    ASSERT_EQUAL_QUIET(h_pairs, d_pairs);
-  }
+  ASSERT_EQUAL_QUIET(h_pairs, d_pairs);
 };
 
 

testing/cuda/pair_sort_by_key.cu

@@ -6,16 +6,11 @@
 #include <thrust/execution_policy.h>
 
 
-template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3>
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2>
 __global__
-void stable_sort_by_key_kernel(ExecutionPolicy exec, Iterator1 keys_first, Iterator1 keys_last, Iterator2 values_first, Iterator3 is_supported)
+void stable_sort_by_key_kernel(ExecutionPolicy exec, Iterator1 keys_first, Iterator1 keys_last, Iterator2 values_first)
 {
-#if (__CUDA_ARCH__ >= 200)
-  *is_supported = true;
   thrust::stable_sort_by_key(exec, keys_first, keys_last, values_first);
-#else
-  *is_supported = false;
-#endif
 }
 
 
@@ -51,21 +46,16 @@ void TestPairStableSortByKeyDevice(ExecutionPolicy exec)
   thrust::device_vector<P>   d_pairs = h_pairs;
   thrust::device_vector<int> d_values = h_values;
 
-  thrust::device_vector<bool> is_supported(1);
-
   // sort on the device
-  stable_sort_by_key_kernel<<<1,1>>>(exec, d_pairs.begin(), d_pairs.end(), d_values.begin(), is_supported.begin());
+  stable_sort_by_key_kernel<<<1,1>>>(exec, d_pairs.begin(), d_pairs.end(), d_values.begin());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
-  if(is_supported[0])
-  {
-    // sort on the host
-    thrust::stable_sort_by_key(h_pairs.begin(), h_pairs.end(), h_values.begin());
+  // sort on the host
+  thrust::stable_sort_by_key(h_pairs.begin(), h_pairs.end(), h_values.begin());
 
-    ASSERT_EQUAL_QUIET(h_pairs,  d_pairs);
-    ASSERT_EQUAL(h_values, d_values);
-  }
+  ASSERT_EQUAL_QUIET(h_pairs,  d_pairs);
+  ASSERT_EQUAL(h_values, d_values);
 };
 
 

testing/cuda/partition.cu

@@ -286,16 +286,11 @@ void TestPartitionCopyStencilDeviceNoSync()
 DECLARE_UNITTEST(TestPartitionCopyStencilDeviceNoSync);
 
 
-template<typename ExecutionPolicy, typename Iterator1, typename Predicate, typename Iterator2, typename Iterator3>
+template<typename ExecutionPolicy, typename Iterator1, typename Predicate, typename Iterator2>
 __global__
-void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Predicate pred, Iterator2 result, Iterator3 is_supported)
+void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Predicate pred, Iterator2 result)
 {
-#if (__CUDA_ARCH__ >= 200)
-  *is_supported = true;
   *result = thrust::stable_partition(exec, first, last, pred);
-#else
-  *is_supported = false;
-#endif
 }
 
 
@@ -313,24 +308,20 @@ void TestStablePartitionDevice(ExecutionPolicy exec)
   data[4] = 2; 
 
   thrust::device_vector<iterator> result(1);
-  thrust::device_vector<bool> is_supported(1);
-
-  stable_partition_kernel<<<1,1>>>(exec, data.begin(), data.end(), is_even<T>(), result.begin(), is_supported.begin());
+
+  stable_partition_kernel<<<1,1>>>(exec, data.begin(), data.end(), is_even<T>(), result.begin());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
-  if(is_supported[0])
-  {
-    thrust::device_vector<T> ref(5);
-    ref[0] = 2;
-    ref[1] = 2;
-    ref[2] = 1;
-    ref[3] = 1;
-    ref[4] = 1;
+  thrust::device_vector<T> ref(5);
+  ref[0] = 2;
+  ref[1] = 2;
+  ref[2] = 1;
+  ref[3] = 1;
+  ref[4] = 1;
 
-    ASSERT_EQUAL(2, (iterator)result[0] - data.begin());
-    ASSERT_EQUAL(ref, data);
-  }
+  ASSERT_EQUAL(2, (iterator)result[0] - data.begin());
+  ASSERT_EQUAL(ref, data);
 }
 
 
@@ -355,16 +346,11 @@ void TestStablePartitionDeviceNoSync()
 DECLARE_UNITTEST(TestStablePartitionDeviceNoSync);
 
 
-template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Predicate, typename Iterator3, typename Iterator4>
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Predicate, typename Iterator3>
 __global__
-void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, Iterator3 result, Iterator4 is_supported)
+void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, Iterator3 result)
 {
-#if (__CUDA_ARCH__ >= 200)
-  *is_supported = true;
   *result = thrust::stable_partition(exec, first, last, stencil_first, pred);
-#else
-  *is_supported = false;
-#endif
 }
 
 
@@ -389,24 +375,20 @@ void TestStablePartitionStencilDevice(ExecutionPolicy exec)
   stencil[4] = 2; 
 
   thrust::device_vector<iterator> result(1);
-  thrust::device_vector<bool> is_supported(1);
-
-  stable_partition_kernel<<<1,1>>>(exec, data.begin(), data.end(), stencil.begin(), is_even<T>(), result.begin(), is_supported.begin());
+
+  stable_partition_kernel<<<1,1>>>(exec, data.begin(), data.end(), stencil.begin(), is_even<T>(), result.begin());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
-  if(is_supported[0])
-  {
-    thrust::device_vector<T> ref(5);
-    ref[0] = 1;
-    ref[1] = 1;
-    ref[2] = 0;
-    ref[3] = 0;
-    ref[4] = 0;
+  thrust::device_vector<T> ref(5);
+  ref[0] = 1;
+  ref[1] = 1;
+  ref[2] = 0;
+  ref[3] = 0;
+  ref[4] = 0;
 
-    ASSERT_EQUAL(2, (iterator)result[0] - data.begin());
-    ASSERT_EQUAL(ref, data);
-  }
+  ASSERT_EQUAL(2, (iterator)result[0] - data.begin());
+  ASSERT_EQUAL(ref, data);
 }
 
 

testing/cuda/sort.cu

@@ -4,16 +4,11 @@
 #include <thrust/execution_policy.h>
 
 
-template<typename ExecutionPolicy, typename Iterator, typename Compare, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator, typename Compare>
 __global__
-void sort_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Compare comp, Iterator2 is_supported)
+void sort_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Compare comp)
 {
-#if (__CUDA_ARCH__ >= 200)
-  *is_supported = true;
   thrust::sort(exec, first, last, comp);
-#else
-  *is_supported = false;
-#endif
 }
 
 
@@ -34,19 +29,14 @@ void TestComparisonSortDevice(ExecutionPolicy exec, const size_t n, Compare comp
   thrust::host_vector<T>   h_data = unittest::random_integers<T>(n);
   thrust::device_vector<T> d_data = h_data;
 
-  thrust::device_vector<bool> is_supported(1);
-
-  sort_kernel<<<1,1>>>(exec, d_data.begin(), d_data.end(), comp, is_supported.begin());
+  sort_kernel<<<1,1>>>(exec, d_data.begin(), d_data.end(), comp);
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
 
-  if(is_supported[0])
-  {
-    thrust::sort(h_data.begin(), h_data.end(), comp);
-
-    ASSERT_EQUAL(h_data, d_data);
-  }
+  thrust::sort(h_data.begin(), h_data.end(), comp);
+
+  ASSERT_EQUAL(h_data, d_data);
 };
 
 
@@ -163,7 +153,7 @@ void TestComparisonSortCudaStreams()
   cudaStreamSynchronize(s);
 
   ASSERT_EQUAL(true, thrust::is_sorted(keys.begin(), keys.end(), my_less<int>()));
-                      
+
   cudaStreamDestroy(s);
 }
 DECLARE_UNITTEST(TestComparisonSortCudaStreams);