Enable min/max macro checks and fix all detected.

We can stop defining NOMINMAX at this point for header tests, too.
The definition is now used only for tests/examples.

Cleaned up the thrust.compiler_interface logic and split out a
thrust.private_compiler_interface to be used only for tests/examples.

cmake/ThrustBuildCompilerTargets.cmake

@@ -23,9 +23,14 @@
 #   unconditionally thrown exceptions.
 
 function(thrust_build_compiler_targets)
+  # Used to compile all headers, tests, and examples.
   set(cxx_compile_definitions)
   set(cxx_compile_options)
 
+  # Used to compile tests and examples.
+  set(private_cxx_compile_definitions)
+  set(private_cxx_compile_options)
+
   thrust_update_system_found_flags()
 
   if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
@@ -72,10 +77,10 @@ function(thrust_build_compiler_targets)
 
     # Some of the async tests require /bigobj to fit all their sections into the
     # object files:
-    append_option_if_available("/bigobj" cxx_compile_options)
+    append_option_if_available("/bigobj" private_cxx_compile_options)
 
     # "Oh right, this is Visual Studio."
-    list(APPEND cxx_compile_definitions "NOMINMAX")
+    list(APPEND private_cxx_compile_definitions "NOMINMAX")
   else()
     append_option_if_available("-Werror" cxx_compile_options)
     append_option_if_available("-Wall" cxx_compile_options)
@@ -120,44 +125,61 @@ function(thrust_build_compiler_targets)
     list(APPEND cxx_compile_options -cppsuffix=cu)
   endif()
 
-  add_library(thrust.compiler_interface INTERFACE)
-
-  foreach (cxx_option IN LISTS cxx_compile_options)
-    target_compile_options(thrust.compiler_interface INTERFACE
-      $<$<COMPILE_LANGUAGE:CXX>:${cxx_option}>
-      $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVCXX>>:${cxx_option}>
-      # Only use -Xcompiler with NVCC, not NVC++.
-      #
-      # CMake can't split genexs, so this can't be formatted better :(
-      # This is:
-      # if (using CUDA and CUDA_COMPILER is NVCC) add -Xcompiler=opt:
-      $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcompiler=${cxx_option}>
-    )
-  endforeach()
-
-  foreach (cxx_definition IN LISTS cxx_compile_definitions)
-    # Add these for both CUDA and CXX targets:
-    target_compile_definitions(thrust.compiler_interface INTERFACE
-      ${cxx_definition}
-    )
-  endforeach()
+  ##############################################################################
+  # Helper function that creates a new interface target with ${target_name},
+  # and adds all provided C++ compile options and definitions.
+  function(_thrust_create_compiler_interface target_name options definitions)
+    add_library(${target_name} INTERFACE)
+
+    foreach (option IN LISTS options)
+      target_compile_options(${target_name} INTERFACE
+        $<$<COMPILE_LANGUAGE:CXX>:${option}>
+
+        # Only use -Xcompiler with NVCC, not NVC++:
+        $<$<COMPILE_LANG_AND_ID:CUDA,NVCXX>:${option}>
+        $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=${option}>
+      )
+    endforeach()
+
+    foreach (definition IN LISTS definitions)
+      # Add these for both CUDA and CXX targets:
+      target_compile_definitions(${target_name} INTERFACE ${definition})
+    endforeach()
+  endfunction()
+
+  ##############################################################################
+  # thrust.compiler_interface is used for all Thrust targets.
+  _thrust_create_compiler_interface(thrust.compiler_interface
+    "${cxx_compile_options}"
+    "${cxx_compile_definitions}"
+  )
 
   # Display warning numbers from nvcc cudafe errors:
   target_compile_options(thrust.compiler_interface INTERFACE
-    # If using CUDA w/ NVCC...
-    $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcudafe=--display_error_number>
+    $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcudafe=--display_error_number>
   )
 
   # Tell NVCC to be quiet about deprecated GPU targets:
   target_compile_options(thrust.compiler_interface INTERFACE
-    # If using CUDA w/ NVCC...
-    $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Wno-deprecated-gpu-targets>
+    $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Wno-deprecated-gpu-targets>
+  )
+
+  ##############################################################################
+  # thrust.private_compiler_interface
+  # Used for only test and example targets.
+  _thrust_create_compiler_interface(thrust.private_compiler_interface
+    "${private_cxx_compile_options}"
+    "${private_cxx_compile_definitions}"
   )
 
+  ##############################################################################
+  # Other Misc Interface Targets:
+  ##############################################################################
+
   # This is kept separate for Github issue #1174.
   add_library(thrust.promote_cudafe_warnings INTERFACE)
   target_compile_options(thrust.promote_cudafe_warnings INTERFACE
-    $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcudafe=--promote_warnings>
+    $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcudafe=--promote_warnings>
   )
 
   # Some of our unit tests unconditionally throw exceptions, and compilers will
