gtbercea updated this revision to Diff 199073.
gtbercea added a comment.
- Add tests. Exclude additional abs definition.
Repository:
rC Clang
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D61765/new/
https://reviews.llvm.org/D61765
Files:
lib/Driver/ToolChains/Clang.cpp
lib/Headers/CMakeLists.txt
lib/Headers/__clang_cuda_cmath.h
lib/Headers/__clang_cuda_math_forward_declares.h
lib/Headers/openmp_wrappers/__clang_openmp_math.h
lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h
lib/Headers/openmp_wrappers/cmath
lib/Headers/openmp_wrappers/math.h
test/Headers/Inputs/include/cstdlib
test/Headers/nvptx_device_cmath_functions.c
test/Headers/nvptx_device_cmath_functions.cpp
test/Headers/nvptx_device_math_functions.c
test/Headers/nvptx_device_math_functions.cpp
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 <cstdlib>
#include <math.h>
+#include <cstdlib>
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 <math.h>
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 <cmath>
+#include <cstdlib>
void test_sqrt(double a1) {
#pragma omp target
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 <cmath>
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: 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 <math.h>
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 <cmath>
Index: lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h
===================================================================
--- /dev/null
+++ lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h
@@ -0,0 +1,25 @@
+/*===---- __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
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#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
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/__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,18 @@
static __inline__ __attribute__((always_inline)) __attribute__((device))
#endif
-__DEVICE__ double abs(double);
-__DEVICE__ float abs(float);
-__DEVICE__ int abs(int);
+#ifndef _OPENMP
+__DEVICE__ long abs(long);
+__DEVICE__ long long abs(long long);
+#else
+#ifndef __cplusplus
__DEVICE__ long abs(long);
__DEVICE__ long long abs(long long);
+#endif
+#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/__clang_cuda_cmath.h
===================================================================
--- lib/Headers/__clang_cuda_cmath.h
+++ lib/Headers/__clang_cuda_cmath.h
@@ -36,8 +36,15 @@
#define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
#endif
+#ifndef _OPENMP
+__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
+__DEVICE__ long abs(long __n) { return ::labs(__n); }
+#else
+#ifndef __cplusplus
__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
__DEVICE__ long abs(long __n) { return ::labs(__n); }
+#endif
+#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/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/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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
