Error executing basic test code for FPGA (oneAPI/OpenCL)

Greetings,

We are running some experiments and we have detected that oneAPI for FGPA is not performing well, quite below the previous pure OpenCL implementation (same system, fpga, etc). We tested with different codes, and in even a simple Matrix Multiplication the OpenCL code/runtime performs much better than in a pure oneAPI one. So, we thought about doing OpenCL interoperability with oneAPI (as in this guide https://www.intel.com/content/www/us/en/developer/articles/technical/interoperability-dpcpp-sycl-opencl.html), to see if DPC++ can avoid all the overhead/wrong optimization steps of oneAPI by using something closer to the OpenCL implementation/runtime optimizations. We don't know what oneAPI is adding during the compilation/execution steps, but now using the FPGA is not worthwhile.

The fpga emulator works with the interoperability, but the real fpga does not. Here are our steps:

Error

user@node:~/interoperability$ ./ingest_fpga
Device: s10gx : Stratix 10 Reference Platform (acls10_ref0)
terminate called after throwing an instance of 'cl::sycl::feature_not_supported'
  what():  Online compilation is not supported by this device -3 (CL_COMPILER_NOT_AVAILABLE)

Code

#include <CL/sycl.hpp>
#include <iostream>
#include <array>

#include <CL/sycl/INTEL/fpga_extensions.hpp>
#include "dpc_common.hpp"

using namespace cl::sycl;

int main()
{
        const size_t szKernelData = 32;
        std::array<float, szKernelData> kernelData;
        kernelData.fill(-99.f);
        range<1> r(szKernelData);
        sycl::INTEL::fpga_selector _device;
        queue q{_device};
        //queue q{fpga_selector()};
        std::cout << "Device: " << q.get_device().get_info<info::device::name>() << "\n";
        program p(q.get_context());

        p.build_with_source(
          R"CLC(
            kernel void sinf_test(global float* data) {
              data[get_global_id(0)] = sin(get_global_id(0)*2*M_PI_F/get_global_size(0)) ;
            }
          )CLC", "-cl-std=CL1.2");
        {
                buffer<float, 1> b(kernelData.data(), r);

                q.submit([&](handler& cgh) {
                        auto b_accessor = b.get_access<access::mode::read_write>(cgh);
                        cgh.set_args(b_accessor);
                        cgh.parallel_for(r, p.get_kernel("sinf_test"));
                });
        }
        for(auto& elem : kernelData)
                std::cout << std::defaultfloat << elem << " " << std::hexfloat << elem << std::endl;
        return 0;
}

Compilation

dpcpp -fintelfpga -Xshardware -fsycl-unnamed-lambda ingest.cpp -std=c++17 -o ingest_fpga
# Output
WARNING: No kernels specified.

Any idea?