Commit Graph

27 Commits

Author SHA1 Message Date
Johannes Doerfert 17c8251bca [OpenMP][CUDA][FIX] Use the new `remquo` overload only for OpenMP
CUDA buildbots complained about a redefinition when I landed D89971.
This is odd and I fail to understand where in the CUDA headers the other
definition is supposed to be. For now, given that CUDA doesn't need the
overload (AFAIKT), we simply restrict it to the OpenMP mode.
2020-10-27 23:52:59 -05:00
Johannes Doerfert b1a90e1599 [OpenMP][CUDA] Add missing overload for `remquo(float,float,int*)`
Reported by Colleen Bertoni <bertoni@anl.gov> after running the OvO test
suite: https://github.com/TApplencourt/OvO/

The template overload is still hidden behind an ifdef for OpenMP. In the
future we probably want to remove the ifdef but that requires further
testing.

Reviewed By: JonChesterfield, tra

Differential Revision: https://reviews.llvm.org/D89971
2020-10-27 19:12:51 -05:00
Johannes Doerfert 97652202d1 [OpenMP] Overload `std::isnan` and friends multiple times for the GPU
`std::isnan` and friends can be found in two variants in the wild, one
returns `bool`, as the standard defines it, one returns `int`, as the C
macros do. So far we kinda hoped the system versions of these functions
will work for people, e.g. they are definitions that can be compiled for
the target. We know that is not the case always so we leverage the
`disable_implicit_base` OpenMP context extension to specialize both
versions of these functions without causing an invalid redeclaration.

Reviewed By: JonChesterfield, tra

Differential Revision: https://reviews.llvm.org/D85879
2020-09-16 13:37:09 -05:00
Johannes Doerfert 7f1e6fcff9 [OpenMP] Use __OPENMP_NVPTX__ instead of _OPENMP in wrapper headers
Due to recent changes we cannot use OpenMP in CUDA files anymore
(PR45533) as the math handling of CUDA is different when _OPENMP is
defined. We actually want this different behavior only if we are
offloading with OpenMP to NVIDIA, thus generating NVPTX. With this patch
we do not interfere with the CUDA math handling except if we are in
NVPTX offloading mode, as indicated by the presence of __OPENMP_NVPTX__.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D78155
2020-07-10 18:53:34 -05:00
Johannes Doerfert f85ae058f5 [OpenMP] Provide math functions in OpenMP device code via OpenMP variants
For OpenMP target regions to piggy back on the CUDA/AMDGPU/... implementation of math functions,
we include the appropriate definitions inside of an `omp begin/end declare variant match(device={arch(nvptx)})` scope.
This way, the vendor specific math functions will become specialized versions of the system math functions.
When a system math function is called and specialized version is available the selection logic introduced in D75779
instead call the specialized version. In contrast to the code path we used so far, the system header is actually included.
This means functions without specialized versions are available and so are macro definitions.

This should address PR42061, PR42798, and PR42799.

Reviewed By: ye-luo

Differential Revision: https://reviews.llvm.org/D75788
2020-04-07 23:33:24 -05:00
Gheorghe-Teodor Bercea 144291e14c [OpenMP][bugfix] Add missing math functions variants for log and abs.
Summary: When including the random header in C++, some of the math functions it relies on are not present in the CUDA headers. We include this variants in this case.

Reviewers: jdoerfert, hfinkel, tra, caomhin

Reviewed By: tra

Subscribers: efriedma, guansong, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D62046

llvm-svn: 361066
2019-05-17 19:15:53 +00:00
Gheorghe-Teodor Bercea 9392bd6987 [OpenMP][Bugfix] Move double and float versions of abs under c++ macro
Summary:
This is a fix for the reported bug:

