[OpenMP][Clang] Support for target math functions

[gtbercea]

Summary

In this patch we propose a temporary solution to resolving math functions for the NVPTX toolchain, temporary until OpenMP variant is supported by Clang.

We intercept the inclusion of math.h and cmath headers and if we are in the OpenMP-NVPTX case, we re-use CUDA's math function resolution mechanism.

Authors:
@gtbercea
@jdoerfert

[gtbercea]

• Minor fixes.

[jdoerfert]

For the record, this is an implementation of the scheme proposed in https://reviews.llvm.org/D60907#1484756.
There are drawbacks, see the TODO, but it will give most people a short term solution until we get OpenMP 5.0 variants.

Finally, there is a remote chance this will cause trouble to people that use math.h/cmath functions, e.g. with the old SSE hack, which are not available anymore.
I don't suspect that to happen but if it does we can, again as a short term solution, selectively extract declarations from the host cmath into the device cmath.

[jdoerfert]

lib/Headers/openmp_wrappers/__clang_openmp_math.h

30

I think this is a leftover we forgot to remove.

[gtbercea]

• Clean-up. Add header.

[hfinkel]

lib/Driver/ToolChains/Clang.cpp

1164

Please add a driver test covering this.

lib/Headers/openmp_wrappers/__clang_openmp_math.h

21

And this is why I wanted to just make __device__ work in OpenMP mode. But, as a temporary solution until we get declare variant working, I'm okay with this.

Modulo the naming nit and the missing driver test, LGTM.

47

I'd prefer that, as this is clang specific, we make this clear by naming this __CLANG_NO_HOST_MATH__.

[ABataev]

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.

[ABataev]

Moreover, I think this will cause troubles even in simple cases. Assume we have target if(cond) construct. In this case we will need to compile the target region for both, the device and the host. If the target region uses some device-specific math functions, it will break the compilation for the host.

[hfinkel]

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

[ABataev]

In D61399#1488299, @hfinkel wrote:

In 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?

[hfinkel]

In D61399#1488309, @ABataev wrote:

In D61399#1488299, @hfinkel wrote:

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

[ABataev]

In D61399#1488329, @hfinkel wrote:

In D61399#1488309, @ABataev wrote:

In D61399#1488299, @hfinkel wrote:

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

[hfinkel]

In D61399#1488366, @ABataev wrote:

In D61399#1488329, @hfinkel wrote:

In D61399#1488309, @ABataev wrote:

In D61399#1488299, @hfinkel wrote:

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

I believe that I understand your point, but two things:

1. I think that you're mistaken on the underlying premise. That code will not meaningfully compile without ifdefs, even if only CUDA-specific devices are the only ones selected. We *always* compile code for the host as well, not for offloading proper, but for the fallback (for execution when the offloading fails). If I emulate this situation by writing this:

#ifdef __NVPTX__
int __nv_floor();
#endif

int main() {
#pragma omp target
__nv_floor();
}

and try to compile using Clang with -fopenmp -fopenmp-targets=nvptx64, the compilation fails:

int1.cpp:8:1: error: use of undeclared identifier '__nv_floor'

and this is because, when we invoke the compilation for the host, there is no declaration for that function. This is true even though nvptx64 is the only target for which the code is being compiled (because we always also compile the host fallback).

2. I believe that the future state -- what we get by following this patch, and then when declare variant is available using that -- gives us all what we want. When we have declare variant, then all of the definitions in these headers will be declared as variants only available on the nvptx device, and so, for a user to use such a function they would need to explicit add a variant that is only available on the host. This would be explicit.

I think that you're getting at what I mentioned earlier, where if you have code that uses some function called, for example, rnorm3d -- this is a non-standard function, and so a user is free to define it, and then they compile with OpenMP target and include math.h, then the version of rnorm3d we bring in from the CUDA header will silently override their version (which would otherwise be implicitly declare target). But I don't think that this will happen either with this patch, instead, they'll get a warning about conflicting definitions (error: redefinition of 'rnorm3d'), and if it's an external-linkage function, then something should give them an error (although we should check this). The more interesting case, I think, actually comes when we switch to using declare variant, because then I think this silent override does occur, so we'd want to add a warning for when a variant delcared in a system header file would silently override a plain (non-variant) function declare/defined elsewhere. I believe that will give us all of the benefits of this while also addressing the concern you highlight.

[gtbercea]

• Address comments.
• Add math and cmath inclusion tests.
• Add driver test.

[ABataev]

In D61399#1488897, @hfinkel wrote:

In D61399#1488366, @ABataev wrote:

In D61399#1488329, @hfinkel wrote:

In D61399#1488309, @ABataev wrote:

In D61399#1488299, @hfinkel wrote:

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

I believe that I understand your point, but two things:

1. I think that you're mistaken on the underlying premise. That code will not meaningfully compile without ifdefs, even if only CUDA-specific devices are the only ones selected. We *always* compile code for the host as well, not for offloading proper, but for the fallback (for execution when the offloading fails). If I emulate this situation by writing this:

#ifdef __NVPTX__
int __nv_floor();
#endif

int main() {
#pragma omp target
__nv_floor();
}

and try to compile using Clang with -fopenmp -fopenmp-targets=nvptx64, the compilation fails:

int1.cpp:8:1: error: use of undeclared identifier '__nv_floor'

and this is because, when we invoke the compilation for the host, there is no declaration for that function. This is true even though nvptx64 is the only target for which the code is being compiled (because we always also compile the host fallback).

2. I believe that the future state -- what we get by following this patch, and then when declare variant is available using that -- gives us all what we want. When we have declare variant, then all of the definitions in these headers will be declared as variants only available on the nvptx device, and so, for a user to use such a function they would need to explicit add a variant that is only available on the host. This would be explicit.

I think that you're getting at what I mentioned earlier, where if you have code that uses some function called, for example, rnorm3d -- this is a non-standard function, and so a user is free to define it, and then they compile with OpenMP target and include math.h, then the version of rnorm3d we bring in from the CUDA header will silently override their version (which would otherwise be implicitly declare target). But I don't think that this will happen either with this patch, instead, they'll get a warning about conflicting definitions (error: redefinition of 'rnorm3d'), and if it's an external-linkage function, then something should give them an error (although we should check this). The more interesting case, I think, actually comes when we switch to using declare variant, because then I think this silent override does occur, so we'd want to add a warning for when a variant delcared in a system header file would silently override a plain (non-variant) function declare/defined elsewhere. I believe that will give us all of the benefits of this while also addressing the concern you highlight.

Ahh, yes, I forgot that we still generate the host version of the target region. Still, I think we need to prvide the default implementation of those non-standard functions (they can be very simple, maybe reporting error is going to be enough), which can be overriden by user.
Also, if I do recall correctly, this solution works only for C++ (because we use our own declaration for the standard math functions, they are automatically marked as non-builtins). For regular C in many cases, instead of the library function calls, the llvm intrinsic is generated. Did you check that it works for С?

[gtbercea]

@ABataev this patch works for both C and C++ and for both math.h and cmath headers.

[ABataev]

In D61399#1489757, @gtbercea wrote:

@ABataev this patch works for both C and C++ and for both math.h and cmath headers.

Did you test it for the builtins? like pow, powf and powl? How are the builtins resolved with this patch?

[gtbercea]

In D61399#1489762, @ABataev wrote:

In D61399#1489757, @gtbercea wrote:

@ABataev this patch works for both C and C++ and for both math.h and cmath headers.

Did you test it for the builtins? like pow, powf and powl? How are the builtins resolved with this patch?

I have. There are tests for this.

[hfinkel]

Still, I think we need to prvide the default implementation of those non-standard functions (they can be very simple, maybe reporting error is going to be enough), which can be overriden by user.

I appreciate your motivation, and I agree with you to some extent. I don't object to having generic versions of useful math functions, but I don't think they should be required. It's not reasonable to make someone add generic versions of every function which happens to appear in a system/target-specific math.h header. NVPTX won't be the only target that has target-optimized functions that get pulled in, even from our own headers, but system headers also have differences anyway depending on what preprocessor macros are defined. In the end, people can write portable code if they stick to what's in the standard, and we should make it reasonably easy for them to step outside of the standard to do what they need to do when the standard subset of available functionality doesn't meet their needs for whatever reason. This is what we do for C/C++, where we provide intrinsics and other system functions for those who can't write their code only using the facilities that C/C++ provide.

In any case, I think that we can figure out how to add generic versions of non-standard math functions in a separate thread. I think that we should move forward with this and then make progress on generic versions separately. It's also possible that we want to fold this discussion into the discussion on an LLVM math library (we've talked about this for some time in the context of vector math libraries, and I'd not thought about accelerators in this context, but maybe this is all related).

[ABataev]

In D61399#1490189, @hfinkel wrote:

Still, I think we need to prvide the default implementation of those non-standard functions (they can be very simple, maybe reporting error is going to be enough), which can be overriden by user.

I appreciate your motivation, and I agree with you to some extent. I don't object to having generic versions of useful math functions, but I don't think they should be required. It's not reasonable to make someone add generic versions of every function which happens to appear in a system/target-specific math.h header. NVPTX won't be the only target that has target-optimized functions that get pulled in, even from our own headers, but system headers also have differences anyway depending on what preprocessor macros are defined. In the end, people can write portable code if they stick to what's in the standard, and we should make it reasonably easy for them to step outside of the standard to do what they need to do when the standard subset of available functionality doesn't meet their needs for whatever reason. This is what we do for C/C++, where we provide intrinsics and other system functions for those who can't write their code only using the facilities that C/C++ provide.

In any case, I think that we can figure out how to add generic versions of non-standard math functions in a separate thread. I think that we should move forward with this and then make progress on generic versions separately. It's also possible that we want to fold this discussion into the discussion on an LLVM math library (we've talked about this for some time in the context of vector math libraries, and I'd not thought about accelerators in this context, but maybe this is all related).

It is up to you. I don't have strong objections if you think this will work as required. Just the tests must be fixed, especially codegen tests.

test/CodeGen/nvptx_device_cmath_functions.c

1 ↗