@@ -168,7 +190,7 @@ function(thrust_build_compiler_targets)
   if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
     target_compile_options(thrust.silence_unreachable_code_warnings INTERFACE
       $<$<COMPILE_LANGUAGE:CXX>:/wd4702>
-      $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcompiler=/wd4702>
+      $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=/wd4702>
     )
   endif()
 
@@ -183,11 +205,11 @@ function(thrust_build_compiler_targets)
     # THRUST_IF_CONSTEXPR to address these warnings.
     target_compile_options(thrust.compiler_interface_cpp11 INTERFACE
       $<$<COMPILE_LANGUAGE:CXX>:/wd4127>
-      $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcompiler=/wd4127>
+      $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=/wd4127>
     )
     target_compile_options(thrust.compiler_interface_cpp14 INTERFACE
       $<$<COMPILE_LANGUAGE:CXX>:/wd4127>
-      $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcompiler=/wd4127>
+      $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=/wd4127>
     )
   endif()
 

cmake/header_test.in

@@ -15,7 +15,10 @@
 // Hacky way to build a string, but it works on all tested platforms.
 #define THRUST_MACRO_CHECK(MACRO, HEADER)                                      \
   THRUST_MACRO_CHECK_IMPL(Identifier MACRO should not be used from Thrust      \
-                            headers due to conflicts with HEADER.)
+                          headers due to conflicts with HEADER macros.)
+#define THRUST_MACRO_FUNCTION_CHECK(MACRO, HEADER)                             \
+  THRUST_MACRO_CHECK_IMPL(Wrap MACRO in parentheses when used in Thrust        \
+                          headers to avoid conflicts with HEADER macros.)
 
 // Use raw platform checks instead of the THRUST_HOST_COMPILER macros since we
 // don't want to #include any headers other than the one being tested.
@@ -42,13 +45,14 @@
 #endif
 
 // complex.h conflicts
-#define I THRUST_MACRO_CHECK('I', complex.h)
+#define I THRUST_MACRO_CHECK("I", complex.h)
 
 // windows.h conflicts
-// Disabling for now; we use min/max in many places, but since most
-// projects build with NOMINMAX this doesn't seem to be high priority to fix.
-//#define min(...) THRUST_MACRO_CHECK('min', windows.h)
-//#define max(...) THRUST_MACRO_CHECK('max', windows.h)
+// Note that these are unlikely to be caught by builds against gcc's
+// stdlib, since their bits/c++config.h header unconditionally undefs both
+// min/max because of the windows.h macros. MSVC builds should catch this, tho.
+#define min(...) THRUST_MACRO_FUNCTION_CHECK("min", windows.h)
+#define max(...) THRUST_MACRO_FUNCTION_CHECK("max", windows.h)
 
 #endif // THRUST_IGNORE_MACRO_CHECKS
 

examples/CMakeLists.txt

