[CUDA][HIP] Fix std::min in wrapper header

[yxsamliu]

The std::min behaves like 'a<b?a:b', which does not match libstdc++/libc++ behavior like 'b<a?b:a' when input is NaN.

Make it consistent with libstdc++/libc++.

Fixes: #93962

Fixes: ROCm/HIP#3502

[llvmbot]

@llvm/pr-subscribers-backend-x86

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

The std::min behaves like 'a<b?a:b', which does not match libstdc++/libc++ behavior like 'b<a?b:a' when input is NaN.

Make it consistent with libstdc++/libc++.

Fixes: #93962

Fixes: ROCm/HIP#3502

---

Full diff: https://github.com/llvm/llvm-project/pull/93976.diff

2 Files Affected:

• (modified) clang/lib/Headers/cuda_wrappers/algorithm (+1-1)
• (added) clang/test/Headers/cuda_wrapper_algorithm.cu (+48)

diff --git a/clang/lib/Headers/cuda_wrappers/algorithm b/clang/lib/Headers/cuda_wrappers/algorithm
index f14a0b00bb046..3f59f28ae35b3 100644
--- a/clang/lib/Headers/cuda_wrappers/algorithm
+++ b/clang/lib/Headers/cuda_wrappers/algorithm
@@ -99,7 +99,7 @@ template <class __T>
 __attribute__((enable_if(true, "")))
 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
 min(const __T &__a, const __T &__b) {
-  return __a < __b ? __a : __b;
+  return __b < __a ? __b : __a;
 }
 
 #pragma pop_macro("_CPP14_CONSTEXPR")
diff --git a/clang/test/Headers/cuda_wrapper_algorithm.cu b/clang/test/Headers/cuda_wrapper_algorithm.cu
new file mode 100644
index 0000000000000..d514285f7e17b
--- /dev/null
+++ b/clang/test/Headers/cuda_wrapper_algorithm.cu
@@ -0,0 +1,48 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+
+// RUN: %clang_cc1 \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple x86_64-unknown-unknown \
+// RUN:   -emit-llvm %s -O1 -o - \
+// RUN:   | FileCheck %s
+
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+
+#include <algorithm>
+
+extern "C" bool cmp(double a, double b) { return a<b; }
+
+// CHECK-LABEL: @test_std_min(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_min() {
+  return std::min(__builtin_nan(""), 0.0);
+}
+
+// CHECK-LABEL: @test_std_min_cmp(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_min_cmp() {
+  return std::min(__builtin_nan(""), 0.0, cmp);
+}
+
+// CHECK-LABEL: @test_std_max(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_max() {
+  return std::max(__builtin_nan(""), 0.0);
+}
+
+// CHECK-LABEL: @test_std_max_cmp(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_max_cmp() {
+  return std::max(__builtin_nan(""), 0.0, cmp);
+}
+

[llvmbot]

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

The std::min behaves like 'a<b?a:b', which does not match libstdc++/libc++ behavior like 'b<a?b:a' when input is NaN.

Make it consistent with libstdc++/libc++.

Fixes: #93962

Fixes: ROCm/HIP#3502

---

Full diff: https://github.com/llvm/llvm-project/pull/93976.diff

2 Files Affected:

• (modified) clang/lib/Headers/cuda_wrappers/algorithm (+1-1)
• (added) clang/test/Headers/cuda_wrapper_algorithm.cu (+48)

diff --git a/clang/lib/Headers/cuda_wrappers/algorithm b/clang/lib/Headers/cuda_wrappers/algorithm
index f14a0b00bb046..3f59f28ae35b3 100644
--- a/clang/lib/Headers/cuda_wrappers/algorithm
+++ b/clang/lib/Headers/cuda_wrappers/algorithm
@@ -99,7 +99,7 @@ template <class __T>
 __attribute__((enable_if(true, "")))
 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
 min(const __T &__a, const __T &__b) {
-  return __a < __b ? __a : __b;
+  return __b < __a ? __b : __a;
 }
 
 #pragma pop_macro("_CPP14_CONSTEXPR")
diff --git a/clang/test/Headers/cuda_wrapper_algorithm.cu b/clang/test/Headers/cuda_wrapper_algorithm.cu
new file mode 100644
index 0000000000000..d514285f7e17b
--- /dev/null
+++ b/clang/test/Headers/cuda_wrapper_algorithm.cu
@@ -0,0 +1,48 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+
+// RUN: %clang_cc1 \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple x86_64-unknown-unknown \
+// RUN:   -emit-llvm %s -O1 -o - \
+// RUN:   | FileCheck %s
+
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+
+#include <algorithm>
+
+extern "C" bool cmp(double a, double b) { return a<b; }
+
+// CHECK-LABEL: @test_std_min(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_min() {
+  return std::min(__builtin_nan(""), 0.0);
+}
+
+// CHECK-LABEL: @test_std_min_cmp(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_min_cmp() {
+  return std::min(__builtin_nan(""), 0.0, cmp);
+}
+
+// CHECK-LABEL: @test_std_max(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_max() {
+  return std::max(__builtin_nan(""), 0.0);
+}
+
+// CHECK-LABEL: @test_std_max_cmp(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" double test_std_max_cmp() {
+  return std::max(__builtin_nan(""), 0.0, cmp);
+}
+

[Artem-B]

LGTM for following the standard c++ library.

Floating point is an endless stream of surprises.

E.g. https://pixorblog.wordpress.com/2016/06/27/some-remarks-about-minmax-functions/

For instance, min() is not commutative and is not equivatent to cmath::fmin().