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

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Wed Jan 6 11:09:10 PST 2021


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:
> > 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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D93638



More information about the cfe-commits mailing list