(On Diff #197911)

1. Provide tests for C++ too.
2. Do not use driver to generate the code, use frontend.
3. Do not include real system header files. Use some stubs instead.

[hfinkel]

It is up to you. I don't have strong objections if you think this will work as required. Just the tests must be fixed, especially codegen tests.

Thanks, Alexey. I think this will work as required, and then we'll be able to update it when we get declare variant. Agreed on the tests (on all points).

[gtbercea]

• Add driver test.

[gtbercea]

• Add new tests. Add stub headers.
• Remove old tests.

[jdoerfert]

Alexey, is this is good to go now?

[ABataev]

lib/Headers/__clang_cuda_cmath.h

54

Why we have this guard here? It does not work for OpenMP? Why?

lib/Headers/__clang_cuda_device_functions.h

49–53

Can we do anything with this in С mode? I mean, to allow it in C.

[gtbercea]

lib/Headers/__clang_cuda_cmath.h

54

Because all the FP_XXX macros are defined in cmath and for OpenMP we can't include it yet because we don't support variant yet.

lib/Headers/__clang_cuda_device_functions.h

49–53

We can if we rename the function.

[ABataev]

lib/Headers/__clang_cuda_cmath.h

54

The better to add TODO or FIXME to fix this once variant construct in supported.

lib/Headers/__clang_cuda_device_functions.h

49–53

Take a look here, probably it will solve the problem:
https://clang.llvm.org/docs/AttributeReference.html#overloadable

[gtbercea]

• Address comments.

[ABataev]

lib/Headers/__clang_cuda_cmath.h

444

I see that the same guard is used lib/Headers/__clang_cuda_device_functions.h, but for different set of functions. Is this ok?

[gtbercea]

lib/Headers/__clang_cuda_cmath.h

444

Yep, it's intentional. Again the variant issue.

[ABataev]

LG with a nit.

lib/Headers/__clang_cuda_cmath.h

444

Then again, add TODO or FIXME

[gtbercea]

• Address comments.

[JDevlieghere]

I've reverted this in rL360192 because it breaks stage 2 builds on GreenDragon. Please see the commit message and the inline comment for more details.

lib/Headers/CMakeLists.txt

36

This doesn't do what you think it would do. The files are copied into the root of the resource directory, which causes stage 2 build failures on GreenDragon.

[hfinkel]

lib/Headers/CMakeLists.txt

36

Can you provide a link to the failure log? Is the problem that the files are not copied into their subdirectory?

[JDevlieghere]

lib/Headers/CMakeLists.txt

36

Correct. There's not much to see, but here's a build that fails because of this: http://green.lab.llvm.org/green/view/LLDB/job/lldb-cmake/25557/console

CMake Error at cmake/modules/HandleLLVMOptions.cmake:497 (message):
  LLVM_ENABLE_MODULES is not supported by this compiler

The cmake log output, which I grabbed from the bot is more descriptive:

Performing C++ SOURCE FILE Test CXX_SUPPORTS_MODULES failed with the following output:
Change Dir: /Users/buildslave/jenkins/workspace/lldb-cmake/lldb-build/CMakeFiles/CMakeTmp

Run Build Command:"/usr/local/bin/ninja" "cmTC_0dd11"
[1/2] Building CXX object CMakeFiles/cmTC_0dd11.dir/src.cxx.o
FAILED: CMakeFiles/cmTC_0dd11.dir/src.cxx.o
/Users/buildslave/jenkins/workspace/lldb-cmake/host-compiler/bin/clang++    -fPIC -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -std=c++11 -DCXX_SUPPORTS_MODULES  -Werror=unguarded-availability-new -fmodules -fmodules-cache-path=/Users/buildslave/jenkins/workspace/lldb-cmake/lldb-build/module.cache -fcxx-modules -isysroot /Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk -mmacosx-version-min=10.9 -o CMakeFiles/cmTC_0dd11.dir/src.cxx.o -c src.cxx
While building module 'Darwin' imported from /Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk/usr/include/assert.h:42:
While building module 'std' imported from /Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk/usr/include/tgmath.h:27:
In file included from <module-includes>:2:
/Users/buildslave/jenkins/workspace/lldb-cmake/host-compiler/bin/../include/c++/v1/ctype.h:38:15: fatal error: cyclic dependency in module 'Darwin': Darwin -> std -> Darwin
#include_next <ctype.h>
              ^
While building module 'Darwin' imported from /Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk/usr/include/assert.h:42:
In file included from <module-includes>:89:
In file included from /Users/buildslave/jenkins/workspace/lldb-cmake/host-compiler/lib/clang/9.0.0/include/tgmath.h:21:
/Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk/usr/include/tgmath.h:27:10: fatal error: could not build module 'std'
#include <math.h>
 ~~~~~~~~^
In file included from src.cxx:2:
In file included from /Users/buildslave/jenkins/workspace/lldb-cmake/host-compiler/bin/../include/c++/v1/cassert:20:
/Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.14.sdk/usr/include/assert.h:42:10: fatal error: could not build module 'Darwin'
#include <sys/cdefs.h>
 ~~~~~~~~^
3 errors generated.
ninja: build stopped: subcommand failed.

Source file was:
#undef NDEBUG
                               #include <cassert>
                               #define NDEBUG
                               #include <cassert>
                               int main() { assert(this code is not compiled); }

[gtbercea]

• Fix move to openmp_wrapper folder. Fix header ordering problem.

[gtbercea]

• Eliminate declarations of functions not needed for math function resolution.

lib/Driver/ToolChain.cpp

bool ToolChain::needsProfileRT(const ArgList &Args) {
   if (needsGCovInstrumentation(Args) ||
       Args.hasArg(options::OPT_fprofile_generate) ||
       Args.hasArg(options::OPT_fprofile_generate_EQ) ||
       Args.hasArg(options::OPT_fcs_profile_generate) ||
       Args.hasArg(options::OPT_fcs_profile_generate_EQ) ||
       Args.hasArg(options::OPT_fprofile_instr_generate) ||
       Args.hasArg(options::OPT_fprofile_instr_generate_EQ) ||
       Args.hasArg(options::OPT_fcreate_profile) ||
-      Args.hasArg(options::OPT_forder_file_instrumentation)) 
+      Args.hasArg(options::OPT_forder_file_instrumentation))
     return true;
 
   return false;
 }
 
 bool ToolChain::needsGCovInstrumentation(const llvm::opt::ArgList &Args) {
   return Args.hasFlag(options::OPT_fprofile_arcs, options::OPT_fno_profile_arcs,
                       false) ||

lib/Driver/ToolChains/Clang.cpp

void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
 
   // Add offload include arguments specific for CUDA.  This must happen before
   // we -I or -include anything else, because we must pick up the CUDA headers
   // from the particular CUDA installation, rather than from e.g.
   // /usr/local/include.
   if (JA.isOffloading(Action::OFK_Cuda))
     getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
 
+  // If we are offloading to a target via OpenMP we need to include the
+  // openmp_wrappers folder which contains alternative system headers.
+  if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
+      getToolChain().getTriple().isNVPTX()){
+    if (!Args.hasArg(options::OPT_nobuiltininc)) {
+      // Add openmp_wrappers/* to our system include path.  This lets us wrap
+      // standard library headers.
+      SmallString<128> P(D.ResourceDir);
+      llvm::sys::path::append(P, "include");
+      llvm::sys::path::append(P, "openmp_wrappers");
+      CmdArgs.push_back("-internal-isystem");
+      CmdArgs.push_back(Args.MakeArgString(P));
+    }
+
+    CmdArgs.push_back("-include");
+    CmdArgs.push_back("__clang_openmp_math.h");
+  }
+
   // Add -i* options, and automatically translate to
   // -include-pch/-include-pth for transparent PCH support. It's
   // wonky, but we include looking for .gch so we can support seamless
   // replacement into a build system already set up to be generating
   // .gch files.
 
   if (getToolChain().getDriver().IsCLMode()) {
     const Arg *YcArg = Args.getLastArg(options::OPT__SLASH_Yc);

lib/Headers/CMakeLists.txt

set(files
   avx512vldqintrin.h
   avx512vlintrin.h
   avx512vpopcntdqvlintrin.h
   avx512vnniintrin.h
   avx512vlvnniintrin.h
   avxintrin.h
   bmi2intrin.h
   bmiintrin.h
   __clang_cuda_builtin_vars.h
   __clang_cuda_cmath.h
   __clang_cuda_complex_builtins.h
   __clang_cuda_device_functions.h
   __clang_cuda_intrinsics.h
   __clang_cuda_libdevice_declares.h
   __clang_cuda_math_forward_declares.h
   __clang_cuda_runtime_wrapper.h
   cetintrin.h
set(cuda_wrapper_files
   cuda_wrappers/complex
   cuda_wrappers/new
 )
 
 set(ppc_wrapper_files
   ppc_wrappers/mmintrin.h
 )
 
+set(openmp_wrapper_files
+  openmp_wrappers/math.h
+  openmp_wrappers/cmath
+  openmp_wrappers/__clang_openmp_math.h
+)
+
 set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
 set(out_files)
 set(generated_files)
 
 function(copy_header_to_output_dir src_dir file)
   set(src ${src_dir}/${file})
   set(dst ${output_dir}/${file})
   add_custom_command(OUTPUT ${dst}
function(clang_generate_header td_option td_file out_file)
   copy_header_to_output_dir(${CMAKE_CURRENT_BINARY_DIR} ${out_file})
   set(out_files ${out_files} PARENT_SCOPE)
   list(APPEND generated_files "${CMAKE_CURRENT_BINARY_DIR}/${out_file}")
   set(generated_files ${generated_files} PARENT_SCOPE)
 endfunction(clang_generate_header)
 
 
 # Copy header files from the source directory to the build directory
-foreach( f ${files} ${cuda_wrapper_files} ${ppc_wrapper_files} )
+foreach( f ${files} ${cuda_wrapper_files} ${ppc_wrapper_files} ${openmp_wrapper_files})
   copy_header_to_output_dir(${CMAKE_CURRENT_SOURCE_DIR} ${f})
 endforeach( f )
 
 # Generate header files and copy them to the build directory
 # Generate arm_neon.h
 clang_generate_header(-gen-arm-neon arm_neon.td arm_neon.h)
 # Generate arm_fp16.h
 clang_generate_header(-gen-arm-fp16 arm_fp16.td arm_fp16.h)
install(
   DESTINATION ${header_install_dir}/cuda_wrappers
   COMPONENT clang-resource-headers)
 
 install(
   FILES ${ppc_wrapper_files}
   DESTINATION ${header_install_dir}/ppc_wrappers
   COMPONENT clang-resource-headers)
 
+install(
+  FILES ${openmp_wrapper_files}
+  DESTINATION ${header_install_dir}/openmp_wrappers
+  COMPONENT clang-resource-headers)
+
 if (NOT LLVM_ENABLE_IDE)
   add_llvm_install_targets(install-clang-resource-headers
                            DEPENDS clang-resource-headers
                            COMPONENT clang-resource-headers)
 endif()

lib/Headers/__clang_cuda_cmath.h

 // It's important that we declare the functions in the global namespace and pull
 // them into namespace std with using statements, as opposed to simply declaring
 // these functions in namespace std, because our device functions need to
 // overload the standard library functions, which may be declared in the global
 // namespace or in std, depending on the degree of conformance of the stdlib
 // implementation.  Declaring in the global namespace and pulling into namespace
 // std covers all of the known knowns.
 
+#ifdef _OPENMP
+#define __DEVICE__ static __attribute__((always_inline))
+#else
 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
+#endif
 
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
 __DEVICE__ float exp(float __x) { return ::expf(__x); }
 __DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
+// TODO: remove when variant is supported
+#ifndef _OPENMP
 __DEVICE__ int fpclassify(float __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
                               FP_ZERO, __x);
 }
 __DEVICE__ int fpclassify(double __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
                               FP_ZERO, __x);
 }
+#endif
 __DEVICE__ float frexp(float __arg, int *__exp) {
   return ::frexpf(__arg, __exp);
 }
 
 // For inscrutable reasons, the CUDA headers define these functions for us on
 // Windows.
 #ifndef _MSC_VER
 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }

 using ::modff;
 using ::nearbyintf;
 using ::nextafterf;
 using ::powf;
 using ::remainderf;
 using ::remquof;
 using ::rintf;
 using ::roundf;
