Add SYCL parallel backend integration

Author: allnes

app/SYCL/CMakeLists.txt

@@ -28,6 +28,7 @@ if(WIN32)
         "--acpp-cpu-cxx=${_sycl_cpu_cxx}"
         -std=c++20
         -fms-runtime-lib=dll
+        -DITLABAI_HAS_SYCL
         -D_USE_MATH_DEFINES
         -I "${CMAKE_SOURCE_DIR}/include"
     )
@@ -58,6 +59,8 @@ if(WIN32)
     target_sources(SYCL_Example PRIVATE "${SYCL_EXAMPLE_KERNEL_OBJECT}")
 else()
     add_library(SYCL_Example_kernel OBJECT sycl_kernel.cpp)
+    target_compile_definitions(SYCL_Example_kernel PRIVATE ITLABAI_HAS_SYCL)
+    target_link_libraries(SYCL_Example_kernel PRIVATE AdaptiveCpp::acpp-rt TBB_unified Kokkos_imported)
     add_sycl_to_target(TARGET SYCL_Example_kernel SOURCES
         sycl_kernel.cpp
     )

app/SYCL/sycl_example.cpp

@@ -1,6 +1,6 @@
 #include <cmath>
 #include <iostream>
-#include <sycl/sycl.hpp>
+#include <string>
 #include <vector>
 
 #include "graph/runtime_options.hpp"
@@ -15,8 +15,8 @@ bool almost_equal(float lhs, float rhs) {
 
 }  // namespace
 
