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); +} + _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
