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

[Artem Belevich via Phabricator via cfe-commits]

tra added inline comments.


================
Comment at: clang/lib/Headers/__clang_hip_runtime_wrapper.h:73-74
+#define __HOST_DEVICE__                                                        \
+  static __host__ __device__ inline __attribute__((always_inline))
+__HOST_DEVICE__ double _Cosh(double x, double y) { return cosh(x) * y; }
+__HOST_DEVICE__ float _FCosh(float x, float y) { return coshf(x) * y; }
----------------
hliao wrote:
> tra wrote:
> > hliao wrote:
> > > tra wrote:
> > > > 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. 
> > > > 
> > > 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.
> > 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.
> 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.
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.



Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D93638/new/

https://reviews.llvm.org/D93638