[clang] [CUDA][HIP] Fix std::min in wrapper header (PR #93976)

[Yaxun Liu via cfe-commits]

https://github.com/yxsamliu created https://github.com/llvm/llvm-project/pull/93976

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: https://github.com/llvm/llvm-project/issues/93962

Fixes: https://github.com/ROCm/HIP/issues/3502

>From ac8100056c81d1b4d3d40c31574be93ca78cf80c Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 31 May 2024 11:33:32 -0400
Subject: [PATCH] [CUDA][HIP] Fix std::min in wrapper header

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: https://github.com/llvm/llvm-project/issues/93962

Fixes: https://github.com/ROCm/HIP/issues/3502
---
 clang/lib/Headers/cuda_wrappers/algorithm    |  2 +-
 clang/test/Headers/cuda_wrapper_algorithm.cu | 48 ++++++++++++++++++++
 2 files changed, 49 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/Headers/cuda_wrapper_algorithm.cu

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);
+}
+