Page MenuHomePhabricator

[HIP] Add default header and include path
ClosedPublic

Authored by yaxunl on Jun 4 2020, 10:30 AM.

Details

Summary

To support std::complex and some other standard C/C++ functions in HIP device code,
they need to be forced to be __host__ __device__ functions by pragmas. This is done
by some clang standard C++ wrapper headers which are shared between cuda-clang and hip-Clang.

For these standard C++ wapper headers to work properly, specific include path order
has to be enforced:

  1. clang C++ wrapper include path
  2. standard C++ include path
  3. clang include path

Also, these C++ wrapper headers require device version of some standard C/C++ functions
must be declared before including them. This needs to be done by including a default
header which declares or defines these device functions. The default header is always
included before any other headers are included by users.

This patch adds the the default header and include path for HIP.

Corresponding HIP runtime change: https://github.com/ROCm-Developer-Tools/HIP/pull/2098

Diff Detail

Event Timeline

yaxunl created this revision.Jun 4 2020, 10:30 AM
yaxunl edited the summary of this revision. (Show Details)Jun 4 2020, 11:25 AM
tra accepted this revision.Jun 4 2020, 11:36 AM

Thank you for the patch. This will make my life a lot easier.

There are a few nits, but it's LGTM in general.

Do I understand it correctly that for detection of HIP installation all we need is to find the bitcode libraries?

clang/lib/Driver/ToolChains/ROCm.h
68

CUDA -> GPU?
Looks like there are number of other mentions of CUDA that should be fixed.

clang/lib/Headers/__clang_hip_math.h
27

Arguments for compiler-internal declarations and local symbols should all be prefixed with __.

932–936

The macros pushed here do not match the set of macros defined below.
E.g. __HIP_OVERLOAD vs __HIP_OVERLOAD1, __DEF_FLOAT_FUN vs. __DEF_FUN1, etc.

This revision is now accepted and ready to land.Jun 4 2020, 11:36 AM
yaxunl marked 6 inline comments as done.Jun 5 2020, 7:51 AM
yaxunl added inline comments.
clang/lib/Driver/ToolChains/ROCm.h
68

will fix when commit

clang/lib/Headers/__clang_hip_math.h
27

will fix when committing

932–936

will fix when commit

This revision was automatically updated to reflect the committed changes.
yaxunl marked 3 inline comments as done.
Herald added a project: Restricted Project. · View Herald TranscriptJun 5 2020, 10:04 AM
thakis added a subscriber: thakis.Jun 5 2020, 12:03 PM

This is still failing on Windows: http://45.33.8.238/win/16970/step_7.txt

Please take a look and revert for now if it takes a while to investigate.

Also still failing on mac: http://45.33.8.238/mac/14992/step_7.txt

I will try to fix it one more time. If it still fails I will revert it.

reverted.

The difficulty for fixing these failures is that they only show up on some systems I do not have access and cannot see the complete output.

thakis added a comment.Jun 5 2020, 6:18 PM

The recommit again broke tests on mac and windows:
http://45.33.8.238/win/17004/step_7.txt
http://45.33.8.238/mac/15014/step_7.txt

This is like the 8th distinct build break today (most not yours, of course). Let's revert and call it a weekend, and try again next week.

yaxunl added a comment.Jun 5 2020, 6:51 PM

reverted.

The difficulty for fixing these failures is that they only show up on some systems I do not have access and cannot see the complete output.

recommitted by 1fa43e0b34d9736f62c6c1b6c371a5e39cd1624d. so far so good.

thakis added a comment.Jun 5 2020, 7:06 PM

One more example failure: http://lab.llvm.org:8011/builders/clang-x64-windows-msvc/builds/16333

(Note I had re-reverted your reland, see the 2 comments above your comment.)

One more example failure: http://lab.llvm.org:8011/builders/clang-x64-windows-msvc/builds/16333

(Note I had re-reverted your reland, see the 2 comments above your comment.)

Sorry I missed that failure. I have fixed that issue and recommitted again. I will monitor it and revert if necessary. Thanks.

This broke a test clang/test/Tooling/clang-check-offload.cpp for a critical Linux distro at Facebook. With this change, the test adds a -include __clang_hip_runtime_wrapper argument. The wrapper includes some standard c++ headers, but our distro don't have those headers in the default include paths, thus causing a break.

I notice this behavior doesn't happen for CUDA tests, which also rely on a similar __clang_cuda_runtime_wrapper. I think what's causing the difference is the different handling of nogpuinc/nogpulib option. My knowledge on this area is limited, so correct me if I'm wrong. CUDA seems to respect nogpuinc and doesn't include its wrapper if the flag is provided: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/Cuda.cpp#L255. But based on this change, HIP does things differently: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/AMDGPU.cpp#L226.

If I modify RocmInstallationDetector::AddHIPIncludeArgs to also respect nogpuinc/nogpulib, the test will pass for us. Is it a mistake for HIP to always include the wrapper file? Could you provide a fix for this issue? Thanks!

This broke a test clang/test/Tooling/clang-check-offload.cpp for a critical Linux distro at Facebook. With this change, the test adds a -include __clang_hip_runtime_wrapper argument. The wrapper includes some standard c++ headers, but our distro don't have those headers in the default include paths, thus causing a break.

I notice this behavior doesn't happen for CUDA tests, which also rely on a similar __clang_cuda_runtime_wrapper. I think what's causing the difference is the different handling of nogpuinc/nogpulib option. My knowledge on this area is limited, so correct me if I'm wrong. CUDA seems to respect nogpuinc and doesn't include its wrapper if the flag is provided: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/Cuda.cpp#L255. But based on this change, HIP does things differently: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/AMDGPU.cpp#L226.

If I modify RocmInstallationDetector::AddHIPIncludeArgs to also respect nogpuinc/nogpulib, the test will pass for us. Is it a mistake for HIP to always include the wrapper file? Could you provide a fix for this issue? Thanks!

Thanks for investigating the issue. It makes sense to respect nogpuinc and nogpulib. fixed by 2580635bd2f3c0527353e4d7823326cd9f92ff7c

This broke a test clang/test/Tooling/clang-check-offload.cpp for a critical Linux distro at Facebook. With this change, the test adds a -include __clang_hip_runtime_wrapper argument. The wrapper includes some standard c++ headers, but our distro don't have those headers in the default include paths, thus causing a break.

I notice this behavior doesn't happen for CUDA tests, which also rely on a similar __clang_cuda_runtime_wrapper. I think what's causing the difference is the different handling of nogpuinc/nogpulib option. My knowledge on this area is limited, so correct me if I'm wrong. CUDA seems to respect nogpuinc and doesn't include its wrapper if the flag is provided: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/Cuda.cpp#L255. But based on this change, HIP does things differently: https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/AMDGPU.cpp#L226.

If I modify RocmInstallationDetector::AddHIPIncludeArgs to also respect nogpuinc/nogpulib, the test will pass for us. Is it a mistake for HIP to always include the wrapper file? Could you provide a fix for this issue? Thanks!

Thanks for investigating the issue. It makes sense to respect nogpuinc and nogpulib. fixed by 2580635bd2f3c0527353e4d7823326cd9f92ff7c

It works! Thanks for the quick fix.