+// TODO: remove once variant is supported
+#ifndef _OPENMP
 using ::scalblnf;
+#endif
 using ::scalbnf;
 using ::sinf;
 using ::sinhf;
 using ::sqrtf;
 using ::tanf;
 using ::tanhf;
 using ::tgammaf;
 using ::truncf;

lib/Headers/__clang_cuda_device_functions.h

 /*===---- __clang_cuda_device_functions.h - CUDA runtime support -----------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_CUDA_DEVICE_FUNCTIONS_H__
 #define __CLANG_CUDA_DEVICE_FUNCTIONS_H__
 
+#ifndef _OPENMP
 #if CUDA_VERSION < 9000
 #error This file is intended to be used with CUDA-9+ only.
 #endif
+#endif
 
 // __DEVICE__ is a helper macro with common set of attributes for the wrappers
 // we implement in this file. We need static in order to avoid emitting unused
 // functions and __forceinline__ helps inlining these wrappers at -O1.
 #pragma push_macro("__DEVICE__")
+#ifdef _OPENMP
+#define __DEVICE__ static __attribute__((always_inline))
+#else
 #define __DEVICE__ static __device__ __forceinline__
+#endif
 
 // libdevice provides fast low precision and slow full-recision implementations
 // for some functions. Which one gets selected depends on
 // __CLANG_CUDA_APPROX_TRANSCENDENTALS__ which gets defined by clang if
 // -ffast-math or -fcuda-approx-transcendentals are in effect.
 #pragma push_macro("__FAST_OR_SLOW")
 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
 #define __FAST_OR_SLOW(fast, slow) fast
 #else
 #define __FAST_OR_SLOW(fast, slow) slow
 #endif
 
 __DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); }
 __DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); }
 __DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); }
 __DEVICE__ unsigned int __brev(unsigned int __a) { return __nv_brev(__a); }
 __DEVICE__ unsigned long long __brevll(unsigned long long __a) {
   return __nv_brevll(__a);
 }
+#if defined(__cplusplus)
 __DEVICE__ void __brkpt() { asm volatile("brkpt;"); }
 __DEVICE__ void __brkpt(int __a) { __brkpt(); }
+#else
+__DEVICE__ void __attribute__((overloadable)) __brkpt(void) { asm volatile("brkpt;"); }
+__DEVICE__ void __attribute__((overloadable)) __brkpt(int __a) { __brkpt(); }
+#endif
 __DEVICE__ unsigned int __byte_perm(unsigned int __a, unsigned int __b,
                                     unsigned int __c) {
   return __nv_byte_perm(__a, __b, __c);
 }
 __DEVICE__ int __clz(int __a) { return __nv_clz(__a); }
 __DEVICE__ int __clzll(long long __a) { return __nv_clzll(__a); }
 __DEVICE__ float __cosf(float __a) { return __nv_fast_cosf(__a); }
 __DEVICE__ double __dAtomicAdd(double *__p, double __v) {

 __DEVICE__ int ilogbf(float __a) { return __nv_ilogbf(__a); }
 __DEVICE__ double j0(double __a) { return __nv_j0(__a); }
 __DEVICE__ float j0f(float __a) { return __nv_j0f(__a); }
 __DEVICE__ double j1(double __a) { return __nv_j1(__a); }
 __DEVICE__ float j1f(float __a) { return __nv_j1f(__a); }
 __DEVICE__ double jn(int __n, double __a) { return __nv_jn(__n, __a); }
 __DEVICE__ float jnf(int __n, float __a) { return __nv_jnf(__n, __a); }
 #if defined(__LP64__) || defined(_WIN64)
-__DEVICE__ long labs(long __a) { return llabs(__a); };
+__DEVICE__ long labs(long __a) { return __nv_llabs(__a); };
 #else
 __DEVICE__ long labs(long __a) { return __nv_abs(__a); };
 #endif
 __DEVICE__ double ldexp(double __a, int __b) { return __nv_ldexp(__a, __b); }
 __DEVICE__ float ldexpf(float __a, int __b) { return __nv_ldexpf(__a, __b); }
 __DEVICE__ double lgamma(double __a) { return __nv_lgamma(__a); }
 __DEVICE__ float lgammaf(float __a) { return __nv_lgammaf(__a); }
 __DEVICE__ long long llabs(long long __a) { return __nv_llabs(__a); }

 __DEVICE__ long lroundf(float __a) { return llroundf(__a); }
 #else
 __DEVICE__ long lrint(double __a) { return (long)rint(__a); }
 __DEVICE__ long lrintf(float __a) { return __float2int_rn(__a); }
 __DEVICE__ long lround(double __a) { return round(__a); }
 __DEVICE__ long lroundf(float __a) { return roundf(__a); }
 #endif
 __DEVICE__ int max(int __a, int __b) { return __nv_max(__a, __b); }
+// These functions shouldn't be declared when including this header
+// for math function resolution purposes.
+#ifndef _OPENMP
 __DEVICE__ void *memcpy(void *__a, const void *__b, size_t __c) {
   return __builtin_memcpy(__a, __b, __c);
 }
 __DEVICE__ void *memset(void *__a, int __b, size_t __c) {
   return __builtin_memset(__a, __b, __c);
 }
+#endif
 __DEVICE__ int min(int __a, int __b) { return __nv_min(__a, __b); }
 __DEVICE__ double modf(double __a, double *__b) { return __nv_modf(__a, __b); }
 __DEVICE__ float modff(float __a, float *__b) { return __nv_modff(__a, __b); }
 __DEVICE__ double nearbyint(double __a) { return __nv_nearbyint(__a); }
 __DEVICE__ float nearbyintf(float __a) { return __nv_nearbyintf(__a); }
 __DEVICE__ double nextafter(double __a, double __b) {
   return __nv_nextafter(__a, __b);
 }
__DEVICE__ float rnormf(int __dim, const float *__t) {
   return __nv_rnormf(__dim, __t);
 }
 __DEVICE__ double round(double __a) { return __nv_round(__a); }
 __DEVICE__ float roundf(float __a) { return __nv_roundf(__a); }
 __DEVICE__ double rsqrt(double __a) { return __nv_rsqrt(__a); }
 __DEVICE__ float rsqrtf(float __a) { return __nv_rsqrtf(__a); }
 __DEVICE__ double scalbn(double __a, int __b) { return __nv_scalbn(__a, __b); }
 __DEVICE__ float scalbnf(float __a, int __b) { return __nv_scalbnf(__a, __b); }
+// TODO: remove once variant is supported
+#ifndef _OPENMP
 __DEVICE__ double scalbln(double __a, long __b) {
   if (__b > INT_MAX)
     return __a > 0 ? HUGE_VAL : -HUGE_VAL;
   if (__b < INT_MIN)
     return __a > 0 ? 0.0 : -0.0;
   return scalbn(__a, (int)__b);
 }
 __DEVICE__ float scalblnf(float __a, long __b) {
   if (__b > INT_MAX)
     return __a > 0 ? HUGE_VALF : -HUGE_VALF;
   if (__b < INT_MIN)
     return __a > 0 ? 0.f : -0.f;
   return scalbnf(__a, (int)__b);
 }
+#endif
 __DEVICE__ double sin(double __a) { return __nv_sin(__a); }
 __DEVICE__ void sincos(double __a, double *__s, double *__c) {
   return __nv_sincos(__a, __s, __c);
 }
 __DEVICE__ void sincosf(float __a, float *__s, float *__c) {
   return __FAST_OR_SLOW(__nv_fast_sincosf, __nv_sincosf)(__a, __s, __c);
 }
 __DEVICE__ void sincospi(double __a, double *__s, double *__c) {

lib/Headers/__clang_cuda_libdevice_declares.h

 /*===-- __clang_cuda_libdevice_declares.h - decls for libdevice functions --===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_CUDA_LIBDEVICE_DECLARES_H__
 #define __CLANG_CUDA_LIBDEVICE_DECLARES_H__
 
+#if defined(__cplusplus)
 extern "C" {
+#endif
 
-__device__ int __nv_abs(int __a);
-__device__ double __nv_acos(double __a);
-__device__ float __nv_acosf(float __a);
-__device__ double __nv_acosh(double __a);
-__device__ float __nv_acoshf(float __a);
-__device__ double __nv_asin(double __a);
-__device__ float __nv_asinf(float __a);
-__device__ double __nv_asinh(double __a);
-__device__ float __nv_asinhf(float __a);
-__device__ double __nv_atan2(double __a, double __b);
-__device__ float __nv_atan2f(float __a, float __b);
-__device__ double __nv_atan(double __a);
-__device__ float __nv_atanf(float __a);
-__device__ double __nv_atanh(double __a);
-__device__ float __nv_atanhf(float __a);
-__device__ int __nv_brev(int __a);
-__device__ long long __nv_brevll(long long __a);
-__device__ int __nv_byte_perm(int __a, int __b, int __c);
-__device__ double __nv_cbrt(double __a);
-__device__ float __nv_cbrtf(float __a);
-__device__ double __nv_ceil(double __a);
-__device__ float __nv_ceilf(float __a);
-__device__ int __nv_clz(int __a);
-__device__ int __nv_clzll(long long __a);
-__device__ double __nv_copysign(double __a, double __b);
-__device__ float __nv_copysignf(float __a, float __b);
-__device__ double __nv_cos(double __a);
-__device__ float __nv_cosf(float __a);
-__device__ double __nv_cosh(double __a);
-__device__ float __nv_coshf(float __a);
-__device__ double __nv_cospi(double __a);
-__device__ float __nv_cospif(float __a);
-__device__ double __nv_cyl_bessel_i0(double __a);
-__device__ float __nv_cyl_bessel_i0f(float __a);
-__device__ double __nv_cyl_bessel_i1(double __a);
-__device__ float __nv_cyl_bessel_i1f(float __a);
-__device__ double __nv_dadd_rd(double __a, double __b);
-__device__ double __nv_dadd_rn(double __a, double __b);
-__device__ double __nv_dadd_ru(double __a, double __b);
-__device__ double __nv_dadd_rz(double __a, double __b);
-__device__ double __nv_ddiv_rd(double __a, double __b);
-__device__ double __nv_ddiv_rn(double __a, double __b);
-__device__ double __nv_ddiv_ru(double __a, double __b);
-__device__ double __nv_ddiv_rz(double __a, double __b);
-__device__ double __nv_dmul_rd(double __a, double __b);
-__device__ double __nv_dmul_rn(double __a, double __b);
-__device__ double __nv_dmul_ru(double __a, double __b);
-__device__ double __nv_dmul_rz(double __a, double __b);
-__device__ float __nv_double2float_rd(double __a);
-__device__ float __nv_double2float_rn(double __a);
-__device__ float __nv_double2float_ru(double __a);
-__device__ float __nv_double2float_rz(double __a);
-__device__ int __nv_double2hiint(double __a);
-__device__ int __nv_double2int_rd(double __a);
-__device__ int __nv_double2int_rn(double __a);
-__device__ int __nv_double2int_ru(double __a);
-__device__ int __nv_double2int_rz(double __a);
-__device__ long long __nv_double2ll_rd(double __a);
-__device__ long long __nv_double2ll_rn(double __a);
-__device__ long long __nv_double2ll_ru(double __a);
-__device__ long long __nv_double2ll_rz(double __a);
-__device__ int __nv_double2loint(double __a);
-__device__ unsigned int __nv_double2uint_rd(double __a);
-__device__ unsigned int __nv_double2uint_rn(double __a);
-__device__ unsigned int __nv_double2uint_ru(double __a);
-__device__ unsigned int __nv_double2uint_rz(double __a);
-__device__ unsigned long long __nv_double2ull_rd(double __a);
-__device__ unsigned long long __nv_double2ull_rn(double __a);
-__device__ unsigned long long __nv_double2ull_ru(double __a);
-__device__ unsigned long long __nv_double2ull_rz(double __a);
-__device__ unsigned long long __nv_double_as_longlong(double __a);
-__device__ double __nv_drcp_rd(double __a);
-__device__ double __nv_drcp_rn(double __a);
-__device__ double __nv_drcp_ru(double __a);
-__device__ double __nv_drcp_rz(double __a);
-__device__ double __nv_dsqrt_rd(double __a);
-__device__ double __nv_dsqrt_rn(double __a);
-__device__ double __nv_dsqrt_ru(double __a);
-__device__ double __nv_dsqrt_rz(double __a);
-__device__ double __nv_dsub_rd(double __a, double __b);
-__device__ double __nv_dsub_rn(double __a, double __b);
-__device__ double __nv_dsub_ru(double __a, double __b);
-__device__ double __nv_dsub_rz(double __a, double __b);
-__device__ double __nv_erfc(double __a);
-__device__ float __nv_erfcf(float __a);
-__device__ double __nv_erfcinv(double __a);
-__device__ float __nv_erfcinvf(float __a);
-__device__ double __nv_erfcx(double __a);
-__device__ float __nv_erfcxf(float __a);
-__device__ double __nv_erf(double __a);
-__device__ float __nv_erff(float __a);
-__device__ double __nv_erfinv(double __a);
-__device__ float __nv_erfinvf(float __a);
-__device__ double __nv_exp10(double __a);
-__device__ float __nv_exp10f(float __a);
-__device__ double __nv_exp2(double __a);
-__device__ float __nv_exp2f(float __a);
-__device__ double __nv_exp(double __a);
-__device__ float __nv_expf(float __a);
-__device__ double __nv_expm1(double __a);
-__device__ float __nv_expm1f(float __a);
-__device__ double __nv_fabs(double __a);
-__device__ float __nv_fabsf(float __a);
-__device__ float __nv_fadd_rd(float __a, float __b);
-__device__ float __nv_fadd_rn(float __a, float __b);
-__device__ float __nv_fadd_ru(float __a, float __b);
-__device__ float __nv_fadd_rz(float __a, float __b);
-__device__ float __nv_fast_cosf(float __a);
-__device__ float __nv_fast_exp10f(float __a);
-__device__ float __nv_fast_expf(float __a);
-__device__ float __nv_fast_fdividef(float __a, float __b);
-__device__ float __nv_fast_log10f(float __a);
-__device__ float __nv_fast_log2f(float __a);
-__device__ float __nv_fast_logf(float __a);
-__device__ float __nv_fast_powf(float __a, float __b);
-__device__ void __nv_fast_sincosf(float __a, float *__s, float *__c);
-__device__ float __nv_fast_sinf(float __a);
-__device__ float __nv_fast_tanf(float __a);
-__device__ double __nv_fdim(double __a, double __b);
-__device__ float __nv_fdimf(float __a, float __b);
-__device__ float __nv_fdiv_rd(float __a, float __b);
-__device__ float __nv_fdiv_rn(float __a, float __b);
-__device__ float __nv_fdiv_ru(float __a, float __b);
-__device__ float __nv_fdiv_rz(float __a, float __b);
-__device__ int __nv_ffs(int __a);
-__device__ int __nv_ffsll(long long __a);
-__device__ int __nv_finitef(float __a);
-__device__ unsigned short __nv_float2half_rn(float __a);
-__device__ int __nv_float2int_rd(float __a);
-__device__ int __nv_float2int_rn(float __a);
-__device__ int __nv_float2int_ru(float __a);
-__device__ int __nv_float2int_rz(float __a);
-__device__ long long __nv_float2ll_rd(float __a);
-__device__ long long __nv_float2ll_rn(float __a);
-__device__ long long __nv_float2ll_ru(float __a);
-__device__ long long __nv_float2ll_rz(float __a);
-__device__ unsigned int __nv_float2uint_rd(float __a);
-__device__ unsigned int __nv_float2uint_rn(float __a);
-__device__ unsigned int __nv_float2uint_ru(float __a);
-__device__ unsigned int __nv_float2uint_rz(float __a);
-__device__ unsigned long long __nv_float2ull_rd(float __a);
-__device__ unsigned long long __nv_float2ull_rn(float __a);
-__device__ unsigned long long __nv_float2ull_ru(float __a);
-__device__ unsigned long long __nv_float2ull_rz(float __a);
-__device__ int __nv_float_as_int(float __a);
-__device__ unsigned int __nv_float_as_uint(float __a);
-__device__ double __nv_floor(double __a);
-__device__ float __nv_floorf(float __a);
-__device__ double __nv_fma(double __a, double __b, double __c);
-__device__ float __nv_fmaf(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rd(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rn(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_ru(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rz(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rd(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rn(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ru(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rz(float __a, float __b, float __c);
-__device__ double __nv_fma_rd(double __a, double __b, double __c);
-__device__ double __nv_fma_rn(double __a, double __b, double __c);
-__device__ double __nv_fma_ru(double __a, double __b, double __c);
-__device__ double __nv_fma_rz(double __a, double __b, double __c);
-__device__ double __nv_fmax(double __a, double __b);
-__device__ float __nv_fmaxf(float __a, float __b);
-__device__ double __nv_fmin(double __a, double __b);
-__device__ float __nv_fminf(float __a, float __b);
-__device__ double __nv_fmod(double __a, double __b);
-__device__ float __nv_fmodf(float __a, float __b);
-__device__ float __nv_fmul_rd(float __a, float __b);
-__device__ float __nv_fmul_rn(float __a, float __b);
-__device__ float __nv_fmul_ru(float __a, float __b);
-__device__ float __nv_fmul_rz(float __a, float __b);
-__device__ float __nv_frcp_rd(float __a);
-__device__ float __nv_frcp_rn(float __a);
-__device__ float __nv_frcp_ru(float __a);
-__device__ float __nv_frcp_rz(float __a);
-__device__ double __nv_frexp(double __a, int *__b);
-__device__ float __nv_frexpf(float __a, int *__b);
-__device__ float __nv_frsqrt_rn(float __a);
-__device__ float __nv_fsqrt_rd(float __a);
-__device__ float __nv_fsqrt_rn(float __a);
-__device__ float __nv_fsqrt_ru(float __a);
-__device__ float __nv_fsqrt_rz(float __a);
-__device__ float __nv_fsub_rd(float __a, float __b);
-__device__ float __nv_fsub_rn(float __a, float __b);
-__device__ float __nv_fsub_ru(float __a, float __b);
-__device__ float __nv_fsub_rz(float __a, float __b);
-__device__ int __nv_hadd(int __a, int __b);
-__device__ float __nv_half2float(unsigned short __h);
-__device__ double __nv_hiloint2double(int __a, int __b);
-__device__ double __nv_hypot(double __a, double __b);
-__device__ float __nv_hypotf(float __a, float __b);
-__device__ int __nv_ilogb(double __a);
-__device__ int __nv_ilogbf(float __a);
-__device__ double __nv_int2double_rn(int __a);
-__device__ float __nv_int2float_rd(int __a);
-__device__ float __nv_int2float_rn(int __a);
-__device__ float __nv_int2float_ru(int __a);
-__device__ float __nv_int2float_rz(int __a);
-__device__ float __nv_int_as_float(int __a);
-__device__ int __nv_isfinited(double __a);
-__device__ int __nv_isinfd(double __a);
-__device__ int __nv_isinff(float __a);
-__device__ int __nv_isnand(double __a);
-__device__ int __nv_isnanf(float __a);
-__device__ double __nv_j0(double __a);
-__device__ float __nv_j0f(float __a);
-__device__ double __nv_j1(double __a);
-__device__ float __nv_j1f(float __a);
-__device__ float __nv_jnf(int __a, float __b);
-__device__ double __nv_jn(int __a, double __b);
-__device__ double __nv_ldexp(double __a, int __b);
-__device__ float __nv_ldexpf(float __a, int __b);
-__device__ double __nv_lgamma(double __a);
-__device__ float __nv_lgammaf(float __a);
-__device__ double __nv_ll2double_rd(long long __a);
-__device__ double __nv_ll2double_rn(long long __a);
-__device__ double __nv_ll2double_ru(long long __a);
-__device__ double __nv_ll2double_rz(long long __a);
-__device__ float __nv_ll2float_rd(long long __a);
-__device__ float __nv_ll2float_rn(long long __a);
-__device__ float __nv_ll2float_ru(long long __a);
-__device__ float __nv_ll2float_rz(long long __a);
-__device__ long long __nv_llabs(long long __a);
-__device__ long long __nv_llmax(long long __a, long long __b);
-__device__ long long __nv_llmin(long long __a, long long __b);
-__device__ long long __nv_llrint(double __a);
-__device__ long long __nv_llrintf(float __a);
-__device__ long long __nv_llround(double __a);
-__device__ long long __nv_llroundf(float __a);
-__device__ double __nv_log10(double __a);
-__device__ float __nv_log10f(float __a);
-__device__ double __nv_log1p(double __a);
-__device__ float __nv_log1pf(float __a);
-__device__ double __nv_log2(double __a);
-__device__ float __nv_log2f(float __a);
-__device__ double __nv_logb(double __a);
-__device__ float __nv_logbf(float __a);
-__device__ double __nv_log(double __a);
-__device__ float __nv_logf(float __a);
-__device__ double __nv_longlong_as_double(long long __a);
-__device__ int __nv_max(int __a, int __b);
-__device__ int __nv_min(int __a, int __b);
-__device__ double __nv_modf(double __a, double *__b);
-__device__ float __nv_modff(float __a, float *__b);
-__device__ int __nv_mul24(int __a, int __b);
-__device__ long long __nv_mul64hi(long long __a, long long __b);
-__device__ int __nv_mulhi(int __a, int __b);
-__device__ double __nv_nan(const signed char *__a);
-__device__ float __nv_nanf(const signed char *__a);
-__device__ double __nv_nearbyint(double __a);
-__device__ float __nv_nearbyintf(float __a);
-__device__ double __nv_nextafter(double __a, double __b);
-__device__ float __nv_nextafterf(float __a, float __b);
-__device__ double __nv_norm3d(double __a, double __b, double __c);
-__device__ float __nv_norm3df(float __a, float __b, float __c);
-__device__ double __nv_norm4d(double __a, double __b, double __c, double __d);
-__device__ float __nv_norm4df(float __a, float __b, float __c, float __d);
-__device__ double __nv_normcdf(double __a);
-__device__ float __nv_normcdff(float __a);
-__device__ double __nv_normcdfinv(double __a);
-__device__ float __nv_normcdfinvf(float __a);
-__device__ float __nv_normf(int __a, const float *__b);
-__device__ double __nv_norm(int __a, const double *__b);
-__device__ int __nv_popc(int __a);
-__device__ int __nv_popcll(long long __a);
-__device__ double __nv_pow(double __a, double __b);
-__device__ float __nv_powf(float __a, float __b);
-__device__ double __nv_powi(double __a, int __b);
-__device__ float __nv_powif(float __a, int __b);
-__device__ double __nv_rcbrt(double __a);
-__device__ float __nv_rcbrtf(float __a);
-__device__ double __nv_rcp64h(double __a);
-__device__ double __nv_remainder(double __a, double __b);
-__device__ float __nv_remainderf(float __a, float __b);
-__device__ double __nv_remquo(double __a, double __b, int *__c);
-__device__ float __nv_remquof(float __a, float __b, int *__c);
-__device__ int __nv_rhadd(int __a, int __b);
-__device__ double __nv_rhypot(double __a, double __b);
-__device__ float __nv_rhypotf(float __a, float __b);
-__device__ double __nv_rint(double __a);
-__device__ float __nv_rintf(float __a);
-__device__ double __nv_rnorm3d(double __a, double __b, double __c);
-__device__ float __nv_rnorm3df(float __a, float __b, float __c);
-__device__ double __nv_rnorm4d(double __a, double __b, double __c, double __d);
-__device__ float __nv_rnorm4df(float __a, float __b, float __c, float __d);
-__device__ float __nv_rnormf(int __a, const float *__b);
-__device__ double __nv_rnorm(int __a, const double *__b);
-__device__ double __nv_round(double __a);
-__device__ float __nv_roundf(float __a);
-__device__ double __nv_rsqrt(double __a);
-__device__ float __nv_rsqrtf(float __a);
-__device__ int __nv_sad(int __a, int __b, int __c);
-__device__ float __nv_saturatef(float __a);
-__device__ double __nv_scalbn(double __a, int __b);
-__device__ float __nv_scalbnf(float __a, int __b);
-__device__ int __nv_signbitd(double __a);
-__device__ int __nv_signbitf(float __a);
-__device__ void __nv_sincos(double __a, double *__b, double *__c);
-__device__ void __nv_sincosf(float __a, float *__b, float *__c);
-__device__ void __nv_sincospi(double __a, double *__b, double *__c);
-__device__ void __nv_sincospif(float __a, float *__b, float *__c);
-__device__ double __nv_sin(double __a);
-__device__ float __nv_sinf(float __a);
-__device__ double __nv_sinh(double __a);
-__device__ float __nv_sinhf(float __a);
-__device__ double __nv_sinpi(double __a);
-__device__ float __nv_sinpif(float __a);
-__device__ double __nv_sqrt(double __a);
-__device__ float __nv_sqrtf(float __a);
-__device__ double __nv_tan(double __a);
-__device__ float __nv_tanf(float __a);
-__device__ double __nv_tanh(double __a);
-__device__ float __nv_tanhf(float __a);
-__device__ double __nv_tgamma(double __a);
-__device__ float __nv_tgammaf(float __a);
-__device__ double __nv_trunc(double __a);
-__device__ float __nv_truncf(float __a);
-__device__ int __nv_uhadd(unsigned int __a, unsigned int __b);
-__device__ double __nv_uint2double_rn(unsigned int __i);
-__device__ float __nv_uint2float_rd(unsigned int __a);
-__device__ float __nv_uint2float_rn(unsigned int __a);
-__device__ float __nv_uint2float_ru(unsigned int __a);
-__device__ float __nv_uint2float_rz(unsigned int __a);
-__device__ float __nv_uint_as_float(unsigned int __a);
-__device__ double __nv_ull2double_rd(unsigned long long __a);
-__device__ double __nv_ull2double_rn(unsigned long long __a);
-__device__ double __nv_ull2double_ru(unsigned long long __a);
-__device__ double __nv_ull2double_rz(unsigned long long __a);
-__device__ float __nv_ull2float_rd(unsigned long long __a);
-__device__ float __nv_ull2float_rn(unsigned long long __a);
-__device__ float __nv_ull2float_ru(unsigned long long __a);
-__device__ float __nv_ull2float_rz(unsigned long long __a);
-__device__ unsigned long long __nv_ullmax(unsigned long long __a,
+#if defined(_OPENMP)
+#define __DEVICE__
+#elif defined(__CUDA__)
+#define __DEVICE__ __device__
+#endif
+
+__DEVICE__ int __nv_abs(int __a);
+__DEVICE__ double __nv_acos(double __a);
+__DEVICE__ float __nv_acosf(float __a);
+__DEVICE__ double __nv_acosh(double __a);
+__DEVICE__ float __nv_acoshf(float __a);
+__DEVICE__ double __nv_asin(double __a);
+__DEVICE__ float __nv_asinf(float __a);
+__DEVICE__ double __nv_asinh(double __a);
+__DEVICE__ float __nv_asinhf(float __a);
+__DEVICE__ double __nv_atan2(double __a, double __b);
+__DEVICE__ float __nv_atan2f(float __a, float __b);
+__DEVICE__ double __nv_atan(double __a);
+__DEVICE__ float __nv_atanf(float __a);
+__DEVICE__ double __nv_atanh(double __a);
+__DEVICE__ float __nv_atanhf(float __a);
+__DEVICE__ int __nv_brev(int __a);
+__DEVICE__ long long __nv_brevll(long long __a);
+__DEVICE__ int __nv_byte_perm(int __a, int __b, int __c);
+__DEVICE__ double __nv_cbrt(double __a);
+__DEVICE__ float __nv_cbrtf(float __a);
+__DEVICE__ double __nv_ceil(double __a);
+__DEVICE__ float __nv_ceilf(float __a);
+__DEVICE__ int __nv_clz(int __a);
+__DEVICE__ int __nv_clzll(long long __a);
+__DEVICE__ double __nv_copysign(double __a, double __b);
+__DEVICE__ float __nv_copysignf(float __a, float __b);
+__DEVICE__ double __nv_cos(double __a);
+__DEVICE__ float __nv_cosf(float __a);
+__DEVICE__ double __nv_cosh(double __a);
+__DEVICE__ float __nv_coshf(float __a);
+__DEVICE__ double __nv_cospi(double __a);
+__DEVICE__ float __nv_cospif(float __a);
+__DEVICE__ double __nv_cyl_bessel_i0(double __a);
+__DEVICE__ float __nv_cyl_bessel_i0f(float __a);
+__DEVICE__ double __nv_cyl_bessel_i1(double __a);
+__DEVICE__ float __nv_cyl_bessel_i1f(float __a);
+__DEVICE__ double __nv_dadd_rd(double __a, double __b);
+__DEVICE__ double __nv_dadd_rn(double __a, double __b);
+__DEVICE__ double __nv_dadd_ru(double __a, double __b);
+__DEVICE__ double __nv_dadd_rz(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rd(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rn(double __a, double __b);
+__DEVICE__ double __nv_ddiv_ru(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rz(double __a, double __b);
+__DEVICE__ double __nv_dmul_rd(double __a, double __b);
+__DEVICE__ double __nv_dmul_rn(double __a, double __b);
+__DEVICE__ double __nv_dmul_ru(double __a, double __b);
+__DEVICE__ double __nv_dmul_rz(double __a, double __b);
+__DEVICE__ float __nv_double2float_rd(double __a);
+__DEVICE__ float __nv_double2float_rn(double __a);
+__DEVICE__ float __nv_double2float_ru(double __a);
+__DEVICE__ float __nv_double2float_rz(double __a);
+__DEVICE__ int __nv_double2hiint(double __a);
+__DEVICE__ int __nv_double2int_rd(double __a);
+__DEVICE__ int __nv_double2int_rn(double __a);
+__DEVICE__ int __nv_double2int_ru(double __a);
+__DEVICE__ int __nv_double2int_rz(double __a);
+__DEVICE__ long long __nv_double2ll_rd(double __a);
+__DEVICE__ long long __nv_double2ll_rn(double __a);
+__DEVICE__ long long __nv_double2ll_ru(double __a);
+__DEVICE__ long long __nv_double2ll_rz(double __a);
+__DEVICE__ int __nv_double2loint(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rd(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rn(double __a);
+__DEVICE__ unsigned int __nv_double2uint_ru(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rz(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rd(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rn(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_ru(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rz(double __a);
+__DEVICE__ unsigned long long __nv_double_as_longlong(double __a);
+__DEVICE__ double __nv_drcp_rd(double __a);
+__DEVICE__ double __nv_drcp_rn(double __a);
+__DEVICE__ double __nv_drcp_ru(double __a);
+__DEVICE__ double __nv_drcp_rz(double __a);
+__DEVICE__ double __nv_dsqrt_rd(double __a);
+__DEVICE__ double __nv_dsqrt_rn(double __a);
+__DEVICE__ double __nv_dsqrt_ru(double __a);
+__DEVICE__ double __nv_dsqrt_rz(double __a);
+__DEVICE__ double __nv_dsub_rd(double __a, double __b);
+__DEVICE__ double __nv_dsub_rn(double __a, double __b);
+__DEVICE__ double __nv_dsub_ru(double __a, double __b);
+__DEVICE__ double __nv_dsub_rz(double __a, double __b);
+__DEVICE__ double __nv_erfc(double __a);
+__DEVICE__ float __nv_erfcf(float __a);
+__DEVICE__ double __nv_erfcinv(double __a);
+__DEVICE__ float __nv_erfcinvf(float __a);
+__DEVICE__ double __nv_erfcx(double __a);
+__DEVICE__ float __nv_erfcxf(float __a);
+__DEVICE__ double __nv_erf(double __a);
+__DEVICE__ float __nv_erff(float __a);
+__DEVICE__ double __nv_erfinv(double __a);
+__DEVICE__ float __nv_erfinvf(float __a);
+__DEVICE__ double __nv_exp10(double __a);
+__DEVICE__ float __nv_exp10f(float __a);
+__DEVICE__ double __nv_exp2(double __a);
+__DEVICE__ float __nv_exp2f(float __a);
+__DEVICE__ double __nv_exp(double __a);
+__DEVICE__ float __nv_expf(float __a);
+__DEVICE__ double __nv_expm1(double __a);
+__DEVICE__ float __nv_expm1f(float __a);
+__DEVICE__ double __nv_fabs(double __a);
+__DEVICE__ float __nv_fabsf(float __a);
+__DEVICE__ float __nv_fadd_rd(float __a, float __b);
+__DEVICE__ float __nv_fadd_rn(float __a, float __b);
+__DEVICE__ float __nv_fadd_ru(float __a, float __b);
+__DEVICE__ float __nv_fadd_rz(float __a, float __b);
+__DEVICE__ float __nv_fast_cosf(float __a);
+__DEVICE__ float __nv_fast_exp10f(float __a);
+__DEVICE__ float __nv_fast_expf(float __a);
+__DEVICE__ float __nv_fast_fdividef(float __a, float __b);
+__DEVICE__ float __nv_fast_log10f(float __a);
+__DEVICE__ float __nv_fast_log2f(float __a);
+__DEVICE__ float __nv_fast_logf(float __a);
+__DEVICE__ float __nv_fast_powf(float __a, float __b);
+__DEVICE__ void __nv_fast_sincosf(float __a, float *__s, float *__c);
+__DEVICE__ float __nv_fast_sinf(float __a);
+__DEVICE__ float __nv_fast_tanf(float __a);
+__DEVICE__ double __nv_fdim(double __a, double __b);
+__DEVICE__ float __nv_fdimf(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rd(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rn(float __a, float __b);
+__DEVICE__ float __nv_fdiv_ru(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rz(float __a, float __b);
+__DEVICE__ int __nv_ffs(int __a);
+__DEVICE__ int __nv_ffsll(long long __a);
+__DEVICE__ int __nv_finitef(float __a);
+__DEVICE__ unsigned short __nv_float2half_rn(float __a);
+__DEVICE__ int __nv_float2int_rd(float __a);
+__DEVICE__ int __nv_float2int_rn(float __a);
+__DEVICE__ int __nv_float2int_ru(float __a);
+__DEVICE__ int __nv_float2int_rz(float __a);
+__DEVICE__ long long __nv_float2ll_rd(float __a);
+__DEVICE__ long long __nv_float2ll_rn(float __a);
+__DEVICE__ long long __nv_float2ll_ru(float __a);
+__DEVICE__ long long __nv_float2ll_rz(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rd(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rn(float __a);
+__DEVICE__ unsigned int __nv_float2uint_ru(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rz(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rd(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rn(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_ru(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rz(float __a);
+__DEVICE__ int __nv_float_as_int(float __a);
+__DEVICE__ unsigned int __nv_float_as_uint(float __a);
+__DEVICE__ double __nv_floor(double __a);
+__DEVICE__ float __nv_floorf(float __a);
+__DEVICE__ double __nv_fma(double __a, double __b, double __c);
+__DEVICE__ float __nv_fmaf(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rd(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rn(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_ru(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rz(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rd(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rn(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ru(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rz(float __a, float __b, float __c);
+__DEVICE__ double __nv_fma_rd(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_rn(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_ru(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_rz(double __a, double __b, double __c);
+__DEVICE__ double __nv_fmax(double __a, double __b);
+__DEVICE__ float __nv_fmaxf(float __a, float __b);
+__DEVICE__ double __nv_fmin(double __a, double __b);
+__DEVICE__ float __nv_fminf(float __a, float __b);
+__DEVICE__ double __nv_fmod(double __a, double __b);
+__DEVICE__ float __nv_fmodf(float __a, float __b);
+__DEVICE__ float __nv_fmul_rd(float __a, float __b);
+__DEVICE__ float __nv_fmul_rn(float __a, float __b);
+__DEVICE__ float __nv_fmul_ru(float __a, float __b);
+__DEVICE__ float __nv_fmul_rz(float __a, float __b);
+__DEVICE__ float __nv_frcp_rd(float __a);
+__DEVICE__ float __nv_frcp_rn(float __a);
+__DEVICE__ float __nv_frcp_ru(float __a);
+__DEVICE__ float __nv_frcp_rz(float __a);
+__DEVICE__ double __nv_frexp(double __a, int *__b);
+__DEVICE__ float __nv_frexpf(float __a, int *__b);
+__DEVICE__ float __nv_frsqrt_rn(float __a);
+__DEVICE__ float __nv_fsqrt_rd(float __a);
+__DEVICE__ float __nv_fsqrt_rn(float __a);
+__DEVICE__ float __nv_fsqrt_ru(float __a);
+__DEVICE__ float __nv_fsqrt_rz(float __a);
+__DEVICE__ float __nv_fsub_rd(float __a, float __b);
+__DEVICE__ float __nv_fsub_rn(float __a, float __b);
+__DEVICE__ float __nv_fsub_ru(float __a, float __b);
+__DEVICE__ float __nv_fsub_rz(float __a, float __b);
+__DEVICE__ int __nv_hadd(int __a, int __b);
+__DEVICE__ float __nv_half2float(unsigned short __h);
+__DEVICE__ double __nv_hiloint2double(int __a, int __b);
+__DEVICE__ double __nv_hypot(double __a, double __b);
+__DEVICE__ float __nv_hypotf(float __a, float __b);
+__DEVICE__ int __nv_ilogb(double __a);
+__DEVICE__ int __nv_ilogbf(float __a);
+__DEVICE__ double __nv_int2double_rn(int __a);
+__DEVICE__ float __nv_int2float_rd(int __a);
+__DEVICE__ float __nv_int2float_rn(int __a);
+__DEVICE__ float __nv_int2float_ru(int __a);
+__DEVICE__ float __nv_int2float_rz(int __a);
+__DEVICE__ float __nv_int_as_float(int __a);
+__DEVICE__ int __nv_isfinited(double __a);
+__DEVICE__ int __nv_isinfd(double __a);
+__DEVICE__ int __nv_isinff(float __a);
+__DEVICE__ int __nv_isnand(double __a);
+__DEVICE__ int __nv_isnanf(float __a);
+__DEVICE__ double __nv_j0(double __a);
+__DEVICE__ float __nv_j0f(float __a);
+__DEVICE__ double __nv_j1(double __a);
+__DEVICE__ float __nv_j1f(float __a);
+__DEVICE__ float __nv_jnf(int __a, float __b);
+__DEVICE__ double __nv_jn(int __a, double __b);
+__DEVICE__ double __nv_ldexp(double __a, int __b);
+__DEVICE__ float __nv_ldexpf(float __a, int __b);
+__DEVICE__ double __nv_lgamma(double __a);
+__DEVICE__ float __nv_lgammaf(float __a);
+__DEVICE__ double __nv_ll2double_rd(long long __a);
+__DEVICE__ double __nv_ll2double_rn(long long __a);
+__DEVICE__ double __nv_ll2double_ru(long long __a);
+__DEVICE__ double __nv_ll2double_rz(long long __a);
+__DEVICE__ float __nv_ll2float_rd(long long __a);
+__DEVICE__ float __nv_ll2float_rn(long long __a);
+__DEVICE__ float __nv_ll2float_ru(long long __a);
+__DEVICE__ float __nv_ll2float_rz(long long __a);
+__DEVICE__ long long __nv_llabs(long long __a);
+__DEVICE__ long long __nv_llmax(long long __a, long long __b);
+__DEVICE__ long long __nv_llmin(long long __a, long long __b);
+__DEVICE__ long long __nv_llrint(double __a);
+__DEVICE__ long long __nv_llrintf(float __a);
+__DEVICE__ long long __nv_llround(double __a);
+__DEVICE__ long long __nv_llroundf(float __a);
+__DEVICE__ double __nv_log10(double __a);
+__DEVICE__ float __nv_log10f(float __a);
+__DEVICE__ double __nv_log1p(double __a);
+__DEVICE__ float __nv_log1pf(float __a);
+__DEVICE__ double __nv_log2(double __a);
+__DEVICE__ float __nv_log2f(float __a);
+__DEVICE__ double __nv_logb(double __a);
+__DEVICE__ float __nv_logbf(float __a);
+__DEVICE__ double __nv_log(double __a);
+__DEVICE__ float __nv_logf(float __a);
+__DEVICE__ double __nv_longlong_as_double(long long __a);
+__DEVICE__ int __nv_max(int __a, int __b);
+__DEVICE__ int __nv_min(int __a, int __b);
+__DEVICE__ double __nv_modf(double __a, double *__b);
+__DEVICE__ float __nv_modff(float __a, float *__b);
+__DEVICE__ int __nv_mul24(int __a, int __b);
+__DEVICE__ long long __nv_mul64hi(long long __a, long long __b);
+__DEVICE__ int __nv_mulhi(int __a, int __b);
+__DEVICE__ double __nv_nan(const signed char *__a);
+__DEVICE__ float __nv_nanf(const signed char *__a);
+__DEVICE__ double __nv_nearbyint(double __a);
+__DEVICE__ float __nv_nearbyintf(float __a);
+__DEVICE__ double __nv_nextafter(double __a, double __b);
+__DEVICE__ float __nv_nextafterf(float __a, float __b);
+__DEVICE__ double __nv_norm3d(double __a, double __b, double __c);
+__DEVICE__ float __nv_norm3df(float __a, float __b, float __c);
+__DEVICE__ double __nv_norm4d(double __a, double __b, double __c, double __d);
+__DEVICE__ float __nv_norm4df(float __a, float __b, float __c, float __d);
+__DEVICE__ double __nv_normcdf(double __a);
+__DEVICE__ float __nv_normcdff(float __a);
+__DEVICE__ double __nv_normcdfinv(double __a);
+__DEVICE__ float __nv_normcdfinvf(float __a);
+__DEVICE__ float __nv_normf(int __a, const float *__b);
+__DEVICE__ double __nv_norm(int __a, const double *__b);
+__DEVICE__ int __nv_popc(int __a);
+__DEVICE__ int __nv_popcll(long long __a);
+__DEVICE__ double __nv_pow(double __a, double __b);
+__DEVICE__ float __nv_powf(float __a, float __b);
+__DEVICE__ double __nv_powi(double __a, int __b);
+__DEVICE__ float __nv_powif(float __a, int __b);
+__DEVICE__ double __nv_rcbrt(double __a);
+__DEVICE__ float __nv_rcbrtf(float __a);
+__DEVICE__ double __nv_rcp64h(double __a);
+__DEVICE__ double __nv_remainder(double __a, double __b);
+__DEVICE__ float __nv_remainderf(float __a, float __b);
+__DEVICE__ double __nv_remquo(double __a, double __b, int *__c);
+__DEVICE__ float __nv_remquof(float __a, float __b, int *__c);
+__DEVICE__ int __nv_rhadd(int __a, int __b);
+__DEVICE__ double __nv_rhypot(double __a, double __b);
+__DEVICE__ float __nv_rhypotf(float __a, float __b);
+__DEVICE__ double __nv_rint(double __a);
+__DEVICE__ float __nv_rintf(float __a);
+__DEVICE__ double __nv_rnorm3d(double __a, double __b, double __c);
+__DEVICE__ float __nv_rnorm3df(float __a, float __b, float __c);
+__DEVICE__ double __nv_rnorm4d(double __a, double __b, double __c, double __d);
+__DEVICE__ float __nv_rnorm4df(float __a, float __b, float __c, float __d);
+__DEVICE__ float __nv_rnormf(int __a, const float *__b);
+__DEVICE__ double __nv_rnorm(int __a, const double *__b);
+__DEVICE__ double __nv_round(double __a);
+__DEVICE__ float __nv_roundf(float __a);
+__DEVICE__ double __nv_rsqrt(double __a);
+__DEVICE__ float __nv_rsqrtf(float __a);
+__DEVICE__ int __nv_sad(int __a, int __b, int __c);
+__DEVICE__ float __nv_saturatef(float __a);
+__DEVICE__ double __nv_scalbn(double __a, int __b);
+__DEVICE__ float __nv_scalbnf(float __a, int __b);
+__DEVICE__ int __nv_signbitd(double __a);
+__DEVICE__ int __nv_signbitf(float __a);
+__DEVICE__ void __nv_sincos(double __a, double *__b, double *__c);
+__DEVICE__ void __nv_sincosf(float __a, float *__b, float *__c);
+__DEVICE__ void __nv_sincospi(double __a, double *__b, double *__c);
+__DEVICE__ void __nv_sincospif(float __a, float *__b, float *__c);
+__DEVICE__ double __nv_sin(double __a);
+__DEVICE__ float __nv_sinf(float __a);
+__DEVICE__ double __nv_sinh(double __a);
+__DEVICE__ float __nv_sinhf(float __a);
+__DEVICE__ double __nv_sinpi(double __a);
+__DEVICE__ float __nv_sinpif(float __a);
+__DEVICE__ double __nv_sqrt(double __a);
+__DEVICE__ float __nv_sqrtf(float __a);
+__DEVICE__ double __nv_tan(double __a);
+__DEVICE__ float __nv_tanf(float __a);
+__DEVICE__ double __nv_tanh(double __a);
+__DEVICE__ float __nv_tanhf(float __a);
+__DEVICE__ double __nv_tgamma(double __a);
+__DEVICE__ float __nv_tgammaf(float __a);
+__DEVICE__ double __nv_trunc(double __a);
+__DEVICE__ float __nv_truncf(float __a);
+__DEVICE__ int __nv_uhadd(unsigned int __a, unsigned int __b);
+__DEVICE__ double __nv_uint2double_rn(unsigned int __i);
+__DEVICE__ float __nv_uint2float_rd(unsigned int __a);
+__DEVICE__ float __nv_uint2float_rn(unsigned int __a);
+__DEVICE__ float __nv_uint2float_ru(unsigned int __a);
+__DEVICE__ float __nv_uint2float_rz(unsigned int __a);
+__DEVICE__ float __nv_uint_as_float(unsigned int __a);
+__DEVICE__ double __nv_ull2double_rd(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_rn(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_ru(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_rz(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rd(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rn(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_ru(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rz(unsigned long long __a);
+__DEVICE__ unsigned long long __nv_ullmax(unsigned long long __a,
                                           unsigned long long __b);
-__device__ unsigned long long __nv_ullmin(unsigned long long __a,
+__DEVICE__ unsigned long long __nv_ullmin(unsigned long long __a,
                                           unsigned long long __b);
-__device__ unsigned int __nv_umax(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_umin(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_umul24(unsigned int __a, unsigned int __b);
-__device__ unsigned long long __nv_umul64hi(unsigned long long __a,
+__DEVICE__ unsigned int __nv_umax(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_umin(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_umul24(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned long long __nv_umul64hi(unsigned long long __a,
                                             unsigned long long __b);
-__device__ unsigned int __nv_umulhi(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_urhadd(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_usad(unsigned int __a, unsigned int __b,
+__DEVICE__ unsigned int __nv_umulhi(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_urhadd(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_usad(unsigned int __a, unsigned int __b,
                                   unsigned int __c);
 #if CUDA_VERSION >= 9000 && CUDA_VERSION < 9020
-__device__ int __nv_vabs2(int __a);
-__device__ int __nv_vabs4(int __a);
-__device__ int __nv_vabsdiffs2(int __a, int __b);
-__device__ int __nv_vabsdiffs4(int __a, int __b);
-__device__ int __nv_vabsdiffu2(int __a, int __b);
-__device__ int __nv_vabsdiffu4(int __a, int __b);
-__device__ int __nv_vabsss2(int __a);
-__device__ int __nv_vabsss4(int __a);
-__device__ int __nv_vadd2(int __a, int __b);
-__device__ int __nv_vadd4(int __a, int __b);
-__device__ int __nv_vaddss2(int __a, int __b);
-__device__ int __nv_vaddss4(int __a, int __b);
-__device__ int __nv_vaddus2(int __a, int __b);
-__device__ int __nv_vaddus4(int __a, int __b);
-__device__ int __nv_vavgs2(int __a, int __b);
-__device__ int __nv_vavgs4(int __a, int __b);
-__device__ int __nv_vavgu2(int __a, int __b);
-__device__ int __nv_vavgu4(int __a, int __b);
-__device__ int __nv_vcmpeq2(int __a, int __b);
-__device__ int __nv_vcmpeq4(int __a, int __b);
-__device__ int __nv_vcmpges2(int __a, int __b);
-__device__ int __nv_vcmpges4(int __a, int __b);
-__device__ int __nv_vcmpgeu2(int __a, int __b);
-__device__ int __nv_vcmpgeu4(int __a, int __b);
-__device__ int __nv_vcmpgts2(int __a, int __b);
-__device__ int __nv_vcmpgts4(int __a, int __b);
-__device__ int __nv_vcmpgtu2(int __a, int __b);
-__device__ int __nv_vcmpgtu4(int __a, int __b);
-__device__ int __nv_vcmples2(int __a, int __b);
-__device__ int __nv_vcmples4(int __a, int __b);
-__device__ int __nv_vcmpleu2(int __a, int __b);
-__device__ int __nv_vcmpleu4(int __a, int __b);
-__device__ int __nv_vcmplts2(int __a, int __b);
-__device__ int __nv_vcmplts4(int __a, int __b);
-__device__ int __nv_vcmpltu2(int __a, int __b);
-__device__ int __nv_vcmpltu4(int __a, int __b);
-__device__ int __nv_vcmpne2(int __a, int __b);
-__device__ int __nv_vcmpne4(int __a, int __b);
-__device__ int __nv_vhaddu2(int __a, int __b);
-__device__ int __nv_vhaddu4(int __a, int __b);
-__device__ int __nv_vmaxs2(int __a, int __b);
-__device__ int __nv_vmaxs4(int __a, int __b);
-__device__ int __nv_vmaxu2(int __a, int __b);
-__device__ int __nv_vmaxu4(int __a, int __b);
-__device__ int __nv_vmins2(int __a, int __b);
-__device__ int __nv_vmins4(int __a, int __b);
-__device__ int __nv_vminu2(int __a, int __b);
-__device__ int __nv_vminu4(int __a, int __b);
-__device__ int __nv_vneg2(int __a);
-__device__ int __nv_vneg4(int __a);
-__device__ int __nv_vnegss2(int __a);
-__device__ int __nv_vnegss4(int __a);
-__device__ int __nv_vsads2(int __a, int __b);
-__device__ int __nv_vsads4(int __a, int __b);
-__device__ int __nv_vsadu2(int __a, int __b);
-__device__ int __nv_vsadu4(int __a, int __b);
-__device__ int __nv_vseteq2(int __a, int __b);
-__device__ int __nv_vseteq4(int __a, int __b);
-__device__ int __nv_vsetges2(int __a, int __b);
-__device__ int __nv_vsetges4(int __a, int __b);
-__device__ int __nv_vsetgeu2(int __a, int __b);
-__device__ int __nv_vsetgeu4(int __a, int __b);
-__device__ int __nv_vsetgts2(int __a, int __b);
-__device__ int __nv_vsetgts4(int __a, int __b);
-__device__ int __nv_vsetgtu2(int __a, int __b);
-__device__ int __nv_vsetgtu4(int __a, int __b);
-__device__ int __nv_vsetles2(int __a, int __b);
-__device__ int __nv_vsetles4(int __a, int __b);
-__device__ int __nv_vsetleu2(int __a, int __b);
-__device__ int __nv_vsetleu4(int __a, int __b);
-__device__ int __nv_vsetlts2(int __a, int __b);
-__device__ int __nv_vsetlts4(int __a, int __b);
-__device__ int __nv_vsetltu2(int __a, int __b);
-__device__ int __nv_vsetltu4(int __a, int __b);
-__device__ int __nv_vsetne2(int __a, int __b);
-__device__ int __nv_vsetne4(int __a, int __b);
-__device__ int __nv_vsub2(int __a, int __b);
-__device__ int __nv_vsub4(int __a, int __b);
-__device__ int __nv_vsubss2(int __a, int __b);
-__device__ int __nv_vsubss4(int __a, int __b);
-__device__ int __nv_vsubus2(int __a, int __b);
-__device__ int __nv_vsubus4(int __a, int __b);
+__DEVICE__ int __nv_vabs2(int __a);
+__DEVICE__ int __nv_vabs4(int __a);
+__DEVICE__ int __nv_vabsdiffs2(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffs4(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffu2(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffu4(int __a, int __b);
+__DEVICE__ int __nv_vabsss2(int __a);
+__DEVICE__ int __nv_vabsss4(int __a);
+__DEVICE__ int __nv_vadd2(int __a, int __b);
+__DEVICE__ int __nv_vadd4(int __a, int __b);
+__DEVICE__ int __nv_vaddss2(int __a, int __b);
+__DEVICE__ int __nv_vaddss4(int __a, int __b);
+__DEVICE__ int __nv_vaddus2(int __a, int __b);
+__DEVICE__ int __nv_vaddus4(int __a, int __b);
+__DEVICE__ int __nv_vavgs2(int __a, int __b);
+__DEVICE__ int __nv_vavgs4(int __a, int __b);
+__DEVICE__ int __nv_vavgu2(int __a, int __b);
+__DEVICE__ int __nv_vavgu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpeq2(int __a, int __b);
+__DEVICE__ int __nv_vcmpeq4(int __a, int __b);
+__DEVICE__ int __nv_vcmpges2(int __a, int __b);
+__DEVICE__ int __nv_vcmpges4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgeu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgeu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgts2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgts4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgtu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgtu4(int __a, int __b);
+__DEVICE__ int __nv_vcmples2(int __a, int __b);
+__DEVICE__ int __nv_vcmples4(int __a, int __b);
+__DEVICE__ int __nv_vcmpleu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpleu4(int __a, int __b);
+__DEVICE__ int __nv_vcmplts2(int __a, int __b);
+__DEVICE__ int __nv_vcmplts4(int __a, int __b);
+__DEVICE__ int __nv_vcmpltu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpltu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpne2(int __a, int __b);
+__DEVICE__ int __nv_vcmpne4(int __a, int __b);
+__DEVICE__ int __nv_vhaddu2(int __a, int __b);
+__DEVICE__ int __nv_vhaddu4(int __a, int __b);
+__DEVICE__ int __nv_vmaxs2(int __a, int __b);
+__DEVICE__ int __nv_vmaxs4(int __a, int __b);
+__DEVICE__ int __nv_vmaxu2(int __a, int __b);
+__DEVICE__ int __nv_vmaxu4(int __a, int __b);
+__DEVICE__ int __nv_vmins2(int __a, int __b);
+__DEVICE__ int __nv_vmins4(int __a, int __b);
+__DEVICE__ int __nv_vminu2(int __a, int __b);
+__DEVICE__ int __nv_vminu4(int __a, int __b);
+__DEVICE__ int __nv_vneg2(int __a);
+__DEVICE__ int __nv_vneg4(int __a);
+__DEVICE__ int __nv_vnegss2(int __a);
+__DEVICE__ int __nv_vnegss4(int __a);
+__DEVICE__ int __nv_vsads2(int __a, int __b);
+__DEVICE__ int __nv_vsads4(int __a, int __b);
+__DEVICE__ int __nv_vsadu2(int __a, int __b);
+__DEVICE__ int __nv_vsadu4(int __a, int __b);
+__DEVICE__ int __nv_vseteq2(int __a, int __b);
+__DEVICE__ int __nv_vseteq4(int __a, int __b);
+__DEVICE__ int __nv_vsetges2(int __a, int __b);
+__DEVICE__ int __nv_vsetges4(int __a, int __b);
+__DEVICE__ int __nv_vsetgeu2(int __a, int __b);
+__DEVICE__ int __nv_vsetgeu4(int __a, int __b);
+__DEVICE__ int __nv_vsetgts2(int __a, int __b);
+__DEVICE__ int __nv_vsetgts4(int __a, int __b);
+__DEVICE__ int __nv_vsetgtu2(int __a, int __b);
+__DEVICE__ int __nv_vsetgtu4(int __a, int __b);
+__DEVICE__ int __nv_vsetles2(int __a, int __b);
+__DEVICE__ int __nv_vsetles4(int __a, int __b);
+__DEVICE__ int __nv_vsetleu2(int __a, int __b);
+__DEVICE__ int __nv_vsetleu4(int __a, int __b);
+__DEVICE__ int __nv_vsetlts2(int __a, int __b);
+__DEVICE__ int __nv_vsetlts4(int __a, int __b);
+__DEVICE__ int __nv_vsetltu2(int __a, int __b);
+__DEVICE__ int __nv_vsetltu4(int __a, int __b);
+__DEVICE__ int __nv_vsetne2(int __a, int __b);
+__DEVICE__ int __nv_vsetne4(int __a, int __b);
+__DEVICE__ int __nv_vsub2(int __a, int __b);
+__DEVICE__ int __nv_vsub4(int __a, int __b);
+__DEVICE__ int __nv_vsubss2(int __a, int __b);
+__DEVICE__ int __nv_vsubss4(int __a, int __b);
+__DEVICE__ int __nv_vsubus2(int __a, int __b);
+__DEVICE__ int __nv_vsubus4(int __a, int __b);
 #endif  // CUDA_VERSION
-__device__ double __nv_y0(double __a);
-__device__ float __nv_y0f(float __a);
-__device__ double __nv_y1(double __a);
-__device__ float __nv_y1f(float __a);
-__device__ float __nv_ynf(int __a, float __b);
-__device__ double __nv_yn(int __a, double __b);
+__DEVICE__ double __nv_y0(double __a);
+__DEVICE__ float __nv_y0f(float __a);
+__DEVICE__ double __nv_y1(double __a);
+__DEVICE__ float __nv_y1f(float __a);
+__DEVICE__ float __nv_ynf(int __a, float __b);
+__DEVICE__ double __nv_yn(int __a, double __b);
+#if defined(__cplusplus)
 } // extern "C"
+#endif
 #endif // __CLANG_CUDA_LIBDEVICE_DECLARES_H__

lib/Headers/__clang_cuda_math_forward_declares.h

 // This file forward-declares of some math functions we (or the CUDA headers)
 // will define later.  We need to do this, and do it before cmath is included,
 // because the standard library may have constexpr math functions.  In the
 // absence of a prior __device__ decl, those constexpr functions may become
 // implicitly host+device.  host+device functions can't be overloaded, so that
 // would preclude the use of our own __device__ overloads for these functions.
 
 #pragma push_macro("__DEVICE__")
+#ifdef _OPENMP
+#define __DEVICE__ static __inline__ __attribute__((always_inline))
+#else
 #define __DEVICE__                                                             \
   static __inline__ __attribute__((always_inline)) __attribute__((device))
+#endif
 
 __DEVICE__ double abs(double);
 __DEVICE__ float abs(float);
 __DEVICE__ int abs(int);
 __DEVICE__ long abs(long);
 __DEVICE__ long long abs(long long);
 __DEVICE__ double acos(double);
 __DEVICE__ float acos(float);

lib/Headers/openmp_wrappers/__clang_openmp_math.h

+/*===---- __clang_openmp_math.h - OpenMP target math support ---------------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#if defined(__NVPTX__) && defined(_OPENMP)
+/// TODO:
+/// We are currently reusing the functionality of the Clang-CUDA code path
+/// as an alternative to the host declarations provided by math.h and cmath.
+/// This is suboptimal.
+///
+/// We should instead declare the device functions in a similar way, e.g.,
+/// through OpenMP 5.0 variants, and afterwards populate the module with the
+/// host declarations by unconditionally including the host math.h or cmath,
+/// respectively. This is actually what the Clang-CUDA code path does, using
+/// __device__ instead of variants to avoid redeclarations and get the desired
+/// overload resolution.
+
+#define __CUDA__
+
+#if defined(__cplusplus)
+  #include <__clang_cuda_math_forward_declares.h>
+#endif
+
+/// Include declarations for libdevice functions.
+#include <__clang_cuda_libdevice_declares.h>
+/// Provide definitions for these functions.
+#include <__clang_cuda_device_functions.h>
+
+#if defined(__cplusplus)
+  #include <__clang_cuda_cmath.h>
+#endif
+
+#undef __CUDA__
+
+/// Magic macro for stopping the math.h/cmath host header from being included.
+#define __CLANG_NO_HOST_MATH__
+
+#endif
+

lib/Headers/openmp_wrappers/cmath

+/*===-------------- cmath - Alternative cmath header -----------------------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __cplusplus
+#include <__clang_openmp_math.h>
+#endif
+
+#ifndef __CLANG_NO_HOST_MATH__
+#include_next <cmath>
+#else
+#undef __CLANG_NO_HOST_MATH__
+#endif

lib/Headers/openmp_wrappers/math.h

+/*===------------- math.h - Alternative math.h header ----------------------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __cplusplus
+#include <__clang_openmp_math.h>
+#endif
+
+#ifndef __CLANG_NO_HOST_MATH__
+#include_next <math.h>
+#else
+#undef __CLANG_NO_HOST_MATH__
+#endif
+

test/Driver/openmp-offload-gpu.c

 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime -fno-openmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_FULL_RUNTIME %s
 // NO_FULL_RUNTIME-NOT: "-{{fno-|f}}openmp-cuda-force-full-runtime"
 
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-teams-reduction-recs-num=2048 2>&1 \
 // RUN:   | FileCheck -check-prefix=CUDA_RED_RECS %s
 // CUDA_RED_RECS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
 // CUDA_RED_RECS-SAME: "-fopenmp-cuda-teams-reduction-recs-num=2048"
+
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=OPENMP_NVPTX_WRAPPERS %s
+// OPENMP_NVPTX_WRAPPERS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
+// OPENMP_NVPTX_WRAPPERS-SAME: "-internal-isystem" "{{.*}}openmp_wrappers"

test/Headers/Inputs/include/cmath

+#pragma once
+
+double sqrt(double);
+double pow(double, double);
+double modf(double, double*);

test/Headers/Inputs/include/limits

+#pragma once
+
+namespace std
+{
+struct __numeric_limits_base
+  {};
+template<typename _Tp>
+  struct numeric_limits : public __numeric_limits_base
+    {};
+}

test/Headers/Inputs/include/math.h

 #pragma once
+
+double sqrt(double);
+double pow(double, double);
+double modf(double, double*);

test/Headers/nvptx_device_cmath_functions.c

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cmath>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}

test/Headers/nvptx_device_cmath_functions.cpp

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -internal-isystem %S/Inputs/include -include stdlib.h -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cmath>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}

test/Headers/nvptx_device_math_functions.c

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <math.h>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}

test/Headers/nvptx_device_math_functions.cpp

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -internal-isystem %S/Inputs/include -include stdlib.h -include limits -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <math.h>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}