ROCm / half

MIT License
11 stars 7 forks source link

[Issue]: inoptimal std::isfinite or `no matching function for call to 'isfinite'` #16

Closed AngryLoki closed 2 weeks ago

AngryLoki commented 2 months ago

Problem Description

Hi,

I noticed a few problems affecting ROCm/half and upstream (v2.2 incl. trunk), so I'll report in both places.

Both are related to std::isfinite. First problem is that it does not work with libc++. Here is a breakdown why: 1) Since C++11 std::isfinite is defined for integral types (having specific traits) 2) Used trait is std::is_arithmetic, it uses std::is_floating_point. Both traits have no specialization for half_float::half 3) What happens in gcc/libstdc++/msvc is that std::isfinite(half_value) does implicit conversion to float and then checks for "isfinite". But it does not work in libc++.

My current understanding is that it is a bug of libc++ (caused by misunderstanding of https://wg21.cmeerw.net/lwg/issue2086), so I reported it to https://github.com/llvm/llvm-project/issues/98816

Anyways, the second problem is that this conversion to float should not be called at all, because check could be much simpler:

namespace std
{
    inline bool isfinite(const half_float::half& h) {
        return half_float::isfinite(h);
    }
}

Here is link to conformance view: https://godbolt.org/z/hhWebb9os

It looks safe, it fixes both problems mentioned above. Could you add it to half.hpp?

This affects ROCm/miopen, which currently fails with libc++ with:

/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/main.cpp:111:19: note: in instantiation of member function 'PoolDriver<half_float::half, double>::PoolDriver' requested here
  111 |         drv = new PoolDriver<float16, double>();
      |                   ^
/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/../src/kernels/hip_float8.hpp:595:13: note: candidate function not viable: no known conversion from 'half_float::half' to 'miopen_f8::hip_f8<miopen_f8::hip_f8_type::fp8>' for 1st argument
  595 | inline bool isfinite(miopen_f8::hip_f8<miopen_f8::hip_f8_type::fp8> x) // NOLINT
      |             ^        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/../src/kernels/hip_float8.hpp:600:13: note: candidate function not viable: no known conversion from 'half_float::half' to 'miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8>' for 1st argument
  600 | inline bool isfinite(miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8> x) // NOLINT
      |             ^        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/usr/include/c++/v1/__math/traits.h:49:80: note: candidate template ignored: requirement 'is_arithmetic<half_float::half>::value' was not satisfied [with _A1 = half_float::half]
   49 | _LIBCPP_NODISCARD_EXT _LIBCPP_CONSTEXPR_SINCE_CXX23 _LIBCPP_HIDE_FROM_ABI bool isfinite(_A1 __x) _NOEXCEPT {
      |                                                                                ^
/usr/include/c++/v1/__math/traits.h:54:80: note: candidate template ignored: requirement 'is_arithmetic<half_float::half>::value' was not satisfied [with _A1 = half_float::half]
   54 | _LIBCPP_NODISCARD_EXT _LIBCPP_CONSTEXPR_SINCE_CXX23 _LIBCPP_HIDE_FROM_ABI bool isfinite(_A1) _NOEXCEPT {
      |                                                                                ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:509:23: note: candidate template ignored: substitution failure [with __T = half_float::half]: no type named 'type' in '__hip_enable_if<false, bool>'
  509 | __HIP_OVERLOAD1(bool, isfinite)
      | ~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:452:7: note: expanded from macro '__HIP_OVERLOAD1'
  451 |       typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type  \
      |                                                                          ~~~~
  452 |       __fn(__T __x) {                                                          \
      |       ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:98:31: note: candidate function not viable: call to __device__ function from __host__ function
   98 | __DEVICE__ __CONSTEXPR__ bool isfinite(float __x) { return ::__finitef(__x); }
      |                               ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:99:31: note: candidate function not viable: call to __device__ function from __host__ function
   99 | __DEVICE__ __CONSTEXPR__ bool isfinite(double __x) { return ::__finite(__x); }
      |                               ^
In file included from /var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/main.cpp:31:
In file included from /var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/conv_driver.hpp:40:
/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/mloConvHost.hpp:1144:37: error: no matching function for call to 'isfinite'
 1144 |                                 || !std::isfinite(g_val);
      |                                     ^~~~~~~~~~~~~
/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/../src/kernels/hip_float8.hpp:595:13: note: candidate function not viable: no known conversion from 'half_float::half' to 'miopen_f8::hip_f8<miopen_f8::hip_f8_type::fp8>' for 1st argument
  595 | inline bool isfinite(miopen_f8::hip_f8<miopen_f8::hip_f8_type::fp8> x) // NOLINT
      |             ^        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/sci-libs/miopen-6.1.1/work/MIOpen-rocm-6.1.1/driver/../src/kernels/hip_float8.hpp:600:13: note: candidate function not viable: no known conversion from 'half_float::half' to 'miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8>' for 1st argument
  600 | inline bool isfinite(miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8> x) // NOLINT
      |             ^        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/usr/include/c++/v1/__math/traits.h:49:80: note: candidate template ignored: requirement 'is_arithmetic<half_float::half>::value' was not satisfied [with _A1 = half_float::half]
   49 | _LIBCPP_NODISCARD_EXT _LIBCPP_CONSTEXPR_SINCE_CXX23 _LIBCPP_HIDE_FROM_ABI bool isfinite(_A1 __x) _NOEXCEPT {
      |                                                                                ^
/usr/include/c++/v1/__math/traits.h:54:80: note: candidate template ignored: requirement 'is_arithmetic<half_float::half>::value' was not satisfied [with _A1 = half_float::half]
   54 | _LIBCPP_NODISCARD_EXT _LIBCPP_CONSTEXPR_SINCE_CXX23 _LIBCPP_HIDE_FROM_ABI bool isfinite(_A1) _NOEXCEPT {
      |                                                                                ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:509:23: note: candidate template ignored: substitution failure [with __T = half_float::half]: no type named 'type' in '__hip_enable_if<false, bool>'
  509 | __HIP_OVERLOAD1(bool, isfinite)
      | ~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:452:7: note: expanded from macro '__HIP_OVERLOAD1'
  451 |       typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type  \
      |                                                                          ~~~~
  452 |       __fn(__T __x) {                                                          \
      |       ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:98:31: note: candidate function not viable: call to __device__ function from __host__ function
   98 | __DEVICE__ __CONSTEXPR__ bool isfinite(float __x) { return ::__finitef(__x); }
      |                               ^
/usr/lib/llvm/18/bin/../../../../lib/clang/18/include/__clang_hip_cmath.h:99:31: note: candidate function not viable: call to __device__ function from __host__ function
   99 | __DEVICE__ __CONSTEXPR__ bool isfinite(double __x) { return ::__finite(__x); }
      |                               ^

Operating System

Gentoo

CPU

-

GPU

AMD Instinct MI300X

ROCm Version

ROCm 6.1.0

ROCm Component

No response

Steps to Reproduce

No response

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

tcgu-amd commented 3 weeks ago

Hi @AngryLoki Thanks for reporting the issue! It seems like I am unable to reproduce your issue on ubuntu. Here are the steps I took:

int main(){ half_float::half v(0.2); std::isfinite(v); return 0; }

- Compile with hipcc and libc++
```bash
hipcc -O2 --stdlib=libc++ test.cpp

All I got was some warning regarding unused return value:

test.cpp:6:9: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result]
    6 |         std::isfinite(v);
      |         ^~~~~~~~~~~~~ ~
1 warning generated when compiling for gfx906.
test.cpp:6:9: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result]
    6 |         std::isfinite(v);
      |         ^~~~~~~~~~~~~ ~

Would you be able to double check on your end and see if the issue persists?

Also, please keep in mind that ROCm does not officially support Gentoo. I will label this issue with "enhancement" to reflect this properly.

Thanks :D

AngryLoki commented 2 weeks ago

Hi, sorry for the lack of updates, as multiple things changed since I reported about this issue:

1) bug with no matching function for call to 'isfinite' was fixed on llvm side - https://github.com/llvm/llvm-project/issues/98816 2) about second part (slow ADT-based code) developer in https://sourceforge.net/p/half/discussion/general/thread/7bee917723/#1046/0082 responded that adding std::isfinite specialization has some issues, so they won't add it.

Note that the part of inoptimal code (std::isfinite(fp16_value_implicitly_casted_to_fp32)) affects all platforms, both libc++/libstdc++. However I encountered this pattern only in miopen tests code, and there are no expectations for test code to be fast. And as https://github.com/ROCm/half/issues/8 plans to synchronize with upstream, while upstream has no plans to specialize isfinite, I am closing this task.