diff options
author | Dimitry Andric <dim@FreeBSD.org> | 2020-07-26 19:36:28 +0000 |
---|---|---|
committer | Dimitry Andric <dim@FreeBSD.org> | 2020-07-26 19:36:28 +0000 |
commit | cfca06d7963fa0909f90483b42a6d7d194d01e08 (patch) | |
tree | 209fb2a2d68f8f277793fc8df46c753d31bc853b /clang/lib/Headers/openmp_wrappers | |
parent | 706b4fc47bbc608932d3b491ae19a3b9cde9497b (diff) |
Notes
Diffstat (limited to 'clang/lib/Headers/openmp_wrappers')
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h (renamed from clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h) | 25 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h | 35 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/cmath | 69 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/complex | 25 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/complex.h | 25 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/math.h | 46 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/new | 70 | ||||
-rw-r--r-- | clang/lib/Headers/openmp_wrappers/time.h | 32 |
8 files changed, 273 insertions, 54 deletions
diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h index a422c98bf97d..406c9748e286 100644 --- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h +++ b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h @@ -1,4 +1,4 @@ -/*===---- __clang_openmp_math_declares.h - OpenMP math declares ------------=== +/*===- __clang_openmp_device_functions.h - OpenMP device function declares -=== * * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. * See https://llvm.org/LICENSE.txt for license information. @@ -7,27 +7,36 @@ *===-----------------------------------------------------------------------=== */ -#ifndef __CLANG_OPENMP_MATH_DECLARES_H__ -#define __CLANG_OPENMP_MATH_DECLARES_H__ +#ifndef __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ +#define __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ #ifndef _OPENMP #error "This file is for OpenMP compilation only." #endif -#if defined(__NVPTX__) && defined(_OPENMP) +#pragma omp begin declare variant match( \ + device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) -#define __CUDA__ - -#if defined(__cplusplus) - #include <__clang_cuda_math_forward_declares.h> +#ifdef __cplusplus +extern "C" { #endif +#define __CUDA__ +#define __OPENMP_NVPTX__ + /// Include declarations for libdevice functions. #include <__clang_cuda_libdevice_declares.h> + /// Provide definitions for these functions. #include <__clang_cuda_device_functions.h> +#undef __OPENMP_NVPTX__ #undef __CUDA__ +#ifdef __cplusplus +} // extern "C" #endif + +#pragma omp end declare variant + #endif diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h deleted file mode 100644 index 5d7ce9a965d3..000000000000 --- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h +++ /dev/null @@ -1,35 +0,0 @@ -/*===---- __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_cmath.h> -#endif - -#undef __CUDA__ - -/// Magic macro for stopping the math.h/cmath host header from being included. -#define __CLANG_NO_HOST_MATH__ - -#endif - diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath index a5183a1d8d1b..bd6011eb6f6d 100644 --- a/clang/lib/Headers/openmp_wrappers/cmath +++ b/clang/lib/Headers/openmp_wrappers/cmath @@ -1,4 +1,4 @@ -/*===-------------- cmath - Alternative cmath header -----------------------=== +/*===-- __clang_openmp_device_functions.h - OpenMP math declares ------ c++ -=== * * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. * See https://llvm.org/LICENSE.txt for license information. @@ -7,10 +7,69 @@ *===-----------------------------------------------------------------------=== */ -#include <__clang_openmp_math.h> +#ifndef __CLANG_OPENMP_CMATH_H__ +#define __CLANG_OPENMP_CMATH_H__ + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." +#endif -#ifndef __CLANG_NO_HOST_MATH__ #include_next <cmath> -#else -#undef __CLANG_NO_HOST_MATH__ + +// Make sure we include our math.h overlay, it probably happend already but we +// need to be sure. +#include <math.h> + +// We (might) need cstdlib because __clang_cuda_cmath.h below declares `abs` +// which might live in cstdlib. +#include <cstdlib> + +#pragma omp begin declare variant match( \ + device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + +#define __CUDA__ +#define __OPENMP_NVPTX__ +#include <__clang_cuda_cmath.h> +#undef __OPENMP_NVPTX__ +#undef __CUDA__ + +// Overloads not provided by the CUDA wrappers but by the CUDA system headers. +// Since we do not include the latter we define them ourselves. +#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow)) + +__DEVICE__ float acosh(float __x) { return ::acoshf(__x); } +__DEVICE__ float asinh(float __x) { return ::asinhf(__x); } +__DEVICE__ float atanh(float __x) { return ::atanhf(__x); } +__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); } +__DEVICE__ float erf(float __x) { return ::erff(__x); } +__DEVICE__ float erfc(float __x) { return ::erfcf(__x); } +__DEVICE__ float exp2(float __x) { return ::exp2f(__x); } +__DEVICE__ float expm1(float __x) { return ::expm1f(__x); } +__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); } +__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); } +__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); } +__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); } +__DEVICE__ long long int llrint(float __x) { return ::llrintf(__x); } +__DEVICE__ long long int llround(float __x) { return ::llroundf(__x); } +__DEVICE__ float log1p(float __x) { return ::log1pf(__x); } +__DEVICE__ float log2(float __x) { return ::log2f(__x); } +__DEVICE__ float logb(float __x) { return ::logbf(__x); } +__DEVICE__ long int lrint(float __x) { return ::lrintf(__x); } +__DEVICE__ long int lround(float __x) { return ::lroundf(__x); } +__DEVICE__ float nextafter(float __x, float __y) { + return ::nextafterf(__x, __y); +} +__DEVICE__ float remainder(float __x, float __y) { + return ::remainderf(__x, __y); +} +__DEVICE__ float scalbln(float __x, long int __y) { + return ::scalblnf(__x, __y); +} +__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); } +__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); } + +#undef __DEVICE__ + +#pragma omp end declare variant + #endif diff --git a/clang/lib/Headers/openmp_wrappers/complex b/clang/lib/Headers/openmp_wrappers/complex new file mode 100644 index 000000000000..1ed0b14879ef --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/complex @@ -0,0 +1,25 @@ +/*===-- complex --- OpenMP complex wrapper for target regions --------- c++ -=== + * + * 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_OPENMP_COMPLEX__ +#define __CLANG_OPENMP_COMPLEX__ + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." +#endif + +// We require std::math functions in the complex builtins below. +#include <cmath> + +#define __CUDA__ +#include <__clang_cuda_complex_builtins.h> +#endif + +// Grab the host header too. +#include_next <complex> diff --git a/clang/lib/Headers/openmp_wrappers/complex.h b/clang/lib/Headers/openmp_wrappers/complex.h new file mode 100644 index 000000000000..829c7a785725 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/complex.h @@ -0,0 +1,25 @@ +/*===-- complex --- OpenMP complex wrapper for target regions --------- c++ -=== + * + * 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_OPENMP_COMPLEX_H__ +#define __CLANG_OPENMP_COMPLEX_H__ + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." +#endif + +// We require math functions in the complex builtins below. +#include <math.h> + +#define __CUDA__ +#include <__clang_cuda_complex_builtins.h> +#endif + +// Grab the host header too. +#include_next <complex.h> diff --git a/clang/lib/Headers/openmp_wrappers/math.h b/clang/lib/Headers/openmp_wrappers/math.h index d2786ecb2424..c64af8b13ece 100644 --- a/clang/lib/Headers/openmp_wrappers/math.h +++ b/clang/lib/Headers/openmp_wrappers/math.h @@ -1,4 +1,4 @@ -/*===------------- math.h - Alternative math.h header ----------------------=== +/*===---- openmp_wrapper/math.h -------- OpenMP math.h intercept ------ c++ -=== * * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. * See https://llvm.org/LICENSE.txt for license information. @@ -7,11 +7,45 @@ *===-----------------------------------------------------------------------=== */ -#include <__clang_openmp_math.h> +// If we are in C++ mode and include <math.h> (not <cmath>) first, we still need +// to make sure <cmath> is read first. The problem otherwise is that we haven't +// seen the declarations of the math.h functions when the system math.h includes +// our cmath overlay. However, our cmath overlay, or better the underlying +// overlay, e.g. CUDA, uses the math.h functions. Since we haven't declared them +// yet we get errors. CUDA avoids this by eagerly declaring all math functions +// (in the __device__ space) but we cannot do this. Instead we break the +// dependence by forcing cmath to go first. While our cmath will in turn include +// this file, the cmath guards will prevent recursion. +#ifdef __cplusplus +#include <cmath> +#endif -#ifndef __CLANG_NO_HOST_MATH__ -#include_next <math.h> -#else -#undef __CLANG_NO_HOST_MATH__ +#ifndef __CLANG_OPENMP_MATH_H__ +#define __CLANG_OPENMP_MATH_H__ + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." #endif +#include_next <math.h> + +// We need limits.h for __clang_cuda_math.h below and because it should not hurt +// we include it eagerly here. +#include <limits.h> + +// We need stdlib.h because (for now) __clang_cuda_math.h below declares `abs` +// which should live in stdlib.h. +#include <stdlib.h> + +#pragma omp begin declare variant match( \ + device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + +#define __CUDA__ +#define __OPENMP_NVPTX__ +#include <__clang_cuda_math.h> +#undef __OPENMP_NVPTX__ +#undef __CUDA__ + +#pragma omp end declare variant + +#endif diff --git a/clang/lib/Headers/openmp_wrappers/new b/clang/lib/Headers/openmp_wrappers/new new file mode 100644 index 000000000000..1387d925b126 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/new @@ -0,0 +1,70 @@ +//===--------- new - OPENMP wrapper for <new> ------------------------------=== +// +// 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_OPENMP_WRAPPERS_NEW +#define __CLANG_OPENMP_WRAPPERS_NEW + +#include_next <new> + +#if defined(__NVPTX__) && defined(_OPENMP) + +#include <cstdlib> + +#pragma push_macro("OPENMP_NOEXCEPT") +#if __cplusplus >= 201103L +#define OPENMP_NOEXCEPT noexcept +#else +#define OPENMP_NOEXCEPT +#endif + +// Device overrides for non-placement new and delete. +inline void *operator new(__SIZE_TYPE__ size) { + if (size == 0) + size = 1; + return ::malloc(size); +} +inline void *operator new(__SIZE_TYPE__ size, + const std::nothrow_t &) OPENMP_NOEXCEPT { + return ::operator new(size); +} + +inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); } +inline void *operator new[](__SIZE_TYPE__ size, const std::nothrow_t &) { + return ::operator new(size); +} + +inline void operator delete(void *ptr)OPENMP_NOEXCEPT { + if (ptr) + ::free(ptr); +} +inline void operator delete(void *ptr, const std::nothrow_t &)OPENMP_NOEXCEPT { + ::operator delete(ptr); +} + +inline void operator delete[](void *ptr) OPENMP_NOEXCEPT { + ::operator delete(ptr); +} +inline void operator delete[](void *ptr, + const std::nothrow_t &) OPENMP_NOEXCEPT { + ::operator delete(ptr); +} + +// Sized delete, C++14 only. +#if __cplusplus >= 201402L +inline void operator delete(void *ptr, __SIZE_TYPE__ size)OPENMP_NOEXCEPT { + ::operator delete(ptr); +} +inline void operator delete[](void *ptr, __SIZE_TYPE__ size) OPENMP_NOEXCEPT { + ::operator delete(ptr); +} +#endif + +#pragma pop_macro("OPENMP_NOEXCEPT") +#endif + +#endif // include guard diff --git a/clang/lib/Headers/openmp_wrappers/time.h b/clang/lib/Headers/openmp_wrappers/time.h new file mode 100644 index 000000000000..c760dd1ed963 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/time.h @@ -0,0 +1,32 @@ +/*===---- time.h - OpenMP time header wrapper ------------------------ c ---=== + * + * 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_OPENMP_TIME_H__ +#define __CLANG_OPENMP_TIME_H__ + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." +#endif + +#if defined(__cplusplus) +#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow)) +#else +#define __DEVICE__ static __attribute__((always_inline, nothrow)) +#endif + +#include_next <time.h> + +#pragma omp begin declare variant match( \ + device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + +__DEVICE__ clock_t clock() { return __nvvm_read_ptx_sreg_clock(); } + +#pragma omp end declare variant + +#endif |