Add support for the different SYCL kernel invocation types

CMakeLists.txt

@@ -903,11 +903,12 @@ if (TARGET ${PLSSVM_SYCL_BACKEND_LIBRARY_NAME})
 choose the SYCL implementation to be used in the SYCL backend: ${PLSSVM_SYCL_BACKEND_NAME_LIST} (default: automatic)
 "
     )
+    string(REPLACE ";" "|" PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST "${PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST}")
     set(PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_MANPAGE_ENTRY
         "
 .TP
 .B --sycl_kernel_invocation_type
-choose the kernel invocation type when using SYCL as backend: automatic|nd_range (default: automatic)
+choose the kernel invocation type when using SYCL as backend: ${PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST} (default: automatic)
 "
     )
 endif ()
@@ -936,8 +937,6 @@ endif ()
 
 # configure the manpage files
 configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-train.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-train.1 @ONLY)
-# update manpage entry since plssvm-predict can't recognize the SYCL kernel invocation type
-set(PLSSVM_SYCL_MANPAGE_ENTRY "${PLSSVM_SYCL_IMPLEMENTATION_TYPE_MANPAGE_ENTRY}")
 configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-predict.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-predict.1 @ONLY)
 configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-scale.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-scale.1 @ONLY)
 

README.md

@@ -346,6 +346,8 @@ If the SYCL backend is available, additional options can be set.
   - `AUTO`: check for DPC++/icpx as implementation for the SYCL backend but **do not** fail if not available
   - `OFF`: do not check for DPC++/icpx as implementation for the SYCL backend
 
+- `PLSSVM_ENABLE_SYCL_HIERARCHICAL_AND_SCOPED_KERNELS` (default: `ON`): enable SYCL's `hierarchical` and AdaptiveCpp's `scoped` kernel invocation types
+
 To use DPC++/icpx for SYCL, simply set the `CMAKE_CXX_COMPILER` to the respective DPC++/icpx clang executable during CMake invocation.
 
 If the SYCL implementation is DPC++/icpx the following additional options are available:
@@ -684,7 +686,7 @@ Usage:
   -b, --backend arg             choose the backend: automatic|openmp|hpx|cuda|hip|opencl|sycl|kokkos|stdpar (default: automatic)
   -p, --target_platform arg     choose the target platform: automatic|cpu|gpu_nvidia|gpu_amd|gpu_intel (default: automatic)
       --sycl_kernel_invocation_type arg
-                                choose the kernel invocation type when using SYCL as backend: automatic|nd_range (default: automatic)
+                                choose the kernel invocation type when using SYCL as backend: automatic|basic|work_group|hierarchical|scoped (default: automatic)
       --sycl_implementation_type arg
                                 choose the SYCL implementation to be used in the SYCL backend: automatic|dpcpp|adaptivecpp (default: automatic)
       --kokkos_execution_space arg
@@ -745,7 +747,7 @@ The `--target_platform=automatic` option works for the different backends as fol
 - `stdpar`: target device must be selected at compile time (using `PLSSVM_TARGET_PLATFORMS`) or using environment variables at runtime
 
 The `--sycl_kernel_invocation_type` and `--sycl_implementation_type` flags are only used if the `--backend` is `sycl`, otherwise a warning is emitted on `stderr`.
-If the `--sycl_kernel_invocation_type` is `automatic`, the `nd_range` invocation type is currently always used.
+If the `--sycl_kernel_invocation_type` is `automatic`, the `work_group` invocation type is currently always used.
 If the `--sycl_implementation_type` is `automatic`, the used SYCL implementation is determined by the `PLSSVM_SYCL_BACKEND_PREFERRED_IMPLEMENTATION` CMake flag.
 If the `--kokkos_execution_space` is `automatic`, uses the best fitting execution space based on the provided and/or available target platforms.
 
@@ -793,6 +795,8 @@ Usage:
 
   -b, --backend arg             choose the backend: automatic|openmp|hpx|cuda|hip|opencl|sycl|kokkos|stdpar (default: automatic)
   -p, --target_platform arg     choose the target platform: automatic|cpu|gpu_nvidia|gpu_amd|gpu_intel (default: automatic)
+      --sycl_kernel_invocation_type arg
+                                choose the kernel invocation type when using SYCL as backend: automatic|basic|work_group|hierarchical|scoped (default: automatic)
       --sycl_implementation_type arg
                                 choose the SYCL implementation to be used in the SYCL backend: automatic|dpcpp|adaptivecpp (default: automatic)
       --kokkos_execution_space arg

bindings/Python/README.md

@@ -332,10 +332,10 @@ The following table lists all PLSSVM enumerations exposed on the Python side:
 
 If a SYCL implementation is available, additional enumerations are available:
 
-| enumeration            | values                              | description                                                                                                                                                                                                                                               |
-|------------------------|-------------------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
-| `ImplementationType`   | `AUTOMATIC`, `DPCPP`, `ADAPTIVECPP` | The different supported SYCL implementation types (default: `AUTOMATIC`). If `AUTOMATIC` is provided, determines the used SYCL implementation based on the value of `-DPLSSVM_SYCL_BACKEND_PREFERRED_IMPLEMENTATION` provided during PLSSVM'S build step. |
-| `KernelInvocationType` | `AUTOMATIC`, `ND_RANGE`             | The different supported SYCL kernel invocation types (default: `AUTOMATIC`). If `AUTOMATIC` is provided, simply uses `ND_RANGE` (only implemented to be able to add new invocation types in the future).                                                  |
+| enumeration            | values                                                       | description                                                                                                                                                                                                                                               |
+|------------------------|--------------------------------------------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
+| `ImplementationType`   | `AUTOMATIC`, `DPCPP`, `ADAPTIVECPP`                          | The different supported SYCL implementation types (default: `AUTOMATIC`). If `AUTOMATIC` is provided, determines the used SYCL implementation based on the value of `-DPLSSVM_SYCL_BACKEND_PREFERRED_IMPLEMENTATION` provided during PLSSVM'S build step. |
+| `KernelInvocationType` | `AUTOMATIC`, `BASIC`, `WORK_GROUP`, `HIERARCHICAL`, `SCOPED` | The different supported SYCL kernel invocation types (default: `AUTOMATIC`). If `AUTOMATIC` is provided, simply uses `WORK_GROUP`.                                                                                                                        |
 
 If the stdpar backend is available, an additional enumeration is available:
 

bindings/Python/backends/sycl.cpp

@@ -48,7 +48,10 @@ void init_sycl(py::module_ &m, const py::exception<plssvm::exception> &base_exce
     py::enum_<plssvm::sycl::kernel_invocation_type> py_enum_invocation(sycl_module, "KernelInvocationType", "Enum class for all possible SYCL kernel invocation types supported in PLSSVM.");
     py_enum_invocation
         .value("AUTOMATIC", plssvm::sycl::kernel_invocation_type::automatic, "use the best kernel invocation type for the current SYCL implementation and target hardware platform")
-        .value("ND_RANGE", plssvm::sycl::kernel_invocation_type::nd_range, "use the nd_range kernel invocation type");
+        .value("BASIC", plssvm::sycl::kernel_invocation_type::basic, "use the basic data parallel kernel invocation type")
+        .value("WORK_GROUP", plssvm::sycl::kernel_invocation_type::work_group, "use the work-group data parallel kernel invocation type")
+        .value("HIERARCHICAL", plssvm::sycl::kernel_invocation_type::hierarchical, "use the hierarchical data parallel kernel invocation type")
+        .value("SCOPED", plssvm::sycl::kernel_invocation_type::scoped, "use the AdaptiveCpp specific scoped parallelism kernel invocation type");
 
     // enable implicit conversion from string to enum
     plssvm::bindings::python::util::register_implicit_str_enum_conversion<plssvm::sycl::kernel_invocation_type>(py_enum_invocation);

docs/resources/dirs.dox

@@ -18,7 +18,7 @@
  * @license This file is part of the PLSSVM project which is released under the MIT license.
  *          See the LICENSE.md file in the project root for full license information.
  *
- * @brief Directory containing the implementation of all four available backends: OpenMP, CUDA, OpenCL, and SYCL.
+ * @brief Directory containing the implementation of all available backends.
  */
 
 /**
@@ -488,6 +488,50 @@
  * @brief Directory containing kernel implementations for the explicit CG algorithm using the SYCL backend.
  */
 
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_explicit/basic
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing basic data parallel kernel implementations for the explicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_explicit/hierarchical
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing hierarchical kernel implementations for the explicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_explicit/scoped
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing scoped-parallelism kernel implementations for the explicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_explicit/work_group
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing work-group data parallel kernel implementations for the explicit CG algorithm using the SYCL backend.
+ */
+
 /**
  * @dir include/plssvm/backends/SYCL/kernel/cg_implicit
  * @author Alexander Van Craen
@@ -499,6 +543,105 @@
  * @brief Directory containing kernel implementations for the implicit CG algorithm using the SYCL backend.
  */
 
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_implicit/basic
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing basic data parallel kernel implementations for the implicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_implicit/hierarchical
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing hierarchical kernel implementations for the implicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_implicit/scoped
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing scoped-parallelism kernel implementations for the implicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/cg_implicit/work_group
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing work-group data parallel kernel implementations for the implicit CG algorithm using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/predict
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing kernel implementations for the predictions using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/predict/basic
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing basic data parallel kernel implementations for the predictions using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/predict/hierarchical
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing hierarchical kernel implementations for the predictions using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/predict/scoped
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing scoped-parallelism kernel implementations for the predictions using the SYCL backend.
+ */
+
+/**
+ * @dir include/plssvm/backends/SYCL/kernel/predict/work_group
+ * @author Alexander Van Craen
+ * @author Marcel Breyer
+ * @copyright 2018-today The PLSSVM project - All Rights Reserved
+ * @license This file is part of the PLSSVM project which is released under the MIT license.
+ *          See the LICENSE.md file in the project root for full license information.
+ *
+ * @brief Directory containing work-group data parallel kernel implementations for the predictions using the SYCL backend.
+ */
+
 /**
  * @dir include/plssvm/backends/SYCL/DPCPP
  * @author Alexander Van Craen

include/plssvm/backends/SYCL/AdaptiveCpp/csvm.hpp

@@ -23,6 +23,7 @@
 #include "plssvm/detail/igor_utility.hpp"                             // plssvm::detail::get_value_from_named_parameter
 #include "plssvm/detail/memory_size.hpp"                              // plssvm::detail::memory_size
 #include "plssvm/detail/type_traits.hpp"                              // PLSSVM_REQUIRES, plssvm::detail::is_one_type_of
+#include "plssvm/exceptions/exceptions.hpp"                           // plssvm::invalid_parameter_exception
 #include "plssvm/mpi/communicator.hpp"                                // plssvm::mpi::communicator
 #include "plssvm/parameter.hpp"                                       // plssvm::parameter, plssvm::detail::{has_only_sycl_parameter_named_args_v, has_only_sycl_named_args_v}
 #include "plssvm/svm/csvc.hpp"                                        // plssvm::csvc
@@ -76,6 +77,14 @@ class csvm : public ::plssvm::detail::gpu_csvm<detail::device_ptr, detail::queue
         if constexpr (parser.has(sycl_kernel_invocation_type)) {
             // compile time check: the value must have the correct type
             invocation_type_ = ::plssvm::detail::get_value_from_named_parameter<sycl::kernel_invocation_type>(parser, sycl_kernel_invocation_type);
+
+#if !defined(PLSSVM_SYCL_HIERARCHICAL_AND_SCOPED_KERNELS_ENABLED)
+            if (invocation_type_ == sycl::kernel_invocation_type::hierarchical) {
+                throw ::plssvm::invalid_parameter_exception{ "The provided sycl::kernel_invocation_type::hierarchical is disabled for the AdaptiveCpp SYCL backend!" };
+            } else if (invocation_type_ == sycl::kernel_invocation_type::scoped) {
+                throw ::plssvm::invalid_parameter_exception{ "he provided sycl::kernel_invocation_type::scoped is disabled for the AdaptiveCpp SYCL backend!" };
+            }
+#endif
         }
         this->init(target);
     }

include/plssvm/backends/SYCL/AdaptiveCpp/detail/utility.hpp

@@ -15,9 +15,11 @@
 
 #include "plssvm/backends/execution_range.hpp"                // plssvm::detail::dim_type
 #include "plssvm/backends/SYCL/AdaptiveCpp/detail/queue.hpp"  // plssvm::adaptivecpp::detail::queue (PImpl)
+#include "plssvm/backends/SYCL/kernel_invocation_types.hpp"   // plssvm::sycl::kernel_invocation_type
+#include "plssvm/detail/utility.hpp"                          // plssvm::detail::unreachable
 #include "plssvm/target_platforms.hpp"                        // plssvm::target_platform
 
-#include "sycl/sycl.hpp"  // sycl::range
+#include "sycl/sycl.hpp"  // sycl::range, sycl::nd_range
 
 #include <string>   // std::string
 #include <utility>  // std::pair
@@ -46,6 +48,30 @@ template <std::size_t I>
     }
 }
 
+/**
+ * @brief Convert the provided @p grid and @p block to the final SYCL execution range.
+ * @tparam invocation_type the SYCL kernel invocation type
+ * @param[in] grid the execution grid
+ * @param[in] block the execution block
+ * @return the SYCL native execution range
+ */
+template <sycl::kernel_invocation_type invocation_type>
+auto get_execution_range(const ::plssvm::detail::dim_type &grid, const ::plssvm::detail::dim_type &block) {
+    const ::sycl::range native_grid = detail::dim_type_to_native<2>(grid);
+    const ::sycl::range native_block = detail::dim_type_to_native<2>(block);
+
+    if constexpr (invocation_type == sycl::kernel_invocation_type::basic) {
+        return ::sycl::range<2>{ native_grid * native_block };
+    } else if constexpr (invocation_type == sycl::kernel_invocation_type::work_group) {
+        return ::sycl::nd_range<2>{ native_grid * native_block, native_block };
+    } else if constexpr (invocation_type == sycl::kernel_invocation_type::hierarchical || invocation_type == sycl::kernel_invocation_type::scoped) {
+        return ::sycl::nd_range<2>{ native_grid, native_block };
+    } else {
+        // can't be reached
+        ::plssvm::detail::unreachable();
+    }
+}
+
 /**
  * @brief Returns the list devices matching the target platform @p target and the actually used target platform
  *        (only interesting if the provided @p target was automatic).