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

sycl/include/CL/sycl/intel/fpga_extensions.hpp

@@ -8,5 +8,6 @@
 
 #pragma once
 #include <CL/sycl/intel/fpga_device_selector.hpp>
+#include <CL/sycl/intel/fpga_lsu.hpp>
 #include <CL/sycl/intel/fpga_reg.hpp>
 #include <CL/sycl/intel/pipes.hpp>

sycl/include/CL/sycl/intel/fpga_lsu.hpp

@@ -0,0 +1,113 @@
+//==-------------- fpga_lsu.hpp --- SYCL FPGA Reg Extensions ---------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "fpga_utils.hpp"
+#include <CL/sycl/detail/defines.hpp>
+#include <CL/sycl/pointers.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace intel {
+constexpr uint8_t BURST_COALESCE = 0x1;
+constexpr uint8_t CACHE = 0x2;
+constexpr uint8_t STATICALLY_COALESCE = 0x4;
+constexpr uint8_t PREFETCH = 0x8;
+
+template <int32_t N> struct burst_coalesce_impl {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t N> struct cache {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t N> struct prefetch_impl {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 0;
+};
+
+template <int32_t N> struct statically_coalesce_impl {
+  static constexpr int32_t value = N;
+  static constexpr int32_t default_value = 1;
+};
+
+template <bool B> using burst_coalesce = burst_coalesce_impl<B>;
+template <bool B> using prefetch = prefetch_impl<B>;
+template <bool B> using statically_coalesce = statically_coalesce_impl<B>;
+
+template <class... mem_access_params> class lsu final {
+public:
+  lsu() = delete;
+
+  template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
+    check_load();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    return *__builtin_intel_fpga_mem((T *)Ptr,
+                                     _burst_coalesce | _cache |
+                                         _dont_statically_coalesce | _prefetch,
+                                     _cache_val);
+#else
+    return *Ptr;
+#endif
+  }
+
+  template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
+    check_store();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    *__builtin_intel_fpga_mem((T *)Ptr,
+                              _burst_coalesce | _cache |
+                                  _dont_statically_coalesce | _prefetch,
+                              _cache_val) = Val;
+#else
+    *Ptr = Val;
+#endif
+  }
+
+private:
+  static constexpr int32_t _burst_coalesce_val =
+      GetValue<burst_coalesce_impl, mem_access_params...>::value;
+  static constexpr uint8_t _burst_coalesce =
+      _burst_coalesce_val == 1 ? BURST_COALESCE : 0;
+
+  static constexpr int32_t _cache_val =
+      GetValue<cache, mem_access_params...>::value;
+  static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0;
+
+  static constexpr int32_t _statically_coalesce_val =
+      GetValue<statically_coalesce_impl, mem_access_params...>::value;
+  static constexpr uint8_t _dont_statically_coalesce =
+      _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0;
+
+  static constexpr int32_t _prefetch_val =
+      GetValue<prefetch_impl, mem_access_params...>::value;
+  static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0;
+
+  static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
+
+  static void check_load() {
+    static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
+                  "unable to implement a cache without a burst coalescer");
+    static_assert(_prefetch == 0 || _burst_coalesce == 0,
+                  "unable to implement a prefetcher and a burst coalescer "
+                  "simulataneously");
+    static_assert(
+        _prefetch == 0 || _cache == 0,
+        "unable to implement a prefetcher and a cache simulataneously");
+  }
+  static void check_store() {
+    static_assert(_cache == 0, "unable to implement a store LSU with a cache.");
+    static_assert(_prefetch == 0,
+                  "unable to implement a store LSU with a prefetcher.");
+  }
+};
+} // namespace intel
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/fpga_utils.hpp

@@ -0,0 +1,33 @@
+//==------------- fpga_utils.hpp --- SYCL FPGA Reg Extensions --------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/detail/defines.hpp>
+#include <CL/sycl/stl.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace intel {
+
+template <template <int32_t> class Type, class T>
+struct MatchType : std::is_same<Type<T::value>, T> {};
+
+template <template <int32_t> class Type, class... T> struct GetValue {
+  static constexpr auto value = Type<0>::default_value;
+};
+
+template <template <int32_t> class Type, class T1, class... T>
+struct GetValue<Type, T1, T...> {
+  static constexpr auto value =
+      std::conditional<MatchType<Type, T1>::value, T1,
+                       GetValue<Type, T...>>::type::value;
+};
+} // namespace intel
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/fpga_tests/fpga_lsu.cpp

@@ -0,0 +1,80 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+//==----------------- fpga_lsu.cpp - SYCL FPGA LSU test --------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <CL/sycl/intel/fpga_extensions.hpp>
+
+// TODO: run is disabled, since no support added in FPGA backend yet. Check
+// implementation correctness from CXX and SYCL languages perspective.
+
+int test_lsu(cl::sycl::queue Queue) {
+  int output_data[2];
+  for (size_t i = 0; i < 2; i++) {
+    output_data[i] = -1;
+  }
+
+  int input_data[2];
+  for (size_t i = 0; i < 2; i++) {
+    input_data[i] = i + 1;
+  }
+
+  {
+    cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
+    cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
+
+    Queue.submit([&](cl::sycl::handler &cgh) {
+      auto output_accessor =
+          output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
+      auto input_accessor =
+          input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
+
+      cgh.single_task<class kernel>([=] {
+        auto input_ptr = input_accessor.get_pointer();
+        auto output_ptr = output_accessor.get_pointer();
+
+        using PrefetchingLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::prefetch<true>,
+                                 cl::sycl::intel::statically_coalesce<false>>;
+
+        using BurstCoalescedLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
+                                 cl::sycl::intel::statically_coalesce<false>>;
+
+        using CachingLSU =
+            cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
+                                 cl::sycl::intel::cache<1024>,
+                                 cl::sycl::intel::statically_coalesce<false>>;
+
+        using PipelinedLSU = cl::sycl::intel::lsu<>;
+
+        int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
+        int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
+
+        BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
+        PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+      });
+    });
+  }
+
+  for (int i = 0; i < 2; i++) {
+    if (output_data[i] != input_data[i]) {
+      std::cout << "Unexpected read from output_data: " << output_data[i]
+                << ", v.s. expected " << input_data[i] << std::endl;
+
+      return -1;
+    }
+  }
+  return 0;
+}
+
+int main() {
+  cl::sycl::queue Queue{cl::sycl::intel::fpga_emulator_selector{}};
+
+  return test_lsu(Queue);
+}