[SYCL] Missing and/or Incorrect xocc SPIR Math builtins

[agozillon]

Some math manglings from xocc's SPIR libraries appear to be incorrect or missing. In the sense that math functions correctly translated to their SPIR mangled names (found in: https://github.com/KhronosGroup/SPIR-Tools/wiki/SPIR-2.0-built-in-functions) result in runtime errors caused by missing symbols.

Two known examples are the mad(float, float, float) (_Z3madfff) function used in the test math_mangle.cpp and fmin(float2, float2) (_Z4fminDv2_fDv2_f) found in the test vector_math.cpp.

In at least fmin's case it appears to be that the mangling of the arguments type is different within xocc's SPIR library.

This issue will require some further investigation, but a preliminary look at the issue makes it seem like it will require modifications to the InSPIRation pass to support these alternate SPIR manglings in certain cases.

[j-stephan]

The issue appears to be fixed at least for cl::sycl::rsqrt - I could compile my N-body example in sw_emu mode without any complaints about missing symbols. Using the master branch here, the master branch of XRT and SDAccel 2019.1, everything on Ubuntu 19.04 and Linux 5.0.

[agozillon]

interesting, not too sure why as it wasn't any knowing fix on my end and it still isn't working on my machine, but I'll take a free win. But i'll have to investigate into why still.

Could you please uncomment this line https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/test/xocc_tests/simple_tests/math_mangle.cpp#L51 and give math_mangle.cpp a try

and also if possible the vector_math.cpp test, after uncommenting the code block here: https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/test/xocc_tests/simple_tests/vector_math.cpp#L27

I have a suspicion they're connected to the rsqrt issue you were having, so it would be nice if they magically fixed themselves too.

As an aside, relating to your nbody example, you'll probably have the handlers copy method breaking when using the master branch now and ICE'ng your compiler in hw_emu, I applied a fix to that on the Next branch recently. It's another bug I encountered when running your example recently, still looking into #32 though.

Edit: If they do work, could you if at all possible ldd the dltmp file inside the hidden .run folder that's generated into the directory you compile your example in. e.g:

ldd /sycl/test/xocc_tests/issue_related/.run/27980/sw_emu/device0/binary_1/dltmp linux-vdso.so.1 (0x00007ffd1ddcd000) libgtk3-nocsd.so.0 => /usr/lib/x86_64-linux-gnu/libgtk3-nocsd.so.0 (0x00007f306fc33000) libstdc++.so.6 => /usr/lib/x86_64-linux-gnu/libstdc++.so.6 (0x00007f306faa8000) libm.so.6 => /lib/x86_64-linux-gnu/libm.so.6 (0x00007f306f91b000) libgcc_s.so.1 => /lib/x86_64-linux-gnu/libgcc_s.so.1 (0x00007f306f901000) libc.so.6 => /lib/x86_64-linux-gnu/libc.so.6 (0x00007f306f717000) libdl.so.2 => /lib/x86_64-linux-gnu/libdl.so.2 (0x00007f306f70f000) libpthread.so.0 => /lib/x86_64-linux-gnu/libpthread.so.0 (0x00007f306f6ee000) /lib64/ld-linux-x86-64.so.2 (0x00007f306fe95000)

and check the versions of the linked against libraries/packages (if at all possible)?

Sorry for the requests, a bit of a pain I know, but it sadly doesn't appear to work on our 19.04 machine so perhaps we're missing some key version of a library (which would be important to add to the documentation).

[j-stephan]

Ugh, never mind. I falsely remembered this being a compile-time issue, now I see it still fails during runtime. The simple_tests do too, unfortunately. Sorry for the noise, it's way past midnight here.

[agozillon]

Ah unfortunate the mystery remains open for now, all good I know the feeling!

[j-stephan]

So, any news on this? I was just hit by this again when trying to run sqrt(double).

I might have some free time in two months or so, I could try and have a look at this. Where is the xocc SPIR library located? Or is this something that has to be done by Xilinx employees?

[agozillon]

I have not looked into this recently and probably won't get a chance to anytime soon. I'm unsure this is fixable externally, from the SYCL compilation perspective we just link against the libspir64-39-hls.bc library to resolve SPIR symbols (https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/tools/sycl-xocc/bin/sycl-xocc#L76). There are some other SPIR libraries in the same directory that I've not tried linking against, perhaps they don't suffer the missing symbols but I find it unlikely.

The reason I think this fix is only internally fixable is that the library compilation is something HLS does when shipping the Xilinx SDK and the source code (that at least I'm aware of) doesn't appear to be shipped with the SDK and I'm relatively sure it's not open source at the moment (The closest thing I can find is this, which appears to be a workaround and nowhere near the full library: Vivado/2019.1/src/lib_hlsmcl.cpp).

The only way that I can think of introspecting this problem externally right now is by disassembling libspir64-39-hls.bc file using llvm, but I think that'll only get you so far. However, it might give enough of a hint that we can open up an internal bug against it with enough information that someone on the HLS team will be interested in fixing it.