Skip to content

Commit ea8f22c

Browse files
breyermlbreyerml
authored
Merge pull request #77 from SC-SGS/SYCL_invocation
2 parents 02b60cf + b28ef53 commit ea8f22c

46 files changed

Lines changed: 6444 additions & 518 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎CMakeLists.txt‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -903,11 +903,12 @@ if (TARGET ${PLSSVM_SYCL_BACKEND_LIBRARY_NAME})
903903
choose the SYCL implementation to be used in the SYCL backend: ${PLSSVM_SYCL_BACKEND_NAME_LIST} (default: automatic)
904904
"
905905
)
906+
string(REPLACE ";" "|" PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST "${PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST}")
906907
set(PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_MANPAGE_ENTRY
907908
"
908909
.TP
909910
.B --sycl_kernel_invocation_type
910-
choose the kernel invocation type when using SYCL as backend: automatic|nd_range (default: automatic)
911+
choose the kernel invocation type when using SYCL as backend: ${PLSSVM_SYCL_KERNEL_INVOCATION_TYPE_NAME_LIST} (default: automatic)
911912
"
912913
)
913914
endif ()
@@ -936,8 +937,6 @@ endif ()
936937

937938
# configure the manpage files
938939
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-train.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-train.1 @ONLY)
939-
# update manpage entry since plssvm-predict can't recognize the SYCL kernel invocation type
940-
set(PLSSVM_SYCL_MANPAGE_ENTRY "${PLSSVM_SYCL_IMPLEMENTATION_TYPE_MANPAGE_ENTRY}")
941940
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-predict.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-predict.1 @ONLY)
942941
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-scale.1.in ${CMAKE_CURRENT_SOURCE_DIR}/docs/plssvm-scale.1 @ONLY)
943942

‎README.md‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -346,6 +346,8 @@ If the SYCL backend is available, additional options can be set.
346346
- `AUTO`: check for DPC++/icpx as implementation for the SYCL backend but **do not** fail if not available
347347
- `OFF`: do not check for DPC++/icpx as implementation for the SYCL backend
348348

349+
- `PLSSVM_ENABLE_SYCL_HIERARCHICAL_AND_SCOPED_KERNELS` (default: `ON`): enable SYCL's `hierarchical` and AdaptiveCpp's `scoped` kernel invocation types
350+
349351
To use DPC++/icpx for SYCL, simply set the `CMAKE_CXX_COMPILER` to the respective DPC++/icpx clang executable during CMake invocation.
350352

351353
If the SYCL implementation is DPC++/icpx the following additional options are available:
@@ -684,7 +686,7 @@ Usage:
684686
-b, --backend arg choose the backend: automatic|openmp|hpx|cuda|hip|opencl|sycl|kokkos|stdpar (default: automatic)
685687
-p, --target_platform arg choose the target platform: automatic|cpu|gpu_nvidia|gpu_amd|gpu_intel (default: automatic)
686688
--sycl_kernel_invocation_type arg
687-
choose the kernel invocation type when using SYCL as backend: automatic|nd_range (default: automatic)
689+
choose the kernel invocation type when using SYCL as backend: automatic|basic|work_group|hierarchical|scoped (default: automatic)
688690
--sycl_implementation_type arg
689691
choose the SYCL implementation to be used in the SYCL backend: automatic|dpcpp|adaptivecpp (default: automatic)
690692
--kokkos_execution_space arg
@@ -745,7 +747,7 @@ The `--target_platform=automatic` option works for the different backends as fol
745747
- `stdpar`: target device must be selected at compile time (using `PLSSVM_TARGET_PLATFORMS`) or using environment variables at runtime
746748

747749
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`.
748-
If the `--sycl_kernel_invocation_type` is `automatic`, the `nd_range` invocation type is currently always used.
750+
If the `--sycl_kernel_invocation_type` is `automatic`, the `work_group` invocation type is currently always used.
749751
If the `--sycl_implementation_type` is `automatic`, the used SYCL implementation is determined by the `PLSSVM_SYCL_BACKEND_PREFERRED_IMPLEMENTATION` CMake flag.
750752
If the `--kokkos_execution_space` is `automatic`, uses the best fitting execution space based on the provided and/or available target platforms.
751753

@@ -793,6 +795,8 @@ Usage:
793795
794796
-b, --backend arg choose the backend: automatic|openmp|hpx|cuda|hip|opencl|sycl|kokkos|stdpar (default: automatic)
795797
-p, --target_platform arg choose the target platform: automatic|cpu|gpu_nvidia|gpu_amd|gpu_intel (default: automatic)
798+
--sycl_kernel_invocation_type arg
799+
choose the kernel invocation type when using SYCL as backend: automatic|basic|work_group|hierarchical|scoped (default: automatic)
796800
--sycl_implementation_type arg
797801
choose the SYCL implementation to be used in the SYCL backend: automatic|dpcpp|adaptivecpp (default: automatic)
798802
--kokkos_execution_space arg

‎bindings/Python/README.md‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -332,10 +332,10 @@ The following table lists all PLSSVM enumerations exposed on the Python side:
332332

333333
If a SYCL implementation is available, additional enumerations are available:
334334

335-
| enumeration | values | description |
336-
|------------------------|-------------------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
337-
| `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. |
338-
| `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). |
335+
| enumeration | values | description |
336+
|------------------------|--------------------------------------------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
337+
| `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. |
338+
| `KernelInvocationType` | `AUTOMATIC`, `BASIC`, `WORK_GROUP`, `HIERARCHICAL`, `SCOPED` | The different supported SYCL kernel invocation types (default: `AUTOMATIC`). If `AUTOMATIC` is provided, simply uses `WORK_GROUP`. |
339339

340340
If the stdpar backend is available, an additional enumeration is available:
341341

‎bindings/Python/backends/sycl.cpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,10 @@ void init_sycl(py::module_ &m, const py::exception<plssvm::exception> &base_exce
4848
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.");
4949
py_enum_invocation
5050
.value("AUTOMATIC", plssvm::sycl::kernel_invocation_type::automatic, "use the best kernel invocation type for the current SYCL implementation and target hardware platform")
51-
.value("ND_RANGE", plssvm::sycl::kernel_invocation_type::nd_range, "use the nd_range kernel invocation type");
51+
.value("BASIC", plssvm::sycl::kernel_invocation_type::basic, "use the basic data parallel kernel invocation type")
52+
.value("WORK_GROUP", plssvm::sycl::kernel_invocation_type::work_group, "use the work-group data parallel kernel invocation type")
53+
.value("HIERARCHICAL", plssvm::sycl::kernel_invocation_type::hierarchical, "use the hierarchical data parallel kernel invocation type")
54+
.value("SCOPED", plssvm::sycl::kernel_invocation_type::scoped, "use the AdaptiveCpp specific scoped parallelism kernel invocation type");
5255

5356
// enable implicit conversion from string to enum
5457
plssvm::bindings::python::util::register_implicit_str_enum_conversion<plssvm::sycl::kernel_invocation_type>(py_enum_invocation);

‎docs/resources/dirs.dox‎

Lines changed: 144 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818
* @license This file is part of the PLSSVM project which is released under the MIT license.
1919
* See the LICENSE.md file in the project root for full license information.
2020
*
21-
* @brief Directory containing the implementation of all four available backends: OpenMP, CUDA, OpenCL, and SYCL.
21+
* @brief Directory containing the implementation of all available backends.
2222
*/
2323

2424
/**
@@ -488,6 +488,50 @@
488488
* @brief Directory containing kernel implementations for the explicit CG algorithm using the SYCL backend.
489489
*/
490490

491+
/**
492+
* @dir include/plssvm/backends/SYCL/kernel/cg_explicit/basic
493+
* @author Alexander Van Craen
494+
* @author Marcel Breyer
495+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
496+
* @license This file is part of the PLSSVM project which is released under the MIT license.
497+
* See the LICENSE.md file in the project root for full license information.
498+
*
499+
* @brief Directory containing basic data parallel kernel implementations for the explicit CG algorithm using the SYCL backend.
500+
*/
501+
502+
/**
503+
* @dir include/plssvm/backends/SYCL/kernel/cg_explicit/hierarchical
504+
* @author Alexander Van Craen
505+
* @author Marcel Breyer
506+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
507+
* @license This file is part of the PLSSVM project which is released under the MIT license.
508+
* See the LICENSE.md file in the project root for full license information.
509+
*
510+
* @brief Directory containing hierarchical kernel implementations for the explicit CG algorithm using the SYCL backend.
511+
*/
512+
513+
/**
514+
* @dir include/plssvm/backends/SYCL/kernel/cg_explicit/scoped
515+
* @author Alexander Van Craen
516+
* @author Marcel Breyer
517+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
518+
* @license This file is part of the PLSSVM project which is released under the MIT license.
519+
* See the LICENSE.md file in the project root for full license information.
520+
*
521+
* @brief Directory containing scoped-parallelism kernel implementations for the explicit CG algorithm using the SYCL backend.
522+
*/
523+
524+
/**
525+
* @dir include/plssvm/backends/SYCL/kernel/cg_explicit/work_group
526+
* @author Alexander Van Craen
527+
* @author Marcel Breyer
528+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
529+
* @license This file is part of the PLSSVM project which is released under the MIT license.
530+
* See the LICENSE.md file in the project root for full license information.
531+
*
532+
* @brief Directory containing work-group data parallel kernel implementations for the explicit CG algorithm using the SYCL backend.
533+
*/
534+
491535
/**
492536
* @dir include/plssvm/backends/SYCL/kernel/cg_implicit
493537
* @author Alexander Van Craen
@@ -499,6 +543,105 @@
499543
* @brief Directory containing kernel implementations for the implicit CG algorithm using the SYCL backend.
500544
*/
501545

546+
/**
547+
* @dir include/plssvm/backends/SYCL/kernel/cg_implicit/basic
548+
* @author Alexander Van Craen
549+
* @author Marcel Breyer
550+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
551+
* @license This file is part of the PLSSVM project which is released under the MIT license.
552+
* See the LICENSE.md file in the project root for full license information.
553+
*
554+
* @brief Directory containing basic data parallel kernel implementations for the implicit CG algorithm using the SYCL backend.
555+
*/
556+
557+
/**
558+
* @dir include/plssvm/backends/SYCL/kernel/cg_implicit/hierarchical
559+
* @author Alexander Van Craen
560+
* @author Marcel Breyer
561+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
562+
* @license This file is part of the PLSSVM project which is released under the MIT license.
563+
* See the LICENSE.md file in the project root for full license information.
564+
*
565+
* @brief Directory containing hierarchical kernel implementations for the implicit CG algorithm using the SYCL backend.
566+
*/
567+
568+
/**
569+
* @dir include/plssvm/backends/SYCL/kernel/cg_implicit/scoped
570+
* @author Alexander Van Craen
571+
* @author Marcel Breyer
572+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
573+
* @license This file is part of the PLSSVM project which is released under the MIT license.
574+
* See the LICENSE.md file in the project root for full license information.
575+
*
576+
* @brief Directory containing scoped-parallelism kernel implementations for the implicit CG algorithm using the SYCL backend.
577+
*/
578+
579+
/**
580+
* @dir include/plssvm/backends/SYCL/kernel/cg_implicit/work_group
581+
* @author Alexander Van Craen
582+
* @author Marcel Breyer
583+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
584+
* @license This file is part of the PLSSVM project which is released under the MIT license.
585+
* See the LICENSE.md file in the project root for full license information.
586+
*
587+
* @brief Directory containing work-group data parallel kernel implementations for the implicit CG algorithm using the SYCL backend.
588+
*/
589+
590+
/**
591+
* @dir include/plssvm/backends/SYCL/kernel/predict
592+
* @author Alexander Van Craen
593+
* @author Marcel Breyer
594+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
595+
* @license This file is part of the PLSSVM project which is released under the MIT license.
596+
* See the LICENSE.md file in the project root for full license information.
597+
*
598+
* @brief Directory containing kernel implementations for the predictions using the SYCL backend.
599+
*/
600+
601+
/**
602+
* @dir include/plssvm/backends/SYCL/kernel/predict/basic
603+
* @author Alexander Van Craen
604+
* @author Marcel Breyer
605+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
606+
* @license This file is part of the PLSSVM project which is released under the MIT license.
607+
* See the LICENSE.md file in the project root for full license information.
608+
*
609+
* @brief Directory containing basic data parallel kernel implementations for the predictions using the SYCL backend.
610+
*/
611+
612+
/**
613+
* @dir include/plssvm/backends/SYCL/kernel/predict/hierarchical
614+
* @author Alexander Van Craen
615+
* @author Marcel Breyer
616+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
617+
* @license This file is part of the PLSSVM project which is released under the MIT license.
618+
* See the LICENSE.md file in the project root for full license information.
619+
*
620+
* @brief Directory containing hierarchical kernel implementations for the predictions using the SYCL backend.
621+
*/
622+
623+
/**
624+
* @dir include/plssvm/backends/SYCL/kernel/predict/scoped
625+
* @author Alexander Van Craen
626+
* @author Marcel Breyer
627+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
628+
* @license This file is part of the PLSSVM project which is released under the MIT license.
629+
* See the LICENSE.md file in the project root for full license information.
630+
*
631+
* @brief Directory containing scoped-parallelism kernel implementations for the predictions using the SYCL backend.
632+
*/
633+
634+
/**
635+
* @dir include/plssvm/backends/SYCL/kernel/predict/work_group
636+
* @author Alexander Van Craen
637+
* @author Marcel Breyer
638+
* @copyright 2018-today The PLSSVM project - All Rights Reserved
639+
* @license This file is part of the PLSSVM project which is released under the MIT license.
640+
* See the LICENSE.md file in the project root for full license information.
641+
*
642+
* @brief Directory containing work-group data parallel kernel implementations for the predictions using the SYCL backend.
643+
*/
644+
502645
/**
503646
* @dir include/plssvm/backends/SYCL/DPCPP
504647
* @author Alexander Van Craen

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

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,7 @@
2323
#include "plssvm/detail/igor_utility.hpp" // plssvm::detail::get_value_from_named_parameter
2424
#include "plssvm/detail/memory_size.hpp" // plssvm::detail::memory_size
2525
#include "plssvm/detail/type_traits.hpp" // PLSSVM_REQUIRES, plssvm::detail::is_one_type_of
26+
#include "plssvm/exceptions/exceptions.hpp" // plssvm::invalid_parameter_exception
2627
#include "plssvm/mpi/communicator.hpp" // plssvm::mpi::communicator
2728
#include "plssvm/parameter.hpp" // plssvm::parameter, plssvm::detail::{has_only_sycl_parameter_named_args_v, has_only_sycl_named_args_v}
2829
#include "plssvm/svm/csvc.hpp" // plssvm::csvc
@@ -76,6 +77,14 @@ class csvm : public ::plssvm::detail::gpu_csvm<detail::device_ptr, detail::queue
7677
if constexpr (parser.has(sycl_kernel_invocation_type)) {
7778
// compile time check: the value must have the correct type
7879
invocation_type_ = ::plssvm::detail::get_value_from_named_parameter<sycl::kernel_invocation_type>(parser, sycl_kernel_invocation_type);
80+
81+
#if !defined(PLSSVM_SYCL_HIERARCHICAL_AND_SCOPED_KERNELS_ENABLED)
82+
if (invocation_type_ == sycl::kernel_invocation_type::hierarchical) {
83+
throw ::plssvm::invalid_parameter_exception{ "The provided sycl::kernel_invocation_type::hierarchical is disabled for the AdaptiveCpp SYCL backend!" };
84+
} else if (invocation_type_ == sycl::kernel_invocation_type::scoped) {
85+
throw ::plssvm::invalid_parameter_exception{ "he provided sycl::kernel_invocation_type::scoped is disabled for the AdaptiveCpp SYCL backend!" };
86+
}
87+
#endif
7988
}
8089
this->init(target);
8190
}

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

Lines changed: 27 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,9 +15,11 @@
1515

1616
#include "plssvm/backends/execution_range.hpp" // plssvm::detail::dim_type
1717
#include "plssvm/backends/SYCL/AdaptiveCpp/detail/queue.hpp" // plssvm::adaptivecpp::detail::queue (PImpl)
18+
#include "plssvm/backends/SYCL/kernel_invocation_types.hpp" // plssvm::sycl::kernel_invocation_type
19+
#include "plssvm/detail/utility.hpp" // plssvm::detail::unreachable
1820
#include "plssvm/target_platforms.hpp" // plssvm::target_platform
1921

20-
#include "sycl/sycl.hpp" // sycl::range
22+
#include "sycl/sycl.hpp" // sycl::range, sycl::nd_range
2123

2224
#include <string> // std::string
2325
#include <utility> // std::pair
@@ -46,6 +48,30 @@ template <std::size_t I>
4648
}
4749
}
4850

51+
/**
52+
* @brief Convert the provided @p grid and @p block to the final SYCL execution range.
53+
* @tparam invocation_type the SYCL kernel invocation type
54+
* @param[in] grid the execution grid
55+
* @param[in] block the execution block
56+
* @return the SYCL native execution range
57+
*/
58+
template <sycl::kernel_invocation_type invocation_type>
59+
auto get_execution_range(const ::plssvm::detail::dim_type &grid, const ::plssvm::detail::dim_type &block) {
60+
const ::sycl::range native_grid = detail::dim_type_to_native<2>(grid);
61+
const ::sycl::range native_block = detail::dim_type_to_native<2>(block);
62+
63+
if constexpr (invocation_type == sycl::kernel_invocation_type::basic) {
64+
return ::sycl::range<2>{ native_grid * native_block };
65+
} else if constexpr (invocation_type == sycl::kernel_invocation_type::work_group) {
66+
return ::sycl::nd_range<2>{ native_grid * native_block, native_block };
67+
} else if constexpr (invocation_type == sycl::kernel_invocation_type::hierarchical || invocation_type == sycl::kernel_invocation_type::scoped) {
68+
return ::sycl::nd_range<2>{ native_grid, native_block };
69+
} else {
70+
// can't be reached
71+
::plssvm::detail::unreachable();
72+
}
73+
}
74+
4975
/**
5076
* @brief Returns the list devices matching the target platform @p target and the actually used target platform
5177
* (only interesting if the provided @p target was automatic).

0 commit comments

Comments
 (0)