Offloading build fails

I am not sure what’s going on with the trunk clang at this moment.

$ cat a.cpp
#include
int main() {
}

$ clang++ -std=c++11 -fopenmp -nu/7.2.0/…/…/…/…/include/c++/7.2.0/bits/std_abs.h:70:3:
error: declaration conflicts with target of using declaration already in
scope
abs(double __x)
^
/p/scratch/cpcp0/pcp0151/opt/clang/201911040249/lib/clang/10.0.0/include/__clang_cuda_math_forward_declares.h:50:25:
note: target of using declaration
DEVICE const double abs(const double);

I think this is bug https://bugs.llvm.org/show_bug.cgi?id=42972. The fix may be adding some #ifdef logic to the cuda header, I’ve worked around it locally by reverting D62046

Can you submit a patch to upstream? In the meantime, I workaround by using libc++.

The ‘right’ fix is to stop cuda putting symbols in namespace std::, but that ship has sailed. Or possibly openmp variants as a future repair.

I’ve ‘raised a concern’ on the guilty commit (https://reviews.llvm.org/rC361066) in the hope that someone else fixes it. My next step would probably be to propose a reverting commit and put the author of that patch down as reviewer.

Cheers

Jon

Switching to libc++ did not help much:
$ cat example.cpp

#include
int main() {
#pragma omp target parallel for
for (int i=0;i<10;i++) {
}
}

$ $ clang++ -fopenmp -fopenmp-targets=nvptx64 -stdlib=libc++ example.cpp
nvlink error : Undefined reference to ‘_Z13checkSPMDModeP5ident’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z25GetLogicalThreadIdInBlockb’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z14GetOmpThreadIdib’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z21GetNumberOfOmpThreadsb’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z12GetOmpTeamIdv’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z19GetNumberOfOmpTeamsv’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z22setExecutionParameters13ExecutionMode11RuntimeMode’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z18GetThreadIdInBlockv’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z25GetNumberOfThreadsInBlockv’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z9GetLaneIdv’ in ‘/tmp/example-b3be69.cubin’
nvlink error : Undefined reference to ‘_Z9GetWarpIdv’ in ‘/tmp/example-b3be69.cubin’
clang-10: error: nvlink command failed with exit code 255 (use -v to see invocation)

Tested on POWER8

That’s actually interesting in itself. Plus it’s a problem with my change so I’m definitely on the hook to fix that one.

I moved some inline functions (the symbols you’re seeing) into a source file and that broke some builds. https://reviews.llvm.org/D69652. cc Alexey who reported the error earlier this evening.

Nvcc 10 on x64 works for me with that change, some versions of nvcc on some architectures don’t work. The workaround I’m suggesting is sed -i support.cu ‘s/INLINE/DEVICE/g’, but without a local repro it’s difficult to be sure what the effect is.

I’d be very interested in what version of the nvcc toolchain you’re using, and in whether replacing INLINE with DEVICE in deviceRTL/nvptx/src/support.cu fixes it. Alternatively, forceinline to inline in deviceRTL/nvptx/src/target_impl.h’s definition of INLINE may fix things.

Thanks,

Jon

I’m on POWER8, the nvcc version is below:

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2018 NVIDIA Corporation
Built on Sat_Aug_25_21:10:00_CDT_2018
Cuda compilation tools, release 10.0, V10.0.130

Jon, I can’t locate the file you mentioned in the latest trunk’s openmp project directory.

Thanks! That probably rules out nvcc version dependent behaviour.

If so, then to reproduce I only need to work out how to run the power8 toolchain. Google suggests nvcc is intended to work as a cross compiler so it may be a case of guessing cmake flags. I’m in Europe so won’t get to it tonight.

Clearest way to point to the file is https://reviews.llvm.org/rG764c8420e4b8fc11a9fa94d00f4ee617aa754cb2, otherwise openmp/libomptarget/deviceRTL/nvptx/src/support.cu

Thanks,

Jon

The "solution" here is not reverting the commits, that will just cause
other problems we need to work around. The solution is what was planned
from the very beginning but requires `#omp begin/end declare variant`
support:

  https://reviews.llvm.org/D60907#1484756
  http://lists.llvm.org/pipermail/openmp-dev/2019-July/002629.html