[OpenMP][Clang][BugFix] Split declares and math functions inclusion.

[gtbercea]

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.

[gtbercea]

• Remove define.

[jdoerfert]

This always includes the declare file but not the define file, correct?

Could we have 4 tests that are compiled in target mode with:

// with and without math.h/cmath (clang/clang++)
#include <math.h>

long abs(long __i) { return (__i < 0 ? -i : i); }

[gtbercea]

• Fix cstdlib issue.

[gtbercea]

• Move back functions.

[jdoerfert]

What about abs tests?

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

18

Why do we need the stdlib includes again?

test/Headers/Inputs/include/cstdlib

3

Where is this used? Are there tests missing?

[gtbercea]

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

18

They are both prone to abs inclusion.

We need them here to control the order in which they are included relative to the forward_declares header.

test/Headers/Inputs/include/cstdlib

3

I'll remove it.

[gtbercea]

• Clean test header.

[gtbercea]

• Mock cstdlib header's abs functions.

[jdoerfert]

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

18

I thought the "not defining abs" in __clang_cuda_math_forward_declares.h was the solution?

[tra]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

I'm not quite sure what's the idea here. It may be worth adding a comment.

It could also be expressed somewhat simpler:

#if !(defined(_OPENMP) && defined(__cplusplus))
...
#endif

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

11

You may want to add include guards.

I'd also make inclusion of the file in non-openmp compilation an error, if it makes sense for OpenMP. It does for CUDA.

[gtbercea]

• Add tests. Exclude additional abs definition.

[gtbercea]

• Simplify conditions. Add guards.

[gtbercea]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

When these two functions definitions are here or in the clang_cuda_cmath.h header then I get the following error (adapted for the clang_cuda_cmath.h case):

/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:166:3: error: declaration conflicts with target of using declaration already in scope
  abs(long __i) { return __builtin_labs(__i); }
  ^
/autofs/home/gbercea/patch-compiler/obj-release/lib/clang/9.0.0/include/__clang_cuda_cmath.h:40:17: note: target of using declaration
__DEVICE__ long abs(long __n) { return ::labs(__n); }
                ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:122:11: note: using declaration
  using ::abs;
          ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:174:3: error: declaration conflicts with target of using declaration already in scope
  abs(long long __x) { return __builtin_llabs (__x); }
  ^
/autofs/home/gbercea/patch-compiler/obj-release/lib/clang/9.0.0/include/__clang_cuda_cmath.h:39:22: note: target of using declaration
__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
                     ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:122:11: note: using declaration
  using ::abs;

[jdoerfert]

Last issue I have (in addition to the check @tra suggested) is the order in which we include math.h and cstdlib. can you flip it in one of the test cases?

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

Long story short, we currently cannot use the overload trick through __device__ and therefore *replace* (not augment) host math headers with the cuda versions which unfortunately mix std math functions with other functions that we don't want/need.

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

11

That is something we should be able to do, error out if _OPENMP is not defined I mean.

[gtbercea]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

This doesn't seem to be happening in the CUDA case. My suspicion is it's because of the device attribute.

[gtbercea]

• Error if not in OpenMP.

[tra]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

It looks like until OpenMP supports some sort of target-based overloading this will not play nicely with libstdc++.
Did you, by any chance, check if the header works with libc++ ? I wonder if we may encounter more conflicts like these.

[tra]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

Correct. __device__ functions overload whatever (implicitly)__host__ functions declared by the standard library, so they coexist w/o problems. Usually. host/device implementation nuances are still observable.

[gtbercea]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

Just did a few quick tests with libstdc++ and it was all good.

[tra]

lib/Headers/__clang_cuda_math_forward_declares.h

30–37

How about libc++? The idea is to make sure the change works with both libraries.

[jdoerfert]

Two small changes and then it is fine with me. @tra ?

1. we need to use ifdef to not define clock
2. Can you switch the include order in test/Headers/nvptx_device_math_functions.cpp?

P.S. I'm currently at the OpenMP standard meeting to get the OpenMP variants fixed.
Once done, we should prioritize the implementation.
Excluding non-math functions in the cuda headers is not perfect...

[tra]

In D61765#1499957, @jdoerfert wrote:

Two small changes and then it is fine with me. @tra ?

LGTM in general. I would still like to confirm that the changes work with libc++.

[gtbercea]

In D61765#1500233, @tra wrote:

In D61765#1499957, @jdoerfert wrote:

