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