[PATCH] D61399: [OpenMP][Clang] Support for target math functions

Alexey Bataev via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu May 2 12:25:01 PDT 2019


ABataev added a comment.

In D61399#1488329 <https://reviews.llvm.org/D61399#1488329>, @hfinkel wrote:

> In D61399#1488309 <https://reviews.llvm.org/D61399#1488309>, @ABataev wrote:
>
> > In D61399#1488299 <https://reviews.llvm.org/D61399#1488299>, @hfinkel wrote:
> >
> > > In D61399#1488262 <https://reviews.llvm.org/D61399#1488262>, @ABataev wrote:
> > >
> > > > I don't like this implementation. Seems to me, it breaks one of the OpenMP standard requirements: the program can be compiled without openmp support. I assume, that with this includes the program won't be able to be compiled without OpenMP support anymore because it may use some device-specific math functions explicitly.
> > > >  Instead, I would like to see some additional, device-scpecific math header file, that must be included explicitly to support some device-specific math functions. And we need to provide default implementations for those extra math functions for all the platforms we're going to support, including default host implementations.
> > >
> > >
> > > Can you provide an example of a conforming program that can't be compiled without OpenMP support? Regardless of the use of any device-specific functions (which isn't covered by the standard, of course, but might be needed in practice), the code still needs to be compilable by the host in order to generate the host-fallback version. This doesn't change that. Thus, any program that uses anything from this math.h, etc. needs to compile for the host, and thus, likely compiles without OpenMP support. Maybe I'm missing your point, however.
> >
> >
> > Assume we have something like this:
> >
> >   #pragma omp target if(cond)
> >   a = __nv_xxxx(....);
> >
> >
> > Instead of `__nv_xxx` you can try to use any Cuda-specific function, which is not the part of the standard `math.h`/`cmath` files. Will it be compilable even with OpenMP?
>
>
> I don't think that this changes that one way or the other. Your example won't work, AFAIK, unless you do something like:
>
>   #pragma omp target if(cond)
>   #ifdef __NVPTX__
>   a = __nv_xxxx(....);
>   #else
>   a = something_on_the_host;
>   #endif
>   
>
> and anything from these headers that doesn't also have a host version will suffer the same fate: if it won't also compile for the host (one way or another), then it won't work.


The problem with this header file is that it allows to use those Cuda-specific functions unconditionally in some cases:

  #pragma omp target
  a = __nv_xxxx(....);

It won't require any target-specific guards to compile this code (if we compile it only for Cuda-specific devices) and we're loosing the consistency here: in some cases target regions will require special device guards, in others, with the same function calls, it is not. And the worst thing, is that we implicitly allow to introduce this kind of incostistency into users code. That's why I would prefer to see a special kind of the include file, NVPTX specific, that must be included explicitly, so the user explictly commanded to use some target-specific math functions, if he really wants it. Plus, maybe, in this files we need force check of the platform and warn users that the functions from this header file must be used using device-specific checks. Or provide some kind of the default implementations for all the platforms, that do not support those math function natively.


Repository:
  rC Clang

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

https://reviews.llvm.org/D61399





More information about the cfe-commits mailing list