Two small changes and then it is fine with me. @tra ?

LGTM in general. I would still like to confirm that the changes work with libc++.

As soon as libc++ the limits header included in

__clang_cuda_cmath.h:15
``` is not found:

__clang_cuda_cmath.h:15:10: fatal error: 'limits' file not found
#include <limits>

Not even CUDA works actually so I'm not sure what the best answer to this problem is.

[gtbercea]

• Exclude clock functions. Reverse inclusion order.

[tra]

In D61765#1500309, @gtbercea wrote:

As soon as libc++ the limits header included in

__clang_cuda_cmath.h:15
``` is not found:

__clang_cuda_cmath.h:15:10: fatal error: 'limits' file not found
#include <limits>

Not even CUDA works actually so I'm not sure what the best answer to this problem is.

Could you give me more details on how you've got this error?

If this change breaks CUDA compilation with libc++, that's going to be a problem. Currently CUDA and clang's headers we ship do work with both libc++ and few versions of libstdc++:
E.g: http://lab.llvm.org:8011/builders/clang-cuda-build/builds/33364/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio

[gtbercea]

In D61765#1500446, @tra wrote:

In D61765#1500309, @gtbercea wrote:

As soon as libc++ the limits header included in

__clang_cuda_cmath.h:15
``` is not found:

__clang_cuda_cmath.h:15:10: fatal error: 'limits' file not found
#include <limits>

Not even CUDA works actually so I'm not sure what the best answer to this problem is.

Could you give me more details on how you've got this error?

If this change breaks CUDA compilation with libc++, that's going to be a problem. Currently CUDA and clang's headers we ship do work with both libc++ and few versions of libstdc++:
E.g: http://lab.llvm.org:8011/builders/clang-cuda-build/builds/33364/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio

It's an error on my side, I don't have libc++ installed so trying to use it will come up with header not found errors.

[gtbercea]

In D61765#1500446, @tra wrote:

In D61765#1500309, @gtbercea wrote:

As soon as libc++ the limits header included in

__clang_cuda_cmath.h:15
``` is not found:

__clang_cuda_cmath.h:15:10: fatal error: 'limits' file not found
#include <limits>

Not even CUDA works actually so I'm not sure what the best answer to this problem is.

Could you give me more details on how you've got this error?

If this change breaks CUDA compilation with libc++, that's going to be a problem. Currently CUDA and clang's headers we ship do work with both libc++ and few versions of libstdc++:
E.g: http://lab.llvm.org:8011/builders/clang-cuda-build/builds/33364/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio

This won't affect CUDA in any way, all we have added is OpenMP specific.

[tra]

In D61765#1500457, @gtbercea wrote:

This won't affect CUDA in any way, all we have added is OpenMP specific.

LGTM for CUDA. I'll leave the question of testing with libc++ to someone more familiar with OpenMP.

[jdoerfert]

In D61765#1500351, @gtbercea wrote:

• Exclude clock functions. Reverse inclusion order.

LGTM from my side. I don't have strong feelings about testing libc++ now, though it is probably a good idea to have such a testbed.
I agree this should not infer with CUDA (at least that is our intention).

P.S.
There will be a ticket to add the OpenMP variant support we need (similar to the __device__ in CUDA) will be written by Tom Scogland and me for a first vote already in this meeting.
That is what we will then need to implement.

lib/Driver/ToolChains/Clang.cpp

if (!Args.hasArg(options::OPT_nobuiltininc)) {
       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");
+    CmdArgs.push_back("__clang_openmp_math_declares.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.
 

lib/Headers/CMakeLists.txt

 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
+  openmp_wrappers/__clang_openmp_math_declares.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})

lib/Headers/__clang_cuda_cmath.h

 // 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
 
+#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); }
 __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); }

lib/Headers/__clang_cuda_math_forward_declares.h

 #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);
+#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);
 __DEVICE__ float acosh(float);
 __DEVICE__ double asin(double);
 __DEVICE__ float asin(float);
 __DEVICE__ double asinh(double);
 __DEVICE__ float asinh(float);
 __DEVICE__ double atan2(double, double);

lib/Headers/openmp_wrappers/__clang_openmp_math.h

 /// 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/__clang_openmp_math_declares.h

+/*===---- __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

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/Headers/Inputs/include/cstdlib

+#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); }
+}

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
+// 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>
 
 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
+// 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
   {
     // 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
+// 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>
 
 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
+// 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
   {
     // 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);
   }
 }