[SYCL][FPGA] Adding a wrapper header for __builtin_intel_fpga_mem

[mohammadfawaz]

This patch adds a new wrapper header for the Clang builtin __builtin_intel_fpga_mem

The SPIRV spec for this featuer is here: https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_fpga_memory_accesses.asciidoc

The Intel FPGA lsu class is implemented in CL/sycl/intel/fpga_lsu.hpp which
is included in CL/sycl/intel/fpga_extensions.hpp.

The class cl::sycl::intel::lsu allows users to explicitly request that the
implementation of a global memory access is configured in a certain way. The
class has two member functions, load() and store() which allow loading from
and storing to a global_ptr, respectively, and is templated on 4 optional paremeters.

[mohammadfawaz]

@MrSidims any updates on this? Are you an approving reviewer?

[MrSidims]

LGTM with a few nits.

[MrSidims]

For the future, it's recommended not to amend patches before pushing, so changes could be tracked from version to version.

[mohammadfawaz]

For the future, it's recommended not to amend patches before pushing, so changes could be tracked from version to version.

Sounds good! Thank you @MrSidims .

[keryell]

Nice to see higher-level constructs for FPGA in SYCL, instead of ugly attributes or #pragma! :-)

With this it is possible then to build higher-level libraries to wrap accessors and provide different behaviors and hiding explicit load/store.

[keryell]

Thanks, the example is simpler.
But I have the feeling that the user-facing API could be streamlined and more obvious.

[mohammadfawaz]

@keryell Thank you for your comments. The code does indeed look better and the use model is more intuitive.

[rbegam]

LGTM

[bader]

@jbrodman, could you take a look at the documentation for this extension, please?

[mohammadfawaz]

@jbrodman, could you take a look at the documentation for this extension, please?
@jbrodman @bader A gentle reminder to take a look at this at your convenience. Thanks!

[MrSidims]

@mohammadfawaz looks like the PR as is is going nowhere, I suggest to split it to 2 PRs: implementation itself and the spec. Implementation PR can be approved (basically it's already approved) and merged right away. Spec can be merged later.

[mohammadfawaz]

@mohammadfawaz looks like the PR as is is going nowhere, I suggest to split it to 2 PRs: implementation itself and the spec. Implementation PR can be approved (basically it's already approved) and merged right away. Spec can be merged later.

Done. I removed the spec from this PR. Will open another one for the spec.

[MrSidims]

Approved. @mohammadfawaz the only thing, that I want to ask you to change is the PR description. @bader usually uses it to create a commit message for a squashed commit during merging, hence I advise to replace "The .md file I added contains an example of how I see this being used. For now, this is only useful for global_ptrs." with a little extract from the spec itself (or just refer to the spec PR, when it's ready, but the first way is slightly preferable).