[SYCL][FPGA] Update LSU builtins template to allow multi_ptr (#3321)

Updated LSU headers to allow multi_ptr instead of global_ptr.
Added address space checks to builtins.

Author: aditikum

sycl/include/CL/sycl/INTEL/fpga_lsu.hpp

@@ -47,7 +47,9 @@ template <class... _mem_access_params> class lsu final {
 public:
   lsu() = delete;
 
-  template <typename _T> static _T load(sycl::global_ptr<_T> Ptr) {
+  template <typename _T, access::address_space _space>
+  static _T load(sycl::multi_ptr<_T, _space> Ptr) {
+    check_space<_space>();
     check_load();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
     return *__builtin_intel_fpga_mem((_T *)Ptr,
@@ -59,7 +61,9 @@ template <class... _mem_access_params> class lsu final {
 #endif
   }
 
-  template <typename _T> static void store(sycl::global_ptr<_T> Ptr, _T Val) {
+  template <typename _T, access::address_space _space>
+  static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) {
+    check_space<_space>();
     check_store();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
     *__builtin_intel_fpga_mem((_T *)Ptr,
@@ -92,6 +96,14 @@ template <class... _mem_access_params> class lsu final {
 
   static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
 
+  template <access::address_space _space> static void check_space() {
+    static_assert(_space == access::address_space::global_space ||
+                      _space == access::address_space::global_device_space ||
+                      _space == access::address_space::global_host_space,
+                  "lsu controls are only supported for global_ptr, "
+                  "device_ptr, and host_ptr objects");
+  }
+
   static void check_load() {
     static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
                   "unable to implement a cache without a burst coalescer");

sycl/test/extensions/fpga.cpp

@@ -59,16 +59,14 @@ int main() {
   /*Check LSU interface*/
   {
     cl::sycl::buffer<int, 1> output_buffer(1);
-    cl::sycl::buffer<int, 1> input_buffer(1);
+    auto *in_ptr = cl::sycl::malloc_host<int>(1, Queue.get_context());
 
     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();
+        cl::sycl::host_ptr<int> input_ptr(in_ptr);
         auto output_ptr = output_accessor.get_pointer();
 
         using PrefetchingLSU =