[SYCL] hw_emu: Syn check fail when using C-style arrays inside kernel

[j-stephan]

During one of my misguided attempts to use burst reads/writes (see the example at the bottom) I managed to break the compiler in hw_emu mode. Error log mentioned in console output:

clang: warning: argument unused during compilation: '-include /opt/xilinx/Vivado/2019.1/bin/../include/clc.h' [-Wunused-command-line-argument]
clang: warning: argument unused during compilation: '-include /opt/xilinx/Vivado/2019.1/bin/../include/etc/autopilot_ssdm_op.h' [-Wunused-command-line-argument]
warning: overriding the module target triple with spir64-unknown-unknown [-Woverride-module]
error: Syn check fail!

Command line and output:

$ clang++ -O3 -std=c++2a -fsycl -fsycl-targets=fpga64-xilinx-unknown-sycldevice main.cpp -o evil_array -lOpenCL -I/opt/xilinx/xrt/include

warning: Linking two modules of different data layouts: '/opt/xilinx/SDx/2019.1/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: /opt/xilinx/SDx/2019.1/bin/../lnx64/lib/libspir64-39-hls.bc' is 'fpga64-xilinx-none' whereas 'llvm-link' is 'spir64'

Invoking xocc Kernel Compilation
xocc: /opt/xilinx/SDx/2019.1/bin/xocc
--target: hw_emu
--platform: xilinx_u200_xdma_201830_2
Compiling kernel: xSYCL9318603472913316893
Outputting file to: /tmp/xSYCL9318603472913316893.xo
Input file is: /tmp/main_kernels-linked.xpirbc
Option Map File Used: '/opt/xilinx/SDx/2019.1/data/sdx/xocc/optMap.xml'

****** xocc v2019.1 (64-bit)
  **** SW Build 2552052 on Fri May 24 14:47:09 MDT 2019
    ** Copyright 1986-2019 Xilinx, Inc. All Rights Reserved.

Attempting to get a license: ap_opencl
Feature available: ap_opencl
INFO: [XOCC 60-1306] Additional information associated with this xocc compile can be found at:
	Reports: /home/jan/workspace/sycl-bugreports/xilinx/local_array/_x/reports/xSYCL9318603472913316893
	Log files: /home/jan/workspace/sycl-bugreports/xilinx/local_array/_x/logs/xSYCL9318603472913316893
INFO: [XOCC 60-585] Compiling for hardware emulation target
INFO: [XOCC 60-1316] Initiating connection to rulecheck server, at Fri Sep 20 17:45:27 2019
Running Rule Check Server on port:46287
INFO: [XOCC 60-1315] Creating rulecheck session with output '/home/jan/workspace/sycl-bugreports/xilinx/local_array/_x/reports/xSYCL9318603472913316893/xocc_compile_xSYCL9318603472913316893_guidance.html', at Fri Sep 20 17:45:28 2019
INFO: [XOCC 60-895]   Target platform: /opt/xilinx/platforms/xilinx_u200_xdma_201830_2/xilinx_u200_xdma_201830_2.xpfm
INFO: [XOCC 60-423]   Target device: xilinx_u200_xdma_201830_2
INFO: [XOCC 60-242] Creating kernel: 'xSYCL9318603472913316893'
ERROR: [XOCC 60-399] clang failed, please see log file for detail: '/home/jan/workspace/sycl-bugreports/xilinx/local_array/_x/xSYCL9318603472913316893/xSYCL9318603472913316893/xSYCL9318603472913316893/xSYCL9318603472913316893_clang.log'
ERROR: [XOCC 60-599] Kernel compilation failed to complete
ERROR: [XOCC 60-592] Failed to finish compilation
Option Map File Used: '/opt/xilinx/SDx/2019.1/data/sdx/xocc/optMap.xml'

****** xocc v2019.1 (64-bit)
  **** SW Build 2552052 on Fri May 24 14:47:09 MDT 2019
    ** Copyright 1986-2019 Xilinx, Inc. All Rights Reserved.

ERROR: [XOCC 60-602] Source file does not exist: /tmp/xSYCL9318603472913316893.xo
ERROR: [XOCC 60-623] Unsupported input file type specified.
/usr/bin/ld: /tmp/main-9bcdb9.o: file not recognized: file truncated
clang-9: error: sycl-link-xocc command failed with exit code 255 (use -v to see invocation)
clang-9: error: linker command failed with exit code 1 (use -v to see invocation)

Code:

#include <CL/sycl.hpp>

struct kernel;

auto loop(cl::sycl::accessor<std::size_t, 3, cl::sycl::access::mode::read_write> s)
{
    for(std::size_t z = 0; z < 42; ++z)
    {
        for(std::size_t y = 0; y < 42; ++y)
        {
            std::size_t in_row[42];
            std::size_t out_row[42];

            for(std::size_t x = 0; x < 42; ++x)
            {
                const auto id = cl::sycl::id<3>{x, y, z};
                in_row[x] = s[id];
            }

            for(std::size_t x = 0; x < 42; ++x)
                out_row[x] = in_row[x] + 42;

            for(std::size_t x = 0; x < 42; ++x)
            {
                const auto id = cl::sycl::id<3>{x, y, z};
                s[id] = out_row[x];
            }
        }
    }
}

auto main() -> int
{
    auto queue = cl::sycl::queue{};

    auto s_buf = cl::sycl::buffer<std::size_t, 3>{
                    cl::sycl::range<3>{42, 42, 42}};

    queue.submit([&](cl::sycl::handler& cgh)
    {
        auto s = s_buf.get_access<cl::sycl::access::mode::read_write>(cgh);

        cgh.single_task<kernel>([=]()
        {
            loop(s);
        });
                                                
    });
    queue.wait();

    return 0;
}

[j-stephan]

Follow-up: This also occurs when I try to use a condition result in a later division if one of the operands is a double. Example:

auto divisor = (condition) ? if_true : if_false;
/* ... */
auto result = dividend / divisor;

If either of dividend or divisor are a double, the compilation will fail with the Syn check fail! message. If they are both integral types the compilation will succeed.

[agozillon]

Thanks for the reports as always, I am trying to update to the most recent upstream (and internal) changes at the moment, so with any luck some of these things will be fixed. If that doesn't help we're also hoping to look into side stepping address space generation in our LLVM-IR which will fix a few of our hw_emu issues that stem from address spaces.

[agozillon]

As a follow up to my previous comment about side stepping address spaces for our FPGA, it won't be possible for a little while unfortunately as it'll take some tweaks by another team most likely and that's not likely to happen for a while.

We did move to a new address space inference method though that Intel swapped to a little while ago and pointed us towards, it seems a little more robust, still not perfect for us though.

[j-stephan]

I'm no longer able to reproduce this.