Compiler crash in hw_emu mode during Aggressive Dead Code Elimination

[j-stephan]

This is not related to #22. I'm using the code from #28 with the following modifications:

• all discard_write accesses are changed to write (see Missing access mode: discard_write #31)
• the call to cl::sycl::rsqrt in line 68 is changed to 1.f / cl::sycl::sqrt as there seems to be a missing symbol for the Alveo U200 emulation:

/tools/Xilinx/SDx/2018.3/data/emulation/unified/cpu_em/generic_pcie/model/genericpciemodel: symbol lookup error: /home/jstephan/Downloads/.run/8220/sw_emu/device0/binary_1/dltmp: undefined symbol: rsqrtf

With the current compiler and OpenCL runtime I am able to successfully run the program in sw_emu mode. When changing to hw_emu the compilation will crash:

$ clang++ -std=c++2a -fsycl -fsycl-targets=fpga64-xilinx-unknown-sycldevice nbody.cpp -o nbody.hw_emu -lOpenCL -I/opt/xilinx/xrt/include/
warning: Linking two modules of different data layouts: '/tools/Xilinx/SDx/2018.3/bin/../lnx64/lib/libspir64-39-hls.bc' is 'e-m:e-i64:64-i128:128-i256:256-i512:512-i1024:1024-i2048:2048-i4096:4096-n8:16:32:64-S128-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024' whereas 'llvm-link' is 'e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024'

warning: Linking two modules of different target triples: /tools/Xilinx/SDx/2018.3/bin/../lnx64/lib/libspir64-39-hls.bc' is 'fpga64-xilinx-none' whereas 'llvm-link' is 'spir64'


****** xocc v2018.3 (64-bit)
  **** SW Build 2405991 on Thu Dec  6 23:36:41 MST 2018
    ** Copyright 1986-2018 Xilinx, Inc. All Rights Reserved.

Attempting to get a license: ap_opencl
Feature available: ap_opencl jstephan@ite196  ~/Downloads 
INFO: [XOCC 60-1306] Additional information associated with this xocc compile can be found at:
	Reports: /home/jstephan/Downloads/_x/reports/x13642406682781150725
	Log files: /home/jstephan/Downloads/_x/logs/x13642406682781150725
INFO: [XOCC 60-585] Compiling for hardware emulation target
Running SDx Rule Check Server on port:40935
INFO: [XOCC 60-895]   Target platform: /opt/xilinx/platforms/xilinx_u200_xdma_201830_1/xilinx_u200_xdma_201830_1.xpfm
WARNING: [XOCC 74-49] Failed to read in file /tools/Xilinx/SDx/2018.3/platforms/achromatic during platform validation.
INFO: [XOCC 60-423]   Target device: xilinx_u200_xdma_201830_1
INFO: [XOCC 60-242] Creating kernel: 'x13642406682781150725'
ERROR: [XOCC 17-1309] Gcc: #18 0x00007f47c721409b __libc_start_main /build/glibc-B9XfQf/glibc-2.28/csu/../csu/libc-start.c:308:16
ERROR: [XOCC 60-398] clang failed
ERROR: [XOCC 60-599] Kernel compilation failed to complete
ERROR: [XOCC 60-592] Failed to finish compilation

The log file reports the following:

clang: warning: argument unused during compilation: '-include /tools/Xilinx/Vivado/2018.3/bin/../include/clc.h' [-Wunused-command-line-argument]
clang: warning: argument unused during compilation: '-include /tools/Xilinx/Vivado/2018.3/bin/../include/etc/autopilot_ssdm_op.h' [-Wunused-command-line-argument]
warning: overriding the module target triple with spir64-unknown-unknown [-Woverride-module]
#0 0x00000000018d831a llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x18d831a)
#1 0x00000000018d645e llvm::sys::RunSignalHandlers() (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x18d645e)
#2 0x00000000018d6582 SignalHandler(int) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x18d6582)
#3 0x00007fdb2daabdd0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12dd0)
#4 0x00000000016be28a (anonymous namespace)::AggressiveDeadCodeElimination::collectLiveScopes(llvm::DILocation const&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x16be28a)
#5 0x00000000016beeb7 (anonymous namespace)::AggressiveDeadCodeElimination::markLive(llvm::Instruction*) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x16beeb7)
#6 0x00000000016bfa39 (anonymous namespace)::AggressiveDeadCodeElimination::markLiveInstructions() (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x16bfa39)
#7 0x00000000016c78e3 (anonymous namespace)::ADCELegacyPass::runOnFunction(llvm::Function&) [clone .part.347] (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x16c78e3)
#8 0x0000000001461e52 llvm::FPPassManager::runOnFunction(llvm::Function&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1461e52)
#9 0x0000000001461ee3 llvm::FPPassManager::runOnModule(llvm::Module&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1461ee3)
#10 0x00000000014627f4 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x14627f4)
#11 0x0000000001a9978a 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> >) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1a9978a)
#12 0x00000000022939ad clang::CodeGenAction::ExecuteAction() (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x22939ad)
#13 0x0000000001e55ef6 clang::FrontendAction::Execute() (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1e55ef6)
#14 0x0000000001e22d56 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1e22d56)
#15 0x0000000001eea523 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x1eea523)
#16 0x0000000000a5ed08 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0xa5ed08)
#17 0x00000000009e43fe main (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0x9e43fe)
#18 0x00007fdb2cd5709b __libc_start_main /build/glibc-B9XfQf/glibc-2.28/csu/../csu/libc-start.c:308:16
#19 0x0000000000a5c399 _start (/tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang+0xa5c399)
Stack dump:
0.      Program arguments: /tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/bin/clang -cc1 -triple spir64-unknown-unknown -emit-llvm-bc -emit-llvm-uselists -disable-free -disable-llvm-verifier -discard-value-names -main-file-name -linked.xpirbc -mrelocation-model static -mthread-model posix -fhls -mhls-ir -disable-O0-optnone -mdisable-fp-elim -no-integrated-as -mconstructor-aliases -dwarf-column-info -debug-info-kind=limited -dwarf-version=4 -debugger-tuning=gdb -coverage-notes-file /home/jstephan/Downloads/_x/x13642406682781150725/x13642406682781150725/x13642406682781150725/x13642406682781150725.clc.00.gcno -resource-dir /tools/Xilinx/Vivado/2018.3/lnx64/tools/clang-3.9-csynth/lib/clang/7.0.0 -Wall -std=CL1.2 -fno-dwarf-directory-asm -fdebug-compilation-dir /home/jstephan/Downloads/_x/x13642406682781150725/x13642406682781150725/x13642406682781150725 -ferror-limit 19 -fmessage-length 0 -cl-kernel-arg-info -fgnu89-inline -fobjc-runtime=gcc -fdiagnostics-show-option -mhls-ir -mlink-opencl-bitcode /tools/Xilinx/Vivado/2018.3/lnx64/lib/libspir64-39-hls.bc -mllvm -hls-bitcode-version=3.1 -mllvm -hls-top-function-name=x13642406682781150725 -disable-llvm-optzns -mllvm -xcl-xmlinfo=/home/jstephan/Downloads/_x/x13642406682781150725/x13642406682781150725/x13642406682781150725/kernel.xml -o /home/jstephan/Downloads/_x/x13642406682781150725/x13642406682781150725/x13642406682781150725/x13642406682781150725.clc.00.bc -x ir /tmp/-linked.xpirbc
1.      Per-module optimization passes
2.      Running pass 'Function Pass Manager' on module '/tmp/-linked.xpirbc'.
3.      Running pass 'Aggressive Dead Code Elimination' on function '@x13642406682781150725'
clang: error: unable to execute command: Segmentation fault (core dumped)
clang: error: clang frontend command failed due to signal (use -v to see invocation)
clang version 7.0.0
Target: spir64-unknown-unknown
Thread model: posix
InstalledDir: /home/jstephan/software/sycl/bin
clang: note: diagnostic msg: PLEASE submit a bug report to http://llvm.org/bugs/ and include the crash backtrace, preprocessed source, and associated run script.
clang: note: diagnostic msg: Error generating preprocessed source(s) - no preprocessable inputs.

[agozillon]

I'll try look into both issues after pull request #30 is merged in, the first issue relating to a missing math function may be related to this: #11 and a "quick fix" for it could be out of my hands.

The second issue I'm not so sure without testing on my machine (which I'll get around to asap) and tinkering a little. It looks like one of xocc's hw_emu specific passes is choking on some IR that's generated in the first kernels entry point though.

As always thanks for reporting the trouble you run into!

[j-stephan]

The next branch also works with this example. Closing.

[agozillon]

nice, thanks for the report! would it be possible to have a copy of the variation of the test that works so I can stick it into our very rudimentary regression suite? You could pull request it in if you'd like or just attach it here as a file when you get a chance and I can add it in.