@@ -92,7 +92,10 @@ function(thrust_add_example target_name_var example_name example_src thrust_targ
   set(example_meta_target thrust.all.example.${example_name})
 
   add_executable(${example_target} "${real_example_src}")
-  target_link_libraries(${example_target} ${thrust_target})
+  target_link_libraries(${example_target} PRIVATE
+    thrust.private_compiler_interface
+    ${thrust_target}
+  )
   target_include_directories(${example_target} PRIVATE "${Thrust_SOURCE_DIR}/examples")
   thrust_clone_target_properties(${example_target} ${thrust_target})
 

testing/CMakeLists.txt

@@ -75,7 +75,10 @@ function(thrust_add_test target_name_var test_name test_src thrust_target)
   set(test_meta_target thrust.all.test.${test_name})
 
   add_executable(${test_target} "${real_test_src}")
-  target_link_libraries(${test_target} PRIVATE ${config_framework_target})
+  target_link_libraries(${test_target} PRIVATE
+    ${config_framework_target}
+    thrust.private_compiler_interface
+  )
   target_include_directories(${test_target} PRIVATE "${Thrust_SOURCE_DIR}/testing")
   thrust_clone_target_properties(${test_target} ${thrust_target})
 

thrust/system/cuda/detail/extrema.h

@@ -265,7 +265,7 @@ namespace __extrema {
 
         // if not enough to fill the device with threadblocks
         // then fill the device with threadblocks
-        reduce_grid_size = static_cast<int>(min(num_tiles, static_cast<size_t>(reduce_device_occupancy)));
+        reduce_grid_size = static_cast<int>((min)(num_tiles, static_cast<size_t>(reduce_device_occupancy)));
 
         typedef AgentLauncher<__reduce::DrainAgent<Size> > drain_agent;
         AgentPlan drain_plan = drain_agent::get_plan();

thrust/system/cuda/detail/merge.h

@@ -170,7 +170,7 @@ namespace __merge {
       Size partition_idx = blockDim.x * blockIdx.x + threadIdx.x;
       if (partition_idx < num_partitions)
       {
-        Size partition_at = thrust::min(partition_idx * items_per_tile,
+        Size partition_at = (thrust::min)(partition_idx * items_per_tile,
                                         keys1_count + keys2_count);
         Size partition_diag = merge_path(keys1,
                                          keys2,
@@ -463,7 +463,7 @@ namespace __merge {
         Size partition_end = merge_partitions[tile_idx + 1];
 
         Size diag0 = ITEMS_PER_TILE * tile_idx;
-        Size diag1 = thrust::min(keys1_count + keys2_count, diag0 + ITEMS_PER_TILE);
+        Size diag1 = (thrust::min)(keys1_count + keys2_count, diag0 + ITEMS_PER_TILE);
 
         // compute bounding box for keys1 & keys2
         //

thrust/system/cuda/detail/reduce.h

@@ -808,7 +808,7 @@ namespace __reduce {
 
         // if not enough to fill the device with threadblocks
         // then fill the device with threadblocks
-        reduce_grid_size = static_cast<int>(min(num_tiles, static_cast<size_t>(reduce_device_occupancy)));
+        reduce_grid_size = static_cast<int>((min)(num_tiles, static_cast<size_t>(reduce_device_occupancy)));
 
         typedef AgentLauncher<DrainAgent<Size> > drain_agent;
         AgentPlan drain_plan = drain_agent::get_plan();

thrust/system/detail/sequential/stable_merge_sort.inl

@@ -97,7 +97,7 @@ void insertion_sort_each(RandomAccessIterator first,
   {
     for(; first < last; first += partition_size)
     {
-      RandomAccessIterator partition_last = thrust::min(last, first + partition_size);
+      RandomAccessIterator partition_last = (thrust::min)(last, first + partition_size);
 
       thrust::system::detail::sequential::insertion_sort(first, partition_last, comp);
     } // end for
@@ -120,7 +120,7 @@ void insertion_sort_each_by_key(RandomAccessIterator1 keys_first,
   {
     for(; keys_first < keys_last; keys_first += partition_size, values_first += partition_size)
     {
-      RandomAccessIterator1 keys_partition_last = thrust::min(keys_last, keys_first + partition_size);
+      RandomAccessIterator1 keys_partition_last = (thrust::min)(keys_last, keys_first + partition_size);
 
       thrust::system::detail::sequential::insertion_sort_by_key(keys_first, keys_partition_last, values_first, comp);
     } // end for
@@ -143,8 +143,8 @@ void merge_adjacent_partitions(sequential::execution_policy<DerivedPolicy> &exec
 {
   for(; first < last; first += 2 * partition_size, result += 2 * partition_size)
   {
-    RandomAccessIterator1 interval_middle = thrust::min(last, first + partition_size);
-    RandomAccessIterator1 interval_last   = thrust::min(last, interval_middle + partition_size);
+    RandomAccessIterator1 interval_middle = (thrust::min)(last, first + partition_size);
+    RandomAccessIterator1 interval_last   = (thrust::min)(last, interval_middle + partition_size);
 
     thrust::merge(exec,
                   first, interval_middle,
@@ -178,8 +178,8 @@ void merge_adjacent_partitions_by_key(sequential::execution_policy<DerivedPolicy
       keys_first < keys_last;
       keys_first += stride, values_first += stride, keys_result += stride, values_result += stride)
   {
-    RandomAccessIterator1 keys_interval_middle = thrust::min(keys_last, keys_first + partition_size);
-    RandomAccessIterator1 keys_interval_last   = thrust::min(keys_last, keys_interval_middle + partition_size);
+    RandomAccessIterator1 keys_interval_middle = (thrust::min)(keys_last, keys_first + partition_size);
+    RandomAccessIterator1 keys_interval_last   = (thrust::min)(keys_last, keys_interval_middle + partition_size);
 
     RandomAccessIterator2 values_first2 = values_first + (keys_interval_middle - keys_first);
 

thrust/system/tbb/detail/reduce_by_key.inl

@@ -198,7 +198,7 @@ template<typename Iterator1, typename Iterator2, typename Iterator3, typename It
     const size_type interval_idx = r.begin();
 
     const size_type offset_to_first = interval_size * interval_idx;
-    const size_type offset_to_last = thrust::min(n, offset_to_first + interval_size);
+    const size_type offset_to_last = (thrust::min)(n, offset_to_first + interval_size);
 
     Iterator1 my_keys_first     = keys_first    + offset_to_first;
     Iterator1 my_keys_last      = keys_first    + offset_to_last;

thrust/system/tbb/detail/reduce_intervals.h

@@ -64,7 +64,7 @@ template<typename RandomAccessIterator1, typename RandomAccessIterator2, typenam
     Size interval_idx = r.begin();
 
     Size offset_to_first = interval_size * interval_idx;
-    Size offset_to_last = thrust::min(n, offset_to_first + interval_size);
+    Size offset_to_last = (thrust::min)(n, offset_to_first + interval_size);
 
     RandomAccessIterator1 my_first = first + offset_to_first;
     RandomAccessIterator1 my_last  = first + offset_to_last;