clang frontend command failed - using std::math

[dosarudaniel]

Describe the bug

clang-13 failed at link time

dadosaru@pcphsft106:~/oneAdePT/debug/examples/Example9.1$ clang++ -o example9.1 CMakeFiles/example9.1.dir/electrons.dp.cpp.o -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice
clang-13: warning: Unknown CUDA version. cuda.h: CUDA_VERSION=11010. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
clang-13: /home/dadosaru/sycl_workspace_may/llvm/llvm/lib/Target/NVPTX/SYCL/LocalAccessorToSharedMemory.cpp:51: virtual bool {anonymous}::LocalAccessorToSharedMemory::runOnModule(llvm::Module&): Assertion `NvvmMetadata && "IR compiled to PTX must have nvvm.annotations"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.      Program arguments: /home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13 -cc1 -triple nvptx64-nvidia-cuda-sycldevice -fsycl-is-device -fdeclare-spirv-builtins -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-strict -sycl-std=2020 -S -disable-free -main-file-name electrons.dp.cpp.o -mrelocation-model static -mframe-pointer=all -fno-rounding-math -fno-verbose-asm -no-integrated-as -aux-target-cpu x86-64 -internal-isystem /home/dadosaru/sycl_workspace_may/llvm/build//bin/../include/sycl -internal-isystem /home/dadosaru/sycl_workspace_may/llvm/build//bin/../include -mlink-builtin-bitcode /home/dadosaru/sycl_workspace_may/llvm/build/lib/clang/13.0.0/../../clc/libspirv-nvptx64--nvidiacl.bc -mlink-builtin-bitcode /usr/local/cuda-11.1/nvvm/libdevice/libdevice.10.bc -target-feature +ptx71 -target-sdk-version=11.1 -target-cpu sm_50 -mllvm -treat-scalable-fixed-error-as-warning -debugger-tuning=gdb -resource-dir /home/dadosaru/sycl_workspace_may/llvm/build/lib/clang/13.0.0 -fno-dwarf-directory-asm -fdebug-compilation-dir=/home/dadosaru/oneAdePT/debug/examples/Example9.1 -ferror-limit 19 -fgnuc-version=4.2.1 -fcolor-diagnostics -o /tmp/electrons-e79761.s -x ir /tmp/electrons-a2abbe.bc
1.      Code generation
2.      Running pass 'localaccessortosharedmemory' on module '/tmp/electrons-8350fd.bc'.
 #0 0x0000561e4883becc llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x2087ecc)
 #1 0x0000561e48839b04 llvm::sys::RunSignalHandlers() (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x2085b04)
 #2 0x0000561e48839c73 SignalHandler(int) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x2085c73)
 #3 0x00007fb5d51d18a0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x128a0)
 #4 0x00007fb5d3e37f47 raise /build/glibc-2ORdQG/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0
 #5 0x00007fb5d3e398b1 abort /build/glibc-2ORdQG/glibc-2.27/stdlib/abort.c:81:0
 #6 0x00007fb5d3e2942a __assert_fail_base /build/glibc-2ORdQG/glibc-2.27/assert/assert.c:89:0
 #7 0x00007fb5d3e294a2 (/lib/x86_64-linux-gnu/libc.so.6+0x304a2)
 #8 0x0000561e4775be06 (anonymous namespace)::LocalAccessorToSharedMemory::runOnModule(llvm::Module&) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0xfa7e06)
 #9 0x0000561e48001901 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x184d901)
#10 0x0000561e48b2ee20 (anonymous namespace)::EmitAssemblyHelper::EmitAssembly(clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x237ae20)
#11 0x0000561e48b30dbe clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::DataLayout const&, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x237cdbe)
#12 0x0000561e497bda5f clang::CodeGenAction::ExecuteAction() (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x3009a5f)
#13 0x0000561e4916f2b9 clang::FrontendAction::Execute() (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x29bb2b9)
#14 0x0000561e4910a761 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x2956761)
#15 0x0000561e49238dc0 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0x2a84dc0)
#16 0x0000561e473f7ad4 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0xc43ad4)
#17 0x0000561e473f40f7 ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0xc400f7)
#18 0x0000561e4736ae1c main (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0xbb6e1c)
#19 0x00007fb5d3e1ab97 __libc_start_main /build/glibc-2ORdQG/glibc-2.27/csu/../csu/libc-start.c:344:0
#20 0x0000561e473f3b5a _start (/home/dadosaru/sycl_workspace_may/llvm/build/bin/clang-13+0xc3fb5a)
clang-13: error: unable to execute command: Aborted (core dumped)
clang-13: error: clang frontend command failed due to signal (use -v to see invocation)
clang version 13.0.0 (git@github.com:dosarudaniel/llvm.git 6a49170027fba61cea649ffdd299ed095b5e3d75)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/dadosaru/sycl_workspace_may/llvm/build//bin
clang-13: note: diagnostic msg: Error generating preprocessed source(s) - no preprocessable inputs.

To Reproduce
When I am trying to execute make example9.1 in this code I get this error:

[ 89%] Linking CXX executable example9.1
cd /home/dadosaru/oneAdePT/debug/examples/Example9.1 && /usr/local/bin/cmake -E cmake_link_script CMakeFiles/example9.1.dir/link.txt --verbose=1
/home/dadosaru/sycl_workspace_may/llvm/build//bin/clang++ -O3 -DNDEBUG CMakeFiles/example9.1.dir/example9.cpp.o CMakeFiles/example9.1.dir/example9.dp.cpp.o CMakeFiles/example9.1.dir/electrons.dp.cpp.o CMakeFiles/example9.1.dir/gammas.dp.cpp.o -o example9.1  -Wl,-rpath,/home/dadosaru/oneAdePT/debug/external/g4hepem/lib:/usr/local/lib:/home/dadosaru/local/lib /home/dadosaru/local/lib/libvecgeom.a /home/dadosaru/local/lib/libvecgeomcuda_static.a /home/dadosaru/local/lib/libvgdml.a ../../external/g4hepem/lib/libg4HepEmInit.so ../../external/g4hepem/lib/libg4HepEmRun.so -fsycl -fsycl-unnamed-lambda -fsycl-targets=nvptx64-nvidia-cuda-sycldevice -Xsycl-target-backend=nvptx64-nvidia-cuda-sycldevice --cuda-gpu-arch=sm_61 -std=c++20 -stdlib=libstdc++ -Wno-unknown-cuda-version -D_ONEADEPT_ -DNDEBUG -DROOT_NO_INT128 -DG4HepEm_CUDA_BUILD -I/home/dadosaru/oneAdePT/../local/include /home/dadosaru/local/lib/libvecgeom.a -ldl -lrt -lpthread /usr/local/lib/libxerces-c.so /home/dadosaru/local/lib/libG4Tree.so /home/dadosaru/local/lib/libG4FR.so /home/dadosaru/local/lib/libG4GMocren.so /home/dadosaru/local/lib/libG4visHepRep.so /home/dadosaru/local/lib/libG4RayTracer.so /home/dadosaru/local/lib/libG4VRML.so /home/dadosaru/local/lib/libG4vis_management.so /home/dadosaru/local/lib/libG4modeling.so /home/dadosaru/local/lib/libG4interfaces.so /home/dadosaru/local/lib/libG4persistency.so /home/dadosaru/local/lib/libG4error_propagation.so /home/dadosaru/local/lib/libG4readout.so /home/dadosaru/local/lib/libG4physicslists.so /home/dadosaru/local/lib/libG4tasking.so /home/dadosaru/local/lib/libG4run.so /home/dadosaru/local/lib/libG4event.so /home/dadosaru/local/lib/libG4tracking.so /home/dadosaru/local/lib/libG4parmodels.so /home/dadosaru/local/lib/libG4processes.so /home/dadosaru/local/lib/libG4analysis.so /usr/lib/x86_64-linux-gnu/libexpat.so /home/dadosaru/local/lib/libG4digits_hits.so /home/dadosaru/local/lib/libG4track.so /home/dadosaru/local/lib/libG4particles.so /home/dadosaru/local/lib/libG4geometry.so /home/dadosaru/local/lib/libG4materials.so /home/dadosaru/local/lib/libG4graphics_reps.so /home/dadosaru/local/lib/libG4intercoms.so /home/dadosaru/local/lib/libG4global.so /home/dadosaru/local/lib/libG4zlib.so /home/dadosaru/local/lib/libG4ptl.so.0.0.2 -pthread /home/dadosaru/local/lib/libG4clhep.so ../../external/g4hepem/lib/libg4HepEmData.so 
fatal error: error in backend: Cannot select: t53: f64 = fcos t52
  t52: f64 = fdiv t51, t37
    t51: f64 = fmul t49, ConstantFP:f64<-2.997925e+10>
      t49: f64 = fmul t46, t48
        t46: f64 = fmul t21, t45
          t21: f64 = select t20, t16, t18
            t20: i1 = setcc t16, t18, setolt:ch
              t16: f64,ch = CopyFromReg t0, Register:f64 %11
                t15: f64 = Register %11
              t18: f64,ch = CopyFromReg t0, Register:f64 %58
                t17: f64 = Register %58
            t16: f64,ch = CopyFromReg t0, Register:f64 %11
              t15: f64 = Register %11
            t18: f64,ch = CopyFromReg t0, Register:f64 %58
              t17: f64 = Register %58
          t45: f64,ch = CopyFromReg t0, Register:f64 %8
            t44: f64 = Register %8
        t48: f64,ch = CopyFromReg t0, Register:f64 %9
          t47: f64 = Register %9
      t50: f64 = ConstantFP<-2.997925e+10>
    t37: f64,ch = CopyFromReg t0, Register:f64 %6
      t36: f64 = Register %6
In function: _ZN22fieldPropagatorConstBz29ComputeStepAndPropagatedStateILb0EEEdddidRN7vecgeom4cuda8Vector3DIdEES5_RKNS2_13NavStateIndexERS6_
clang-13: error: clang frontend command failed with exit code 70 (use -v to see invocation)

I tried to link only one object file and the error above appeared when I executed:
after the CMakeFiles/example9.1.dir/electrons.dp.cpp.o is produced:

cd examples/Example9.1
clang++ -o example9.1 CMakeFiles/example9.1.dir/electrons.dp.cpp.o -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice

Environment:

• OS: Ubuntu 18.04.5 LTS
• Target device and vendor: NVIDIA GPU: GeForce RTX 2080 SUPER
• DPC++ version: 6a49170 (may release)

Did anyone encounter any of this before? Do you have any ideas that I could try ?

electrons-a2abbe.bc.txt

[steffenlarsen]

This looks similar to #3346 which was fixed with #3535 ~2 days after the DPC++ version you are using. Very unfortunate timing!

[dosarudaniel]

I managed to reproduce the error in this simple example.

Even if I am using the latest version of clang++ - c2d6cfa (which has the #3535 fix) when I execute make test12 I get:

fatal error: error in backend: Cannot select: t3: f64 = fcos t2
  t2: f64,ch = CopyFromReg t0, Register:f64 %27
    t1: f64 = Register %27
In function: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE55_16clES2_EUlNS0_7nd_itemILi3EEEE58_19
clang-13: error: clang frontend command failed with exit code 70 (use -v to see invocation)
clang version 13.0.0 (https://github.com/intel/llvm c2d6cfa0ffff2f45465e822447af318d01811b68)

I found a workaround to this problem with this piece of code such that when I am compiling for device I will always use the sycl implementation of cos . (Try to uncomment line 8 and 17 from the example )

#if (defined( __SYCL_DEVICE_ONLY__))
#define cos sycl::cos
#else
#define cos std::cos
#endif

Any other suggestions to fix this ? I am afraid it might interfere with math.h definitions in a larger project.

[steffenlarsen]

I see! Looks like there are two separate bugs in the report, the first having been fixed in #3535. I apologize for my confusion.

Yes, you will have to use the sycl::cos inside the kernels. However, these should also work as intended on host. If you really want to make a distinction between sycl::* and std::* then what you are doing with __SYCL_DEVICE_ONLY__ is the best solution I can think of off the top of my head. Maybe @alexbatashev knows another solution?

[alexbatashev]

I think, for now using sycl:: builtins is a go-to way.
However, we have an extension, that could make life easier for those, who port CUDA/C++ code to SYCL. If you look closely, those std:: math functions result in a call instruction in final binary (proof). libdevice takes advantage of that fact. It implements quite a few things from C and C++ standard libraries. In theory, one could simply build that library for PTX target and link it with the binary the same way we currently do it for SPIR-V target. I guess, adding a new target would be as easy as adding a few lines of CMake magic here. I am not sure if the driver would pick it up automatically for the PTX targets, but you can always manually feed the compiler with fallback object files.

[dosarudaniel]

As a follow up, I used this hack

#if (defined( __SYCL_DEVICE_ONLY__))
#define cos sycl::cos
#else
#define cos std::cos
#endif

for the following math functions:
log, exp, cos, sin, pow, frexp, ldexp, modf

Other used functions which did not need any special treatment:
abs, fabs, sqrt

[hdelan]

Hi @dosarudaniel , this should have been fixed in this PR #6482 . Can you check that this is no longer an issue? Thanks

[hdelan]

I am going to assume that this has been fixed by #6482 . I will close the ticket. Please reopen the issue if you have any further problems with this.