[CUDA][HIP] Fix std::min in wrapper header (#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: #93962

Fixes: ROCm/hip#3502

Author: yxsamliu

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")

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