Fix some windows.h collisions with min/max.

Tried adding these to the header test macro checks, but this introduced
new issues on non-msvc builds. We can revist the header tests later,
this PR just fixes the collisions.

cmake/header_test.in

@@ -15,7 +15,7 @@
 // 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.)
 
 // 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.
@@ -45,11 +45,13 @@
 #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 small THRUST_MACRO_CHECK('small', windows.h)
+// We can't enable these checks without breaking some builds -- some standard
+// library implementations unconditionally `#undef` these macros, which then
+// causes random failures later.
+// Leaving these commented out as a warning: Here be dragons.
 //#define min(...) THRUST_MACRO_CHECK('min', windows.h)
 //#define max(...) THRUST_MACRO_CHECK('max', windows.h)
-#define small THRUST_MACRO_CHECK('small', windows.h)
 
 #endif // THRUST_IGNORE_MACRO_CHECKS
 

thrust/system/cuda/detail/extrema.h

@@ -268,7 +268,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;