Add ability to place Thrust in a custom namespace.

This provides a workaround for downstream projects that encounter
a variety of issues from dynamically linking multiple libraries that
use Thrust.

See the new `thrust/detail/config/namespace.h` header for details.

Added several tests and checks to validate that this behavior is correct,
and the `__THRUST_DEFINE_HAS_MEMBER_FUNCTION` utility has been rewritten
to WAR an nvcc bug when the old implementation was used with objects in
an anonymous namespace.

New tests:
  - testing/namespace_anonymous.cu
  - testing/namespace_wrapped.cu
  - testing/cmake/check_namespace.cmake

cmake/ThrustHeaderTesting.cmake

@@ -118,6 +118,11 @@ foreach(thrust_target IN LISTS THRUST_TARGETS)
   set(headertest_target ${config_prefix}.headers)
   add_library(${headertest_target} OBJECT ${headertest_srcs})
   target_link_libraries(${headertest_target} PUBLIC ${thrust_target})
+  # Wrap Thrust/CUB in a custom namespace to check proper use of ns macros:
+  target_compile_definitions(${headertest_target} PRIVATE
+    "THRUST_WRAPPED_NAMESPACE=wrapped_thrust"
+    "CUB_WRAPPED_NAMESPACE=wrapped_cub"
+  )
   thrust_clone_target_properties(${headertest_target} ${thrust_target})
 
   # Disable macro checks on TBB; the TBB atomic implementation uses `I` and

testing/cmake/CMakeLists.txt

@@ -26,3 +26,12 @@ if (THRUST_CPP_FOUND AND THRUST_CUDA_FOUND)
       ${extra_cmake_flags}
   )
 endif()
+
+# Check that namespace macros are used correctly:
+add_test(
+  NAME thrust.test.cmake.check_namespace
+  COMMAND
+    "${CMAKE_COMMAND}"
+      -D "Thrust_SOURCE_DIR=${Thrust_SOURCE_DIR}"
+      -P "${CMAKE_CURRENT_LIST_DIR}/check_namespace.cmake"
+)

testing/cmake/check_namespace.cmake