-void run_sycl_kernel(sycl::queue& queue, const float* input, float* output,
-                     std::size_t count);
+std::string sycl_device_name();
+void run_sycl_kernel(const float* input, float* output, std::size_t count);
 
 int main() {
   try {
@@ -40,13 +40,10 @@ int main() {
     }
 
     const std::size_t count = relu_output.size();
-    sycl::queue queue(sycl::default_selector_v);
-    std::cout << "SYCL device: "
-              << queue.get_device().get_info<sycl::info::device::name>()
-              << '\n';
+    std::cout << "SYCL device: " << sycl_device_name() << '\n';
 
     std::vector<float> sycl_output(count);
-    run_sycl_kernel(queue, relu_output.data(), sycl_output.data(), count);
+    run_sycl_kernel(relu_output.data(), sycl_output.data(), count);
 
     const std::vector<float> expected_sycl = {1.0F, 1.0F, 1.0F, 3.0F, 5.0F};
     for (std::size_t i = 0; i < count; ++i) {

app/SYCL/sycl_kernel.cpp

@@ -1,20 +1,27 @@
 #include <cstddef>
-#include <sycl/sycl.hpp>
-class SyclExampleTransformKernel;
+#include <string>
 
-void run_sycl_kernel(sycl::queue& queue, const float* input, float* output,
-                     std::size_t count) {
-  sycl::buffer<const float, 1> input_buffer(input, sycl::range<1>(count));
-  sycl::buffer<float, 1> output_buffer(output, sycl::range<1>(count));
+#include "parallel/parallel.hpp"
 
-  queue.submit([&](sycl::handler& handler) {
-    auto input_acc = input_buffer.get_access<sycl::access::mode::read>(handler);
-    auto output_acc =
-        output_buffer.get_access<sycl::access::mode::write>(handler);
-    handler.parallel_for<SyclExampleTransformKernel>(sycl::range<1>(count),
-                                                     [=](sycl::id<1> index) {
-      output_acc[index] = input_acc[index] * 2.0F + 1.0F;
-    });
-  });
-  queue.wait_and_throw();
+std::string sycl_device_name() {
+  return it_lab_ai::parallel::sycl_device_name();
+}
+
+void run_sycl_kernel(const float* input, float* output, std::size_t count) {
+  auto shared_input = it_lab_ai::parallel::make_sycl_shared_buffer<float>(count);
+  auto shared_output =
+      it_lab_ai::parallel::make_sycl_shared_buffer<float>(count);
+
+  shared_input.copy_from(input);
+  const float* input_data = shared_input.data();
+  float* output_data = shared_output.data();
+
+  it_lab_ai::parallel::parallel_for(
+      count,
+      [=](std::size_t index) {
+        output_data[index] = input_data[index] * 2.0F + 1.0F;
+      },
+      it_lab_ai::parallel::Backend::kSycl);
+
+  shared_output.copy_to(output);
 }

include/parallel/backends.hpp

@@ -9,9 +9,16 @@
 #include <cstdint>
 #include <functional>
 #include <limits>
+#include <stdexcept>
+#include <string>
 #include <thread>
+#include <utility>
 #include <vector>
 
+#ifdef ITLABAI_HAS_SYCL
+#  include <sycl/sycl.hpp>
+#endif
+
 namespace it_lab_ai {
 namespace parallel {
 
@@ -20,7 +27,8 @@ enum class Backend : std::uint8_t {
   kThreads = 1,
   kTbb = 2,
   kOmp = 3,
-  kKokkos = 4
+  kKokkos = 4,
+  kSycl = 5
 };
 
 struct Options {
@@ -149,5 +157,100 @@ inline void impl_kokkos(std::size_t count,
   Kokkos::fence();
 }
 
+#ifdef ITLABAI_HAS_SYCL
+inline sycl::queue& sycl_queue() {
+  static sycl::queue queue{sycl::default_selector_v};
+  return queue;
+}
+
+inline std::string sycl_device_name() {
+  return sycl_queue().get_device().get_info<sycl::info::device::name>();
+}
+
+template <typename T>
+class SyclSharedBuffer {
+ public:
+  explicit SyclSharedBuffer(std::size_t count)
+      : count_(count), data_(count == 0 ? nullptr : sycl::malloc_shared<T>(
+                                                       count, sycl_queue())) {}
+
+  SyclSharedBuffer(const SyclSharedBuffer&) = delete;
+  SyclSharedBuffer& operator=(const SyclSharedBuffer&) = delete;
+
+  SyclSharedBuffer(SyclSharedBuffer&& other) noexcept
+      : count_(other.count_), data_(other.data_) {
+    other.count_ = 0;
+    other.data_ = nullptr;
+  }
+
+  SyclSharedBuffer& operator=(SyclSharedBuffer&& other) noexcept {
+    if (this != &other) {
+      release();
+      count_ = other.count_;
+      data_ = other.data_;
+      other.count_ = 0;
+      other.data_ = nullptr;
+    }
+    return *this;
+  }
+
+  ~SyclSharedBuffer() { release(); }
+
+  T* data() { return data_; }
+  const T* data() const { return data_; }
+  std::size_t size() const { return count_; }
+
+  void copy_from(const T* source) {
+    for (std::size_t i = 0; i < count_; ++i) {
+      data_[i] = source[i];
+    }
+  }
+
+  void copy_to(T* destination) const {
+    for (std::size_t i = 0; i < count_; ++i) {
+      destination[i] = data_[i];
+    }
+  }
+
+ private:
+  void release() {
+    if (data_ != nullptr) {
+      sycl::free(data_, sycl_queue());
+      data_ = nullptr;
+      count_ = 0;
+    }
+  }
+
+  std::size_t count_;
+  T* data_;
+};
+
+template <typename T>
+inline SyclSharedBuffer<T> make_sycl_shared_buffer(std::size_t count) {
+  return SyclSharedBuffer<T>(count);
+}
+
+template <typename Func>
+inline void impl_sycl(std::size_t count, Func&& func, const Options&) {
+  if (count == 0) {
+    return;
+  }
+
+  auto& queue = sycl_queue();
+  queue.parallel_for(sycl::range<1>(count),
+                     [func = std::forward<Func>(func)](sycl::id<1> index) {
+    func(static_cast<std::size_t>(index[0]));
+  });
+  queue.wait_and_throw();
+}
+#else
+template <typename Func>
+[[noreturn]] inline void impl_sycl(std::size_t, Func&&, const Options&) {
+  throw std::runtime_error(
+      "SYCL parallel backend was requested, but this target was not built with "
+      "ITLABAI_HAS_SYCL");
+}
+#endif
+
 }  // namespace parallel
 }  // namespace it_lab_ai

include/parallel/parallel.hpp

@@ -30,7 +30,7 @@ inline Backend select_backend(const Options& opt, std::size_t n) {
 
   if (opt.backend == Backend::kSeq || opt.backend == Backend::kThreads ||
       opt.backend == Backend::kTbb || opt.backend == Backend::kOmp ||
-      opt.backend == Backend::kKokkos) {
+      opt.backend == Backend::kKokkos || opt.backend == Backend::kSycl) {
     return opt.backend;
   }
 
@@ -62,6 +62,9 @@ inline void parallel_for(std::size_t count, Func&& func,
     case Backend::kKokkos:
       impl_kokkos(count, std::forward<Func>(func), opt);
       break;
+    case Backend::kSycl:
+      impl_sycl(count, std::forward<Func>(func), opt);
+      break;
   }
 }