From 69079ebd163d6a98d09eb3d07d41c31598c3377f Mon Sep 17 00:00:00 2001
From: Allison Vacanti <alliepiper16@gmail.com>
Date: Mon, 25 Oct 2021 17:29:44 -0400
Subject: [PATCH] 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               | 12 +++++++-----
 test/bfloat16.h                    |  2 +-
 test/half.h                        |  4 ++--
 test/test_device_segmented_sort.cu | 14 +++++++-------
 4 files changed, 17 insertions(+), 15 deletions(-)

diff --git a/cmake/header_test.in b/cmake/header_test.in
index e851b07355..660ba40b7c 100644
--- a/cmake/header_test.in
+++ b/cmake/header_test.in
@@ -31,17 +31,19 @@
 
 // Hacky way to build a string, but it works on all tested platforms.
 #define CUB_MACRO_CHECK(MACRO, HEADER)                                         \
-  CUB_MACRO_CHECK_IMPL(Identifier MACRO should not be used from Thrust         \
-                            headers due to conflicts with HEADER.)
+  CUB_MACRO_CHECK_IMPL(Identifier MACRO should not be used from CUB            \
+                       headers due to conflicts with HEADER macros.)
 
 // complex.h conflicts
 #define I CUB_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 CUB_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(...) CUB_MACRO_CHECK('min', windows.h)
 //#define max(...) CUB_MACRO_CHECK('max', windows.h)
-#define small CUB_MACRO_CHECK('small', windows.h)
 
 #include <cub/${header}>
diff --git a/test/bfloat16.h b/test/bfloat16.h
index 671dbcdf81..aa9474dfdd 100644
--- a/test/bfloat16.h
+++ b/test/bfloat16.h
@@ -193,7 +193,7 @@ struct bfloat16_t
 
     /// numeric_traits<bfloat16_t>::max
     __host__ __device__ __forceinline__
-    static bfloat16_t max() {
+    static bfloat16_t (max)() {
         uint16_t max_word = 0x7F7F;
         return reinterpret_cast<bfloat16_t&>(max_word);
     }
diff --git a/test/half.h b/test/half.h
index dbec425d7a..d390a9ecf5 100644
--- a/test/half.h
+++ b/test/half.h
@@ -275,7 +275,7 @@ struct half_t
 
     /// numeric_traits<half_t>::max
     __host__ __device__ __forceinline__
-    static half_t max() {
+    static half_t (max)() {
         uint16_t max_word = 0x7BFF;
         return reinterpret_cast<half_t&>(max_word);
     }
@@ -315,7 +315,7 @@ std::ostream& operator<<(std::ostream &out, const __half &x)
 template <>
 struct CUB_NS_QUALIFIER::FpLimits<half_t>
 {
-    static __host__ __device__ __forceinline__ half_t Max() { return half_t::max(); }
+    static __host__ __device__ __forceinline__ half_t Max() { return (half_t::max)(); }
 
     static __host__ __device__ __forceinline__ half_t Lowest() { return half_t::lowest(); }
 };
diff --git a/test/test_device_segmented_sort.cu b/test/test_device_segmented_sort.cu
index bf2e8f4c62..5360279c81 100644
--- a/test/test_device_segmented_sort.cu
+++ b/test/test_device_segmented_sort.cu
@@ -419,7 +419,7 @@ public:
   InputDescription& add(const SizeGroupDescription &group)
   {
     if (static_cast<std::size_t>(group.segment_size) <
-        static_cast<std::size_t>(std::numeric_limits<KeyT>::max()))
+        static_cast<std::size_t>((std::numeric_limits<KeyT>::max)()))
     {
       for (int i = 0; i < group.segments; i++)
       {
@@ -1244,8 +1244,8 @@ void RandomizeInput(thrust::host_vector<KeyT> &h_keys,
 {
   for (std::size_t i = 0; i < h_keys.size(); i++)
   {
-    h_keys[i] = RandomValue(std::numeric_limits<KeyT>::max());
-    h_values[i] = RandomValue(std::numeric_limits<ValueT>::max());
+    h_keys[i] = RandomValue((std::numeric_limits<KeyT>::max)());
+    h_values[i] = RandomValue((std::numeric_limits<ValueT>::max)());
   }
 }
 
@@ -1255,8 +1255,8 @@ void RandomizeInput(thrust::host_vector<half_t> &h_keys,
 {
   for (std::size_t i = 0; i < h_keys.size(); i++)
   {
-    h_keys[i] = RandomValue(std::numeric_limits<int>::max());
-    h_values[i] = RandomValue(std::numeric_limits<std::uint32_t>::max());
+    h_keys[i] = RandomValue((std::numeric_limits<int>::max)());
+    h_values[i] = RandomValue((std::numeric_limits<std::uint32_t>::max)());
   }
 }
 #endif
@@ -1267,8 +1267,8 @@ void RandomizeInput(thrust::host_vector<bfloat16_t> &h_keys,
 {
   for (std::size_t i = 0; i < h_keys.size(); i++)
   {
-    h_keys[i] = RandomValue(std::numeric_limits<int>::max());
-    h_values[i] = RandomValue(std::numeric_limits<std::uint32_t>::max());
+    h_keys[i] = RandomValue((std::numeric_limits<int>::max)());
+    h_values[i] = RandomValue((std::numeric_limits<std::uint32_t>::max)());
   }
 }
 #endif