@@ -0,0 +1,88 @@
+# Check all files in thrust to make sure that they use
+# THRUST_NAMESPACE_BEGIN/END instead of bare `namespace thrust {}` declarations.
+#
+# This is run as a ctest test named `thrust.test.cmake.check_namespace`, or
+# manually with:
+# cmake -D "Thrust_SOURCE_DIR=<thrust project root>" -P check_namespace.cmake
+
+cmake_minimum_required(VERSION 3.15)
+
+set(exclusions
+  # This defines the macros and must have bare namespace declarations:
+  thrust/detail/config/namespace.h
+)
+
+function(count_substrings input search_regex output_var)
+  string(REGEX MATCHALL "${search_regex}" matches "${input}")
+  list(LENGTH matches num_matches)
+  set(${output_var} ${num_matches} PARENT_SCOPE)
+endfunction()
+
+set(bare_ns_regex "namespace[ \n\r\t]+thrust[ \n\r\t]*\\{")
+
+# Validation check for the above regex:
+count_substrings([=[
+namespace thrust{
+namespace thrust {
+namespace  thrust  {
+ namespace thrust {
+namespace thrust
+{
+namespace
+thrust
+{
+]=]
+  ${bare_ns_regex} valid_count)
+if (NOT valid_count EQUAL 6)
+  message(FATAL_ERROR "Validation of bare namespace regex failed: "
+                      "Matched ${valid_count} times, expected 6.")
+endif()
+
+set(exit_code 0)
+file(GLOB_RECURSE thrust_srcs
+  RELATIVE "${Thrust_SOURCE_DIR}"
+  "${Thrust_SOURCE_DIR}/thrust/*.h"
+  "${Thrust_SOURCE_DIR}/thrust/*.inl"
+  "${Thrust_SOURCE_DIR}/testing/*.cu"
+  "${Thrust_SOURCE_DIR}/examples/*.cu"
+)
+
+foreach(src ${thrust_srcs})
+  if (${src} IN_LIST exclusions)
+    continue()
+  endif()
+
+  file(READ "${Thrust_SOURCE_DIR}/${src}" src_contents)
+
+  count_substrings("${src_contents}" "${bare_ns_regex}" bare_ns_count)
+  count_substrings("${src_contents}" THRUST_NS_PREFIX prefix_count)
+  count_substrings("${src_contents}" THRUST_NS_POSTFIX postfix_count)
+  count_substrings("${src_contents}" THRUST_NAMESPACE_BEGIN begin_count)
+  count_substrings("${src_contents}" THRUST_NAMESPACE_END end_count)
+
+  if (NOT bare_ns_count EQUAL 0)
+    message("'${src}' contains 'namespace thrust {...}'. Replace with THRUST_NAMESPACE macros.")
+    set(exit_code 1)
+  endif()
+
+  if (NOT prefix_count EQUAL 0)
+    message("'${src}' contains 'THRUST_NS_PREFIX'. Replace with THRUST_NAMESPACE macros.")
+    set(exit_code 1)
+  endif()
+
+  if (NOT postfix_count EQUAL 0)
+    message("'${src}' contains 'THRUST_NS_POSTFIX'. Replace with THRUST_NAMESPACE macros.")
+    set(exit_code 1)
+  endif()
+
+  if (NOT begin_count EQUAL end_count)
+    message("'${src}' namespace macros are unbalanced:")
+    message(" - THRUST_NAMESPACE_BEGIN occurs ${begin_count} times.")
+    message(" - THRUST_NAMESPACE_END   occurs ${end_count} times.")
+    set(exit_code 1)
+  endif()
+endforeach()
+
+if (NOT exit_code EQUAL 0)
+  message(FATAL_ERROR "Errors detected.")
+endif()

testing/copy.cu

@@ -733,15 +733,14 @@ struct only_set_when_expected_it
     }
 };
 
-namespace thrust
-{
+THRUST_NAMESPACE_BEGIN
 namespace detail
 {
 // We need this type to pass as a non-const ref for unary_transform_functor
 // to compile:
 template <>
 struct is_non_const_reference<only_set_when_expected_it> : thrust::true_type {};
-}
+} // end namespace detail
 
 template<>
 struct iterator_traits<only_set_when_expected_it>
@@ -750,7 +749,7 @@ struct iterator_traits<only_set_when_expected_it>
     typedef only_set_when_expected_it reference;
     typedef thrust::random_access_device_iterator_tag iterator_category;
 };
-}
+THRUST_NAMESPACE_END
 
 void TestCopyWithBigIndexesHelper(int magnitude)
 {

testing/mr_disjoint_pool.cu

@@ -32,7 +32,8 @@ struct alloc_id
     }
 };
 
-namespace thrust { namespace detail {
+THRUST_NAMESPACE_BEGIN
+namespace detail {
 template<>
 struct pointer_traits<alloc_id>
 {
@@ -48,7 +49,10 @@ struct pointer_traits<alloc_id>
         return reinterpret_cast<void *>(id.alignment);
     }
 };
-}}
+
+} // end namespace detail
+
+THRUST_NAMESPACE_END
 
 class dummy_resource final : public thrust::mr::memory_resource<alloc_id>
 {

testing/namespace_anonymous.cu

@@ -0,0 +1,35 @@
+// Put all of Thrust and CUB inside anonymous namespaces.
+#define THRUST_CUB_USE_ANON_NAMESPACE
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/transform.h>
+
+#include <unittest/unittest.h>
+
+// Test that we can use a few common utilities and algorithms from an anon
+// namespace at runtime. More extensive testing is performed by the header
+// tests and the check_namespace.cmake test.
+void TestAnonymousNamespace()
+{
+  const std::size_t n = 2048;
+  const auto in_1_begin = thrust::make_constant_iterator<int>(12);
+  const auto in_2_begin = thrust::make_counting_iterator<int>(1024);
+  // Check that the qualifier resolves properly:
+  THRUST_NS_QUALIFIER::device_vector<int> d_out(n);
+
+  thrust::transform(in_1_begin, in_1_begin + n,
+                    in_2_begin,
+                    d_out.begin(),
+                    thrust::plus<>{});
+
+  thrust::host_vector<int> h_out(d_out);
+
+  for (std::size_t i = 0; i < n; ++i)
+  {
+    ASSERT_EQUAL(h_out[i], static_cast<int>(i) + 1024 + 12);
+  }
+}
+DECLARE_UNITTEST(TestAnonymousNamespace);

testing/namespace_wrapped.cu

@@ -0,0 +1,43 @@
+// Wrap thrust and cub in different enclosing namespaces
+// (In practice, you probably want these to be the same, in which case just
+// set THRUST_CUB_WRAPPED_NAMESPACE to set both).
+#define THRUST_WRAPPED_NAMESPACE wrap_thrust
+#define CUB_WRAPPED_NAMESPACE    wrap_cub
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/transform.h>
+
+#include <unittest/unittest.h>
+
+// Test that we can use a few common utilities and algorithms from a wrapped
+// namespace at runtime. More extensive testing is performed by the header
+// tests and the check_namespace.cmake test.
+void TestWrappedNamespace()
+{
+  const std::size_t n = 2048;
+
+  const auto in_1_begin =
+    ::wrap_thrust::thrust::make_constant_iterator<int>(12);
+  const auto in_2_begin =
+    ::wrap_thrust::thrust::make_counting_iterator<int>(1024);
+
+  // Check that the qualifier resolves properly:
+  THRUST_NS_QUALIFIER::device_vector<int> d_out(n);
+
+  ::wrap_thrust::thrust::transform(in_1_begin,
+                                   in_1_begin + n,
+                                   in_2_begin,
+                                   d_out.begin(),
+                                   ::wrap_thrust::thrust::plus<>{});
+
+  ::wrap_thrust::thrust::host_vector<int> h_out(d_out);
+
+  for (std::size_t i = 0; i < n; ++i)
+  {
+    ASSERT_EQUAL(h_out[i], static_cast<int>(i) + 1024 + 12);
+  }
+}
+DECLARE_UNITTEST(TestWrappedNamespace);

testing/scan.cu

@@ -583,15 +583,14 @@ struct only_set_when_expected_it
     }
 };
 
-namespace thrust
-{
+THRUST_NAMESPACE_BEGIN
 template<>
 struct iterator_traits<only_set_when_expected_it>
 {
     typedef long long value_type;
     typedef only_set_when_expected_it reference;
 };
-}
+THRUST_NAMESPACE_END
 
 void TestInclusiveScanWithBigIndexesHelper(int magnitude)
 {