Index: lib/Driver/ToolChains/Clang.cpp =================================================================== --- lib/Driver/ToolChains/Clang.cpp +++ lib/Driver/ToolChains/Clang.cpp @@ -1166,7 +1166,7 @@ } CmdArgs.push_back("-include"); - CmdArgs.push_back("__clang_openmp_math.h"); + CmdArgs.push_back("__clang_openmp_math_declares.h"); } // Add -i* options, and automatically translate to Index: lib/Headers/CMakeLists.txt =================================================================== --- lib/Headers/CMakeLists.txt +++ lib/Headers/CMakeLists.txt @@ -132,6 +132,7 @@ openmp_wrappers/math.h openmp_wrappers/cmath openmp_wrappers/__clang_openmp_math.h + openmp_wrappers/__clang_openmp_math_declares.h ) set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include) Index: lib/Headers/__clang_cuda_cmath.h =================================================================== --- lib/Headers/__clang_cuda_cmath.h +++ lib/Headers/__clang_cuda_cmath.h @@ -36,8 +36,10 @@ #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) #endif +#if !(defined(_OPENMP) && defined(__cplusplus)) __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } __DEVICE__ long abs(long __n) { return ::labs(__n); } +#endif __DEVICE__ float abs(float __x) { return ::fabsf(__x); } __DEVICE__ double abs(double __x) { return ::fabs(__x); } __DEVICE__ float acos(float __x) { return ::acosf(__x); } Index: lib/Headers/__clang_cuda_math_forward_declares.h =================================================================== --- lib/Headers/__clang_cuda_math_forward_declares.h +++ lib/Headers/__clang_cuda_math_forward_declares.h @@ -27,11 +27,13 @@ static __inline__ __attribute__((always_inline)) __attribute__((device)) #endif -__DEVICE__ double abs(double); -__DEVICE__ float abs(float); -__DEVICE__ int abs(int); +#if !(defined(_OPENMP) && defined(__cplusplus)) __DEVICE__ long abs(long); __DEVICE__ long long abs(long long); +#endif +__DEVICE__ int abs(int); +__DEVICE__ double abs(double); +__DEVICE__ float abs(float); __DEVICE__ double acos(double); __DEVICE__ float acos(float); __DEVICE__ double acosh(double); Index: lib/Headers/openmp_wrappers/__clang_openmp_math.h =================================================================== --- lib/Headers/openmp_wrappers/__clang_openmp_math.h +++ lib/Headers/openmp_wrappers/__clang_openmp_math.h @@ -23,15 +23,6 @@ #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 Index: lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h =================================================================== --- /dev/null +++ lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h @@ -0,0 +1,29 @@ +/*===---- __clang_openmp_math_declares.h - OpenMP math declares ------------=== + * + * 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_MATH_DECLARES_H__ +#define __CLANG_OPENMP_MATH_DECLARES_H__ + +#if defined(__NVPTX__) && defined(_OPENMP) + +#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> + +#undef __CUDA__ + +#endif +#endif Index: lib/Headers/openmp_wrappers/cmath =================================================================== --- lib/Headers/openmp_wrappers/cmath +++ lib/Headers/openmp_wrappers/cmath @@ -7,9 +7,7 @@ *===-----------------------------------------------------------------------=== */ -#ifndef __cplusplus #include <__clang_openmp_math.h> -#endif #ifndef __CLANG_NO_HOST_MATH__ #include_next Index: lib/Headers/openmp_wrappers/math.h =================================================================== --- lib/Headers/openmp_wrappers/math.h +++ lib/Headers/openmp_wrappers/math.h @@ -7,9 +7,7 @@ *===-----------------------------------------------------------------------=== */ -#ifndef __cplusplus #include <__clang_openmp_math.h> -#endif #ifndef __CLANG_NO_HOST_MATH__ #include_next Index: test/Headers/Inputs/include/cstdlib =================================================================== --- /dev/null +++ test/Headers/Inputs/include/cstdlib @@ -0,0 +1,16 @@ +#pragma once + +extern int abs (int __x) __attribute__ ((__const__)) ; +extern long int labs (long int __x) __attribute__ ((__const__)) ; + +namespace std +{ + +using ::abs; + +inline long +abs(long __i) { return __builtin_labs(__i); } + +inline long long +abs(long long __x) { return __builtin_llabs (__x); } +} Index: test/Headers/nvptx_device_cmath_functions.c =================================================================== --- test/Headers/nvptx_device_cmath_functions.c +++ test/Headers/nvptx_device_cmath_functions.c @@ -4,7 +4,7 @@ // 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 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.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 Index: test/Headers/nvptx_device_cmath_functions.cpp =================================================================== --- test/Headers/nvptx_device_cmath_functions.cpp +++ test/Headers/nvptx_device_cmath_functions.cpp @@ -4,9 +4,10 @@ // 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 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.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 +#include void test_sqrt(double a1) { #pragma omp target Index: test/Headers/nvptx_device_math_functions.c =================================================================== --- test/Headers/nvptx_device_math_functions.c +++ test/Headers/nvptx_device_math_functions.c @@ -4,7 +4,7 @@ // 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 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.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 Index: test/Headers/nvptx_device_math_functions.cpp =================================================================== --- test/Headers/nvptx_device_math_functions.cpp +++ test/Headers/nvptx_device_math_functions.cpp @@ -4,9 +4,11 @@ // 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 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -internal-isystem %S/Inputs/include -include stdlib.h -include limits -include cstdlib -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 #include +#include void test_sqrt(double a1) { #pragma omp target