[hip] Enable HIP compilation with `<complex`> on MSVC.

[hliao]

Summary

• MSVC has different <complex> implementation which calls into functions declared in <ymath.h>. Instead of through builtins or inline functions, <ymath.h> functions are provided through external libraries. To get <complex> compiled with HIP on MSVC, we need to
  • Wrap <ymath.h> to force its functions to be __host__ __device__.
  • Provide its function definitions for the device compilation.

[hliao]

Beyond the enabling of the compilation with <complex> on Windows, I really have the concern on the current approach supporting <complex> compilation in the device compilation. The device compilation should not relies on the host STL implementation. That results in inconsistent compilation results across various platforms, especially Linux vs. Windows.
BTW, the use of <complex> in CUDA cannot be compiled with NVCC directly even with --expt-relaxed-constexpr, c.f. https://godbolt.org/z/3f79co

[hliao]

Fix typo.

[jdoerfert]

Disclaimer: I request changes because of the next sentence, other than that I have no objection but also cannot review this.
All cuda_wrapper headers say something about complex in the first row, copy & paste error. All have the wrong license text (I think).

---

The device compilation should not relies on the host STL implementation.

The OpenMP one doesn't, at least not as much as the CUDA one. Unsure how that works on windows though.

[hliao]

Fix typo again.

[hliao]

Fix license.

[hliao]

In D93638#2466097, @jdoerfert wrote:

Disclaimer: I request changes because of the next sentence, other than that I have no objection but also cannot review this.
All cuda_wrapper headers say something about complex in the first row, copy & paste error. All have the wrong license text (I think).

Fixed.

---

The device compilation should not relies on the host STL implementation.

The OpenMP one doesn't, at least not as much as the CUDA one. Unsure how that works on windows though.

This patch doesn't apply to OpenMP.

[jdoerfert]

In D93638#2466333, @hliao wrote:

In D93638#2466097, @jdoerfert wrote:

The device compilation should not relies on the host STL implementation.

The OpenMP one doesn't, at least not as much as the CUDA one. Unsure how that works on windows though.

This patch doesn't apply to OpenMP.

I'm aware.

[hliao]

These functions are pure C functions.

[hliao]

Fix the cmake to distribute that header wrapper.

[hliao]

PING

[yaxunl]

clang/lib/Headers/cuda_wrappers/ymath.h

16 ↗

(On Diff #313211)

I am wondering whether we only want to do this for windows, since ymath.h may be an ordinary users' header file on linux.

[tra]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

I don't think we want to provide a HD implementation.
This will potentially change the meaning of these functions on the host side vs what they do in a C++ compilation.
It should probably be just __device__.

Next question is -- do we want to provide the definitions, or would just declarations be sufficient?
In other words -- are these functions needed for some standard C++ library functionality that we expect to work on the GPU?
If it's just to aid with overload resolution, declarations should do.

[hliao]

Only mark HD attributes in ymath.h wrapper header when compiled with MSVC.

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

These functions are declared in ymath.h and, in the host compilation, are resolved by linking MSVC RT libraries. For the device function, as we already mark all prototypes in ymath.h as HD, we have to implement them as HD to match the declaration. But, those definitions should be only available in the device compilation to avoid conflicts in the host compilation.

[tra]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

You don't need the definitions to avoid conflicting declarations.

I'm still not convinced that HD is needed.
Did you try just making these declarations __device__ and remove the ymath.h wrapper?
Basically I'm trying to find the bare minimum we need to do to make it work.

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

We don't call these functions directly. They are called in MSVC's <complex>. As functions in <complex> are marked as HD, we need to mark these functions in ymath.h as HD to pass the compilation.

[tra]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

I assume that we're dealing with this file:
https://github.com/microsoft/STL/blob/master/stl/inc/ymath.h

I don't think we need a wrapper for ymath.
It may be sufficient to define or declare __device__ _Cosh() and other functions and let overload resolution pick the right function.
I think it would be a better approach than providing an inline __host__ __device__ definition for those functions and effectively replacing MSVC-provided host-side implementation of those functions.

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

ymath.h could be included before <complex>. That implies _Cosh could be declared as H only and results in the compilation failure.
BTW, I don't think replacing host-side implementation is a good idea as the host compilation should be kept consistent with the host compiler as much as possible.

[tra]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

How? Isn't clang_hip_runtime_wrapper.h included before anything in the user source file is processed? If the clang_hip_runtime_wrapper.h is not included, first, then the ymath.h wrapper will not work either as it needs __device__ macros.

replacing host-side implementation is a good idea

While consistency between host/device is good, we should not introduce a possible inconsistency between host-side TUs.
Considering vastly larger amounts of host-side code compiled as C++ (e.g. TF has way more C++ code than HIP/CUDA) and correspondingly more reliance on every possible detail of the implementation, I would err on the side of not changing host-side behavior in any way at all, if possible.

It's reasonably safe to add an overload (it may still be observable, but it's usually possible to add it in a way that does not affect the host). Replacing host-side things is more risky, as it greatly expands the opportunities for things to go wrong.

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

<ymath.h> is also included in other headers, which is not wrapped. If we don't wrap <ymath.h>, there's a chance that it's included as it is. That's why we have to wrap <ymath.h> to ensure all functions marked with HD. Do I miss anything?

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

I am wondering whether we could assume <ymath.h> is an internal header *only*.

[hliao]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

73–74 ↗

(On Diff #313211)

It's turned out that <ymath.h> (an internal header) is included in other headers, which is not wrapped like <complex>. The sequence including <ymath.h> using MSVC 2017 is from <algorithm>, <xmemory>, <xmemory0>, <limits>, and then <ymath.h>. As <algorithm> is included before <complex>, without wrapping <ymath.h>, we cannot overload _Cosh (pure C function.).

[hliao]

Forget that C function could be overloaded on Clang with overloadable
extension. With that, we don't need to mark functions from <ymath.h> as HD.
Instead, we could provide their device-side implementation directly.

[tra]

In D93638#2483913, @hliao wrote:

Forget that C function could be overloaded on Clang with overloadable
extension. With that, we don't need to mark functions from <ymath.h> as HD.
Instead, we could provide their device-side implementation directly.

Ah. It was the "C" part that I was missing. I was only thinking of C++/HIP/CUDA. I think we're on the same page now. LGTM.

[yaxunl]

LGTM. Thanks.

clang/lib/Headers/__clang_hip_cmath.h

 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
 
+// Define device-side math functions from <ymath.h> on MSVC.
+#if defined(_MSC_VER)
+#if defined(__cplusplus)
+extern "C" {
+#endif // defined(__cplusplus)
+__DEVICE__ __attribute__((overloadable)) double _Cosh(double x, double y) {
+  return cosh(x) * y;
+}
+__DEVICE__ __attribute__((overloadable)) float _FCosh(float x, float y) {
+  return coshf(x) * y;
+}
+__DEVICE__ __attribute__((overloadable)) short _Dtest(double *p) {
+  return fpclassify(*p);
+}
+__DEVICE__ __attribute__((overloadable)) short _FDtest(float *p) {
+  return fpclassify(*p);
+}
+__DEVICE__ __attribute__((overloadable)) double _Sinh(double x, double y) {
+  return sinh(x) * y;
+}
+__DEVICE__ __attribute__((overloadable)) float _FSinh(float x, float y) {
+  return sinhf(x) * y;
+}
+#if defined(__cplusplus)
+}
+#endif // defined(__cplusplus)
+#endif // defined(_MSC_VER)
+
 #pragma pop_macro("__DEVICE__")
 
 #endif // __CLANG_HIP_CMATH_H__