[[ https://bugs.llvm.org/show_bug.cgi?id=41861 | 41861 ]]

abs functions need to be moved under the c++ macro to avoid conflicts with included headers.

Reviewers: tra, jdoerfert, hfinkel, ABataev, caomhin

Reviewed By: jdoerfert

Subscribers: guansong, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61959

llvm-svn: 360809
2019-05-15 20:28:23 +00:00
Gheorghe-Teodor Bercea 7641f310d7 [OpenMP][bugfix] Fix issues with C++ 17 compilation when handling math functions
Summary: In OpenMP device offloading we must ensure that unde C++ 17, the inclusion of cstdlib will works correctly.

Reviewers: ABataev, tra, jdoerfert, hfinkel, caomhin

Reviewed By: jdoerfert

Subscribers: Hahnfeld, guansong, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61949

llvm-svn: 360804
2019-05-15 20:18:21 +00:00
Gheorghe-Teodor Bercea 946957189d [OpenMP][Clang][BugFix] Split declares and math functions inclusion.
Summary: This patches fixes an issue in which the __clang_cuda_cmath.h header is being included even when cmath or math.h headers are not included.

Reviewers: jdoerfert, ABataev, hfinkel, caomhin, tra

Reviewed By: tra

Subscribers: tra, mgorny, guansong, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61765

llvm-svn: 360626
2019-05-13 22:11:44 +00:00
Gheorghe-Teodor Bercea e62c693c8e [OpenMP][Clang] Support for target math functions
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

Reviewers: hfinkel, caomhin, ABataev, tra

Reviewed By: hfinkel, ABataev, tra

Subscribers: JDevlieghere, mgorny, guansong, cfe-commits, jdoerfert

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61399

llvm-svn: 360265
2019-05-08 15:52:33 +00:00
Jonas Devlieghere fe608c938c Revert "[OpenMP][Clang] Support for target math functions"
This commit appears to be breaking stage-2 builds on GreenDragon. The
OpenMP wrappers for cmath and math.h are copied into the root of the
resource directory and cause a cyclic dependency in module 'Darwin':
Darwin -> std -> Darwin. This blows up when CMake is testing for modules
support and breaks all stage 2 module builds, including the ThinLTO bot
and all LLDB bots.

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

llvm-svn: 360192
2019-05-07 21:08:15 +00:00
Gheorghe-Teodor Bercea 1e28a668bc [OpenMP][Clang] Support for target math functions
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

Reviewers: hfinkel, caomhin, ABataev, tra

Reviewed By: hfinkel, ABataev, tra

Subscribers: mgorny, guansong, cfe-commits, jdoerfert

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61399

llvm-svn: 360063
2019-05-06 18:19:15 +00:00
Evgeny Mankov 88aa3d7237 [CUDA][Windows] Restrict long double device functions declarations to Windows
As agreed in D60220, make long double declarations unobservable on non-windows platforms.

[Testing]
{Windows 10, Ubuntu 16.04.5}/{Visual C++ 2017 15.9.11 & 2019 16.0.1, gcc+ 5.4.0}/CUDA {8.0, 9.0, 9.1, 9.2, 10.0, 10.1}

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D60818

llvm-svn: 358654
2019-04-18 10:08:55 +00:00
Chandler Carruth 4cf5743b77 Move the builtin headers to use the new license file header.
Summary:
These all had somewhat custom file headers with different text from the
ones I searched for previously, and so I missed them. Thanks to Hal and
Kristina and others who prompted me to fix this, and sorry it took so
long.

Reviewers: hfinkel

Subscribers: mcrosier, javed.absar, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D60406

llvm-svn: 357941
2019-04-08 20:51:30 +00:00
Evgeny Mankov 66a8b07cd9 [CUDA][Windows] Last fix for the clang Bug 38811 "Clang fails to compile with CUDA-9.x on Windows" (https://bugs.llvm.org/show_bug.cgi?id=38811).
[IMPORTANT]
With that last fix, CUDA has just started being compiling by clang on Windows after nearly a year and two clang’s major releases (7 and 8).
As long as the last LLVM release, in which clang was compiling CUDA on Windows successfully, was 6.0.1, this fix and two previous have to be included into upcoming 7.1.0 and 8.0.1 releases.

[How to repro]
clang++.exe -x cuda "c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\0_Simple\simplePrintf\simplePrintf.cu" -I"c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\common\inc" --cuda-gpu-arch=sm_50 --cuda-path="C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0" -L"c:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\lib\x64" -lcudart.lib  -v

[Output]
In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.9\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:390:11: error: no matching function for call to '__isinfl'
  return (__isinfl(a) != 0);
          ^~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2662:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __isinfl(long double a))
             ^
In file included from <built-in>:1:
In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.9\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:438:11: error: no matching function for call to '__isnanl'
  return (__isnanl(a) != 0);
          ^~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2672:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __isnanl(long double a))
             ^
In file included from <built-in>:1:
In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.9\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:327:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:486:11: error: no matching function for call to '__finitel'
  return (__finitel(a) != 0);
          ^~~~~~~~~
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0/include\crt/math_functions.hpp:2652:14: note: candidate function not viable: call to __host__ function from __device__ function
__func__(int __finitel(long double a))
             ^
3 errors generated when compiling for sm_50.

[Solution]
Add missing long double device functions' declarations. Provide only declarations to prevent any use of long double on the device side, because CUDA does not support long double on the device side.

[Testing]
{Windows 10, Ubuntu 16.04.5}/{Visual C++ 2017 15.9.9, gcc+ 5.4.0}/CUDA {8.0, 9.0, 9.1, 9.2, 10.0, 10.1}

Reviewed by: Artem Belevich

Differential Revision: http://reviews.llvm.org/D60220

llvm-svn: 357779
2019-04-05 16:51:10 +00:00
Justin Lebar 370c766e40 [CUDA] Remove implementations of nexttoward.
Summary:
__builtin_nexttoward lowers to a libcall, e.g. nexttowardf(), that CUDA
does not have.

Rather than try to implement it, we simply remove these functions --
nvcc doesn't support them either, and nextafter, which does work, does
essentially the same thing on GPUs, because GPUs don't have long double.

Reviewers: tra

Subscribers: cfe-commits, sanjoy

Differential Revision: https://reviews.llvm.org/D40152

llvm-svn: 318494
2017-11-17 01:15:43 +00:00
Justin Lebar 11d5116904 [CUDA] Don't define functions that the CUDA headers themselves define on Windows.
Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D28324

llvm-svn: 291137
2017-01-05 16:53:55 +00:00
Justin Lebar 50fe985349 [CUDA] Wrapper header changes necessary to support MacOS.
Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D26780

llvm-svn: 287288
2016-11-18 00:41:35 +00:00
Justin Lebar e9eb792a0f [CUDA] Declare our __device__ math functions in the same inline namespace as our standard library.
Summary:
Currently we declare our inline __device__ math functions in namespace
std.  But libstdc++ and libc++ declare these functions in an inline
namespace inside namespace std.  We need to match this because, in a
later patch, we want to get e.g. <complex> to use our device overloads,
and it only will if those overloads are in the right inline namespace.

Reviewers: tra

Subscribers: cfe-commits, jhen

Differential Revision: https://reviews.llvm.org/D24977

llvm-svn: 283678
2016-10-08 22:16:03 +00:00
Justin Lebar e3612a039f [CUDA] Make __clang_cuda_cmath.h compatible with libc++.
Summary:
We need to add a bunch more "using"s, which weren't necessary with
libstdc++.

Once this is in I can check in a test to the test-suite.

Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D24588

llvm-svn: 281544
2016-09-14 21:50:14 +00:00
Justin Lebar cb20a09f54 [CUDA] Improve handling of math functions.
Summary:
A bunch of related changes here to our CUDA math headers.

- The second arg to nexttoward is a double (well, technically, long
  double, but we don't have that), not a float.

- Add a forward-declare of llround(float), which is defined in the CUDA
  headers.  We need this for the same reason we need most of the other
  forward-declares: To prevent a constexpr function in our standard
  library from becoming host+device.

- Add nexttowardf implementation.

- Pull "foobarf" functions defined by the CUDA headers in the global
  namespace into namespace std.  This lets you do e.g. std::sinf.

- Add overloads for math functions accepting integer types.  This lets
  you do e.g. std::sin(0) without having an ambiguity between the
  overload that takes a float and the one that takes a double.

With these changes, we pass testcases derived from libc++ for cmath and
math.h.  We can check these testcases in to the test-suite once support
for CUDA lands there.

Reviewers: tra

Subscribers: cfe-commits

Differential Revision: https://reviews.llvm.org/D23627

llvm-svn: 279140
2016-08-18 20:43:13 +00:00
Justin Lebar 25c36fd61b [CUDA] Tweak math forward declares so we're compatible with libstdc++4.9.
Summary:
See comments in patch; we were assuming that some stdlib math functions
would be defined in namespace std, when in fact the spec says they
should be defined in the global namespace.  libstdc++4.9 became more
conforming and broke us.

This new implementation seems to cover the known knowns.

Reviewers: rsmith

Subscribers: cfe-commits, tra

Differential Revision: http://reviews.llvm.org/D18882

llvm-svn: 265751
2016-04-07 23:55:53 +00:00
Justin Lebar 50e5f184d8 [CUDA] Add missing #undef __DEVICE__ to CUDA shim header.
llvm-svn: 264742
2016-03-29 16:24:23 +00:00
Artem Belevich 7f522b7876 Added missing '__'.
llvm-svn: 260719
2016-02-12 20:26:43 +00:00
Eric Christopher 39a84d0b9b Update functions in clang supplied headers to use the compiler reserved
namespace for arguments.

llvm-svn: 260647
2016-02-12 02:22:53 +00:00
Artem Belevich 2aad2b3500 [CUDA] Bug 26497 : Remove wrappers for variants provided by CUDA headers.
... and pull global-scope ones into std namespace with using-declaration.

Differential Revision: http://reviews.llvm.org/D16932

llvm-svn: 259944
2016-02-05 22:54:05 +00:00
Artem Belevich c5f41a34e5 [CUDA] Implemented device-side support functions in <cmath>.
CUDA expects math functions in std:: namespace to work on device side.
In order to make it work with clang without allowing device-side code
generation for functions w/o appropriate target attributes, this patch
provides device-side implementations for <cmath> functions. Most of
them call global-scope math functions provided by CUDA headers. In few
cases we use clang builtins.

Tested out-of tree by compiling and running thrust's unit_tests.
https://github.com/thrust/thrust/tree/master/testing

Differential Revision: http://reviews.llvm.org/D16593

llvm-svn: 258880
2016-01-26 23:37:29 +00:00