[SYCL][FPGA] Adding function wrapper for function-level loop fusion attributes

[aaronkintel]

The function attributes [[intel::loop_fuse(N)]] and
[[intel::loop_fuse_independent(N)]] are being replaced by equivalent
function wrappers by the inclusion of a header file.

These attributes can be accessed now by
sycl::ext::intel::fpga_loop_fuse(F) and
sycl::ext::intel::fpga_loop_fuse_independent(F) for a function or
lambda F.

There is a lit test for the functionality of the loop fusion attributes. The attributes are tested on functions elsewhere (e.g. here and here).

[keryell]

How is it used in real life with the default parameter in first position?
Do you have an example?

[aaronkintel]

The default first parameter mirrors the behaviour of HLS' #pragma loop_fuse (see here). Apparently, I cannot attach *.cpp files here, but I got the following code to compile:

#include <CL/sycl.hpp>
#include <sycl/ext/intel/fpga_extensions.hpp>

constexpr int v = 123456;

int foo() {
  return 42;
}

int main() {
#if defined(FPGA_EMULATOR)
  sycl::ext::intel::fpga_emulator_selector device_selector;
#else
  sycl::ext::intel::fpga_selector device_selector;
#endif

  sycl::queue Queue(device_selector);
  sycl::buffer<int, 1> Buf{sycl::range{1}};

  Queue.submit([&](sycl::handler &CGH) {
    sycl::accessor Acc(Buf, CGH, sycl::write_only);
    CGH.single_task<class Test>([=]() {

      int foo_val = 0;
      sycl::ext::intel::fpga_loop_fuse_independent([&]{ foo_val = foo(); });
      sycl::ext::intel::fpga_loop_fuse<v>([&]{ Acc[0] = foo_val; });
    });
  });

  Queue.wait();

  auto Acc = Buf.template get_access<sycl::access::mode::read_write>();
  assert(Acc[0] == 42 && "Value mismatch");

  return 0;
}

[keryell]

This looks like a good starting test for this feature.
Would it makes sense to have loops somewhere in the example, if it is about loops?

[aaronkintel]

In this compiler, the loop_fuse and loop_fuse_independent attributes add function metadata but do not touch any loops. The implementation and testing of the loop fusion capabilities are in the FPGA compiler.

[MrSidims]

The patch is LGTM, but I also wonder why and whether the first template parameter should be defaulted.

[MrSidims]

@s-kanaev could you please review from runtime-reviewers perspective?

[s-kanaev]

Seems legit.
Should there be any sort of test for this feature?

[aaronkintel]

Seems legit. Should there be any sort of test for this feature?

Updated description.

[aaronkintel]

@MrSidims Would you mind taking another look at my changes?