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

[via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2024-06-03T11:06:44-04:00
New Revision: 987e1b2ae3348a86b3f625119184a0f431c33bc7

URL: https://github.com/llvm/llvm-project/commit/987e1b2ae3348a86b3f625119184a0f431c33bc7
DIFF: https://github.com/llvm/llvm-project/commit/987e1b2ae3348a86b3f625119184a0f431c33bc7.diff

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

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

Added: 
    clang/test/Headers/cuda_wrapper_algorithm.cu

Modified: 
    clang/lib/Headers/cuda_wrappers/algorithm

Removed: 
    


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