Page MenuHomePhabricator

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

Authored by gtbercea on May 9 2019, 3:33 PM.

Diff Detail

Repository
rC Clang

Event Timeline

gtbercea created this revision.May 9 2019, 3:33 PM
Herald added a project: Restricted Project. · View Herald TranscriptMay 9 2019, 3:33 PM
gtbercea updated this revision to Diff 198929.May 9 2019, 3:48 PM
  • Remove define.

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 updated this revision to Diff 199041.May 10 2019, 11:10 AM
  • Fix cstdlib issue.
gtbercea updated this revision to Diff 199043.May 10 2019, 11:29 AM
  • Move back functions.

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 marked 2 inline comments as done.May 10 2019, 12:44 PM
gtbercea added inline comments.
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 updated this revision to Diff 199062.May 10 2019, 12:48 PM
  • Clean test header.
gtbercea updated this revision to Diff 199064.May 10 2019, 1:02 PM
  • Mock cstdlib header's abs functions.
gtbercea marked 2 inline comments as done.May 10 2019, 1:02 PM
jdoerfert added inline comments.May 10 2019, 1:02 PM
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 added a subscriber: tra.May 10 2019, 1:22 PM
tra added inline comments.
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 updated this revision to Diff 199073.May 10 2019, 1:28 PM
  • Add tests. Exclude additional abs definition.
gtbercea updated this revision to Diff 199075.May 10 2019, 1:37 PM
  • Simplify conditions. Add guards.
gtbercea marked 2 inline comments as done.May 10 2019, 1:39 PM
gtbercea marked an inline comment as done.May 10 2019, 1:41 PM
gtbercea added inline comments.
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;

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 marked an inline comment as done.May 10 2019, 1:49 PM
gtbercea added inline comments.
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 updated this revision to Diff 199080.May 10 2019, 1:52 PM
  • Error if not in OpenMP.
tra added inline comments.May 10 2019, 1:53 PM
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 added inline comments.May 10 2019, 1:57 PM
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 marked an inline comment as done.May 10 2019, 2:27 PM
gtbercea added inline comments.
lib/Headers/__clang_cuda_math_forward_declares.h
30–37

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

tra added inline comments.May 10 2019, 2:36 PM
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.

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 added a comment.May 13 2019, 10:29 AM

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++.

In D61765#1500233, @tra 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 updated this revision to Diff 199304.May 13 2019, 12:13 PM
  • Exclude clock functions. Reverse inclusion order.
tra added a comment.May 13 2019, 1:06 PM

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

In D61765#1500446, @tra 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.

In D61765#1500446, @tra 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 accepted this revision.May 13 2019, 2:31 PM

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.

This revision is now accepted and ready to land.May 13 2019, 2:31 PM
  • 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.

This revision was automatically updated to reflect the committed changes.