Enable min/max macro checks and fix all detected.

We can stop defining NOMINMAX at this point, too.
NVIDIA · Oct 26, 2021 · 7502661 · 7502661
1 parent 851ae3a
commit 7502661
Show file tree

Hide file tree

Showing 32 changed files with 127 additions and 102 deletions.
diff --git a/cmake/ThrustBuildCompilerTargets.cmake b/cmake/ThrustBuildCompilerTargets.cmake
@@ -73,9 +73,6 @@ 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)
-
-    # "Oh right, this is Visual Studio."
-    list(APPEND cxx_compile_definitions "NOMINMAX")
   else()
     append_option_if_available("-Werror" cxx_compile_options)
     append_option_if_available("-Wall" cxx_compile_options)

diff --git a/cmake/header_test.in b/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
 

diff --git a/dependencies/cub b/dependencies/cub
diff --git a/examples/bounding_box.cu b/examples/bounding_box.cu
@@ -56,10 +56,10 @@ struct bbox_reduction : public thrust::binary_function<bbox,bbox,bbox>
   bbox operator()(bbox a, bbox b)
   {
     // lower left corner
-    point2d ll(thrust::min(a.lower_left.x, b.lower_left.x), thrust::min(a.lower_left.y, b.lower_left.y));
+    point2d ll((thrust::min)(a.lower_left.x, b.lower_left.x), (thrust::min)(a.lower_left.y, b.lower_left.y));
 
     // upper right corner
-    point2d ur(thrust::max(a.upper_right.x, b.upper_right.x), thrust::max(a.upper_right.y, b.upper_right.y));
+    point2d ur((thrust::max)(a.upper_right.x, b.upper_right.x), (thrust::max)(a.upper_right.y, b.upper_right.y));
 
     return bbox(ll, ur);
   }

diff --git a/examples/discrete_voronoi.cu b/examples/discrete_voronoi.cu
@@ -222,7 +222,7 @@ int main(void)
   seeds.swap(temp);
 
   //JFA : main loop with k=n/2, n/4, ..., 1
-  for(int k = thrust::max(m,n) / 2; k > 0; k /= 2)
+  for(int k = (thrust::max)(m,n) / 2; k > 0; k /= 2)
   {
     jfa(seeds,temp,k,m,n);
     seeds.swap(temp);

diff --git a/examples/minmax.cu b/examples/minmax.cu
@@ -46,8 +46,8 @@ struct minmax_binary_op
   minmax_pair<T> operator()(const minmax_pair<T>& x, const minmax_pair<T>& y) const
   {
     minmax_pair<T> result;
-    result.min_val = thrust::min(x.min_val, y.min_val);
-    result.max_val = thrust::max(x.max_val, y.max_val);
+    result.min_val = (thrust::min)(x.min_val, y.min_val);
+    result.max_val = (thrust::max)(x.max_val, y.max_val);
     return result;
   }
 };

diff --git a/examples/padded_grid_reduction.cu b/examples/padded_grid_reduction.cu
@@ -54,8 +54,8 @@ struct reduce_tuple :
     { 
       if(thrust::get<0>(t0) && thrust::get<0>(t1)) // both valid
         return Tuple(true, 
-            thrust::min(thrust::get<1>(t0), thrust::get<1>(t1)),
-            thrust::max(thrust::get<2>(t0), thrust::get<2>(t1)));
+            (thrust::min)(thrust::get<1>(t0), thrust::get<1>(t1)),
+            (thrust::max)(thrust::get<2>(t0), thrust::get<2>(t1)));
       else if (thrust::get<0>(t0))
         return t0;
       else if (thrust::get<0>(t1))

diff --git a/examples/set_operations.cu b/examples/set_operations.cu
@@ -73,7 +73,7 @@ template <typename Vector>
 void SetIntersection(const Vector& A, const Vector& B)
 {
   // intersection output is at most min(A.size(), B.size())
-  Vector C(thrust::min(A.size(), B.size()));
+  Vector C((thrust::min)(A.size(), B.size()));
 
   // set_union returns an iterator C_end denoting the end of input
   typename Vector::iterator C_end;

diff --git a/examples/summary_statistics.cu b/examples/summary_statistics.cu
@@ -31,8 +31,8 @@ struct summary_stats_data
     void initialize()
     {
       n = mean = M2 = M3 = M4 = 0;
-      min = std::numeric_limits<T>::max();
-      max = std::numeric_limits<T>::min();
+      min = (std::numeric_limits<T>::max)();
+      max = (std::numeric_limits<T>::min)();
     }
 
     T variance()   { return M2 / (n - 1); }
@@ -89,8 +89,8 @@ struct summary_stats_binary_op
 
         //Basic number of samples (n), min, and max
         result.n   = n;
-        result.min = thrust::min(x.min, y.min);
-        result.max = thrust::max(x.max, y.max);
+        result.min = (thrust::min)(x.min, y.min);
+        result.max = (thrust::max)(x.max, y.max);
 
         result.mean = x.mean + delta * y.n / n;
 

diff --git a/testing/counting_iterator.cu b/testing/counting_iterator.cu
@@ -220,7 +220,7 @@ void TestCountingIteratorDifference(void)
     typedef thrust::counting_iterator<thrust::detail::uint64_t> Iterator;
     typedef thrust::iterator_difference<Iterator>::type Difference;
 
-    Difference diff = std::numeric_limits<thrust::detail::uint32_t>::max() + 1;
+    Difference diff = (std::numeric_limits<thrust::detail::uint32_t>::max)() + 1;
 
     Iterator first(0);
     Iterator last = first + diff;

diff --git a/testing/cuda/fill.cu b/testing/cuda/fill.cu
@@ -17,39 +17,39 @@ void TestFillDevice(ExecutionPolicy exec, size_t n)
   thrust::host_vector<T>   h_data = unittest::random_integers<T>(n);
   thrust::device_vector<T> d_data = h_data;
 
-  thrust::fill(h_data.begin() + std::min((size_t)1, n), h_data.begin() + std::min((size_t)3, n), (T) 0);
+  thrust::fill(h_data.begin() + (std::min)((size_t)1, n), h_data.begin() + (std::min)((size_t)3, n), (T) 0);
 
-  fill_kernel<<<1,1>>>(exec, d_data.begin() + std::min((size_t)1, n), d_data.begin() + std::min((size_t)3, n), (T) 0);
+  fill_kernel<<<1,1>>>(exec, d_data.begin() + (std::min)((size_t)1, n), d_data.begin() + (std::min)((size_t)3, n), (T) 0);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  thrust::fill(h_data.begin() + std::min((size_t)117, n), h_data.begin() + std::min((size_t)367, n), (T) 1);
+  thrust::fill(h_data.begin() + (std::min)((size_t)117, n), h_data.begin() + (std::min)((size_t)367, n), (T) 1);
 
-  fill_kernel<<<1,1>>>(exec, d_data.begin() + std::min((size_t)117, n), d_data.begin() + std::min((size_t)367, n), (T) 1);
+  fill_kernel<<<1,1>>>(exec, d_data.begin() + (std::min)((size_t)117, n), d_data.begin() + (std::min)((size_t)367, n), (T) 1);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  thrust::fill(h_data.begin() + std::min((size_t)8, n), h_data.begin() + std::min((size_t)259, n), (T) 2);
+  thrust::fill(h_data.begin() + (std::min)((size_t)8, n), h_data.begin() + (std::min)((size_t)259, n), (T) 2);
 
-  fill_kernel<<<1,1>>>(exec, d_data.begin() + std::min((size_t)8, n), d_data.begin() + std::min((size_t)259, n), (T) 2);
+  fill_kernel<<<1,1>>>(exec, d_data.begin() + (std::min)((size_t)8, n), d_data.begin() + (std::min)((size_t)259, n), (T) 2);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  thrust::fill(h_data.begin() + std::min((size_t)3, n), h_data.end(), (T) 3);
+  thrust::fill(h_data.begin() + (std::min)((size_t)3, n), h_data.end(), (T) 3);
 
-  fill_kernel<<<1,1>>>(exec, d_data.begin() + std::min((size_t)3, n), d_data.end(), (T) 3);
+  fill_kernel<<<1,1>>>(exec, d_data.begin() + (std::min)((size_t)3, n), d_data.end(), (T) 3);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -97,43 +97,43 @@ void TestFillNDevice(ExecutionPolicy exec, size_t n)
   thrust::host_vector<T>   h_data = unittest::random_integers<T>(n);
   thrust::device_vector<T> d_data = h_data;
 
-  size_t begin_offset = std::min<size_t>(1,n);
+  size_t begin_offset = (std::min)<size_t>(1,n);
 
-  thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)3, n) - begin_offset, (T) 0);
+  thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)3, n) - begin_offset, (T) 0);
 
-  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, std::min((size_t)3, n) - begin_offset, (T) 0);
+  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, (std::min)((size_t)3, n) - begin_offset, (T) 0);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  begin_offset = std::min<size_t>(117, n);
+  begin_offset = (std::min)<size_t>(117, n);
 
-  thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)367, n) - begin_offset, (T) 1);
+  thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)367, n) - begin_offset, (T) 1);
 
-  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, std::min((size_t)367, n) - begin_offset, (T) 1);
+  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, (std::min)((size_t)367, n) - begin_offset, (T) 1);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  begin_offset = std::min<size_t>(8, n);
+  begin_offset = (std::min)<size_t>(8, n);
 
-  thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)259, n) - begin_offset, (T) 2);
+  thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)259, n) - begin_offset, (T) 2);
 
-  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, std::min((size_t)259, n) - begin_offset, (T) 2);
+  fill_n_kernel<<<1,1>>>(exec, d_data.begin() + begin_offset, (std::min)((size_t)259, n) - begin_offset, (T) 2);
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
   }
 
   ASSERT_EQUAL(h_data, d_data);
 
-  begin_offset = std::min<size_t>(3, n);
+  begin_offset = (std::min)<size_t>(3, n);
 
   thrust::fill_n(h_data.begin() + begin_offset, h_data.size() - begin_offset, (T) 3);
 

diff --git a/testing/cuda/for_each.cu b/testing/cuda/for_each.cu
@@ -69,7 +69,7 @@ __global__ void for_each_kernel(ExecutionPolicy exec, Iterator first, Iterator l
 template<typename T>
 void TestForEachDeviceSeq(const size_t n)
 {
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
 
@@ -100,7 +100,7 @@ DECLARE_VARIABLE_UNITTEST(TestForEachDeviceSeq);
 template<typename T>
 void TestForEachDeviceDevice(const size_t n)
 {
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
 
@@ -145,7 +145,7 @@ void for_each_n_kernel(ExecutionPolicy exec, Iterator first, Size n, Function f)
 template<typename T>
 void TestForEachNDeviceSeq(const size_t n)
 {
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
 
@@ -176,7 +176,7 @@ DECLARE_VARIABLE_UNITTEST(TestForEachNDeviceSeq);
 template<typename T>
 void TestForEachNDeviceDevice(const size_t n)
 {
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
 

diff --git a/testing/cuda/gather.cu b/testing/cuda/gather.cu
@@ -14,7 +14,7 @@ void gather_kernel(ExecutionPolicy exec, Iterator1 map_first, Iterator1 map_last
 template<typename T, typename ExecutionPolicy>
 void TestGatherDevice(ExecutionPolicy exec, const size_t n)
 {
-  const size_t source_size = std::min((size_t) 10, 2 * n);
+  const size_t source_size = (std::min)((size_t) 10, 2 * n);
 
   // source vectors to gather from
   thrust::host_vector<T>   h_source = unittest::random_samples<T>(source_size);
@@ -107,7 +107,7 @@ struct is_even_gather_if
 template<typename T, typename ExecutionPolicy>
 void TestGatherIfDevice(ExecutionPolicy exec, const size_t n)
 {
-  const size_t source_size = std::min((size_t) 10, 2 * n);
+  const size_t source_size = (std::min)((size_t) 10, 2 * n);
 
   // source vectors to gather from
   thrust::host_vector<T>   h_source = unittest::random_samples<T>(source_size);

diff --git a/testing/cuda/scatter.cu b/testing/cuda/scatter.cu
@@ -15,7 +15,7 @@ template<typename ExecutionPolicy>
 void TestScatterDevice(ExecutionPolicy exec)
 {
   size_t n = 1000;
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<int> h_input(n, 1);
   thrust::device_vector<int> d_input(n, 1);
@@ -73,7 +73,7 @@ template<typename ExecutionPolicy>
 void TestScatterIfDevice(ExecutionPolicy exec)
 {
   size_t n = 1000;
-  const size_t output_size = std::min((size_t) 10, 2 * n);
+  const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
   thrust::host_vector<int> h_input(n, 1);
   thrust::device_vector<int> d_input(n, 1);

diff --git a/testing/fill.cu b/testing/fill.cu
@@ -92,23 +92,23 @@ void TestFill(size_t n)
     thrust::host_vector<T>   h_data = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_data = h_data;
 
-    thrust::fill(h_data.begin() + std::min((size_t)1, n), h_data.begin() + std::min((size_t)3, n), (T) 0);
-    thrust::fill(d_data.begin() + std::min((size_t)1, n), d_data.begin() + std::min((size_t)3, n), (T) 0);
+    thrust::fill(h_data.begin() + (std::min)((size_t)1, n), h_data.begin() + (std::min)((size_t)3, n), (T) 0);
+    thrust::fill(d_data.begin() + (std::min)((size_t)1, n), d_data.begin() + (std::min)((size_t)3, n), (T) 0);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    thrust::fill(h_data.begin() + std::min((size_t)117, n), h_data.begin() + std::min((size_t)367, n), (T) 1);
-    thrust::fill(d_data.begin() + std::min((size_t)117, n), d_data.begin() + std::min((size_t)367, n), (T) 1);
+    thrust::fill(h_data.begin() + (std::min)((size_t)117, n), h_data.begin() + (std::min)((size_t)367, n), (T) 1);
+    thrust::fill(d_data.begin() + (std::min)((size_t)117, n), d_data.begin() + (std::min)((size_t)367, n), (T) 1);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    thrust::fill(h_data.begin() + std::min((size_t)8, n), h_data.begin() + std::min((size_t)259, n), (T) 2);
-    thrust::fill(d_data.begin() + std::min((size_t)8, n), d_data.begin() + std::min((size_t)259, n), (T) 2);
+    thrust::fill(h_data.begin() + (std::min)((size_t)8, n), h_data.begin() + (std::min)((size_t)259, n), (T) 2);
+    thrust::fill(d_data.begin() + (std::min)((size_t)8, n), d_data.begin() + (std::min)((size_t)259, n), (T) 2);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    thrust::fill(h_data.begin() + std::min((size_t)3, n), h_data.end(), (T) 3);
-    thrust::fill(d_data.begin() + std::min((size_t)3, n), d_data.end(), (T) 3);
+    thrust::fill(h_data.begin() + (std::min)((size_t)3, n), h_data.end(), (T) 3);
+    thrust::fill(d_data.begin() + (std::min)((size_t)3, n), d_data.end(), (T) 3);
 
     ASSERT_EQUAL(h_data, d_data);
 
@@ -216,25 +216,25 @@ void TestFillN(size_t n)
     thrust::host_vector<T>   h_data = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_data = h_data;
 
-    size_t begin_offset = std::min<size_t>(1,n);
-    thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)3, n) - begin_offset, (T) 0);
-    thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)3, n) - begin_offset, (T) 0);
+    size_t begin_offset = (std::min)<size_t>(1,n);
+    thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)3, n) - begin_offset, (T) 0);
+    thrust::fill_n(d_data.begin() + begin_offset, (std::min)((size_t)3, n) - begin_offset, (T) 0);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    begin_offset = std::min<size_t>(117, n);
-    thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)367, n) - begin_offset, (T) 1);
-    thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)367, n) - begin_offset, (T) 1);
+    begin_offset = (std::min)<size_t>(117, n);
+    thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)367, n) - begin_offset, (T) 1);
+    thrust::fill_n(d_data.begin() + begin_offset, (std::min)((size_t)367, n) - begin_offset, (T) 1);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    begin_offset = std::min<size_t>(8, n);
-    thrust::fill_n(h_data.begin() + begin_offset, std::min((size_t)259, n) - begin_offset, (T) 2);
-    thrust::fill_n(d_data.begin() + begin_offset, std::min((size_t)259, n) - begin_offset, (T) 2);
+    begin_offset = (std::min)<size_t>(8, n);
+    thrust::fill_n(h_data.begin() + begin_offset, (std::min)((size_t)259, n) - begin_offset, (T) 2);
+    thrust::fill_n(d_data.begin() + begin_offset, (std::min)((size_t)259, n) - begin_offset, (T) 2);
 
     ASSERT_EQUAL(h_data, d_data);
 
-    begin_offset = std::min<size_t>(3, n);
+    begin_offset = (std::min)<size_t>(3, n);
     thrust::fill_n(h_data.begin() + begin_offset, h_data.size() - begin_offset, (T) 3);
     thrust::fill_n(d_data.begin() + begin_offset, d_data.size() - begin_offset, (T) 3);
 

diff --git a/testing/for_each.cu b/testing/for_each.cu
@@ -194,7 +194,7 @@ DECLARE_UNITTEST(TestForEachNSimpleAnySystem);
 template <typename T>
 void TestForEach(const size_t n)
 {
-    const size_t output_size = std::min((size_t) 10, 2 * n);
+    const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
     thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
 
@@ -227,7 +227,7 @@ DECLARE_VARIABLE_UNITTEST(TestForEach);
 template <typename T>
 void TestForEachN(const size_t n)
 {
-    const size_t output_size = std::min((size_t) 10, 2 * n);
+    const size_t output_size = (std::min)((size_t) 10, 2 * n);
 
     thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
+11 −7		cmake/header_test.in
+1 −1		test/bfloat16.h
+2 −2		test/half.h
+5 −5		test/test_block_scan.cu
+7 −7		test/test_device_segmented_sort.cu