From 3917d113cd420ea0be25640263f24dd2a5b9c85d Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 11 Oct 2024 08:47:00 -0700 Subject: [PATCH 1/2] [SYCL][NFC] Optmize handler.hpp compilation for device [1/N] Outlined runtime checks related to `-fsycl-id-queries-fit-in-int` into a separate header file which is only used in host compilation. --- .../sycl/detail/id_queries_fit_in_int.hpp | 112 ++++++++++++++++++ sycl/include/sycl/handler.hpp | 102 +++------------- .../include_deps/sycl_detail_core.hpp.cpp | 1 + 3 files changed, 129 insertions(+), 86 deletions(-) create mode 100644 sycl/include/sycl/detail/id_queries_fit_in_int.hpp diff --git a/sycl/include/sycl/detail/id_queries_fit_in_int.hpp b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp new file mode 100644 index 0000000000000..2c77b44d5e7a2 --- /dev/null +++ b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp @@ -0,0 +1,112 @@ +//==-------------------- id_queries_fit_in_int.hpp -------------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// Our SYCL implementation has a special mode (introduced for performance +// reasons) in which it assume that all result of all id queries (i.e. global +// sizes, work-group sizes, local id, global id, etc.) fit within MAX_INT. +// +// This header contains corresponding helper functions related to this mode. +// +//===----------------------------------------------------------------------===// + +#pragma once + +// We only use those helpers to throw an exception if user selected a range that +// would violate the assumption. That can only happen on host and therefore to +// optimize our headers, the helpers below are only available for host +// compilation. +#ifndef __SYCL_DEVICE_ONLY__ + +#include + +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +#if __SYCL_ID_QUERIES_FIT_IN_INT__ +template struct NotIntMsg; + +template struct NotIntMsg> { + constexpr static const char *Msg = + "Provided range is out of integer limits. Pass " + "`-fno-sycl-id-queries-fit-in-int' to disable range check."; +}; + +template struct NotIntMsg> { + constexpr static const char *Msg = + "Provided offset is out of integer limits. Pass " + "`-fno-sycl-id-queries-fit-in-int' to disable offset check."; +}; +#endif + +#if __SYCL_ID_QUERIES_FIT_IN_INT__ +template +typename std::enable_if_t::value || + std::is_same::value> +checkValueRangeImpl(ValT V) { + static constexpr size_t Limit = + static_cast((std::numeric_limits::max)()); + if (V > Limit) + throw sycl::exception(make_error_code(errc::nd_range), NotIntMsg::Msg); +} +#endif + +template +typename std::enable_if_t> || + std::is_same_v>> +checkValueRange([[maybe_unused]] const T &V) { +#if __SYCL_ID_QUERIES_FIT_IN_INT__ + for (size_t Dim = 0; Dim < Dims; ++Dim) + checkValueRangeImpl(V[Dim]); + + { + unsigned long long Product = 1; + for (size_t Dim = 0; Dim < Dims; ++Dim) { + Product *= V[Dim]; + // check value now to prevent product overflow in the end + checkValueRangeImpl(Product); + } + } +#endif +} + +template +void checkValueRange([[maybe_unused]] const range &R, + [[maybe_unused]] const id &O) { +#if __SYCL_ID_QUERIES_FIT_IN_INT__ + checkValueRange(R); + checkValueRange(O); + + for (size_t Dim = 0; Dim < Dims; ++Dim) { + unsigned long long Sum = R[Dim] + O[Dim]; + + checkValueRangeImpl>(Sum); + } +#endif +} + +template +typename std::enable_if_t>> +checkValueRange([[maybe_unused]] const T &V) { +#if __SYCL_ID_QUERIES_FIT_IN_INT__ + checkValueRange(V.get_global_range()); + checkValueRange(V.get_local_range()); + checkValueRange(V.get_offset()); + + checkValueRange(V.get_global_range(), V.get_offset()); +#endif +} + +} // namespace detail +} // namespace _V1 +} // namespace sycl + +#endif diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 4f443a2103eb4..6aff9b978cf66 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include @@ -227,22 +228,6 @@ __SYCL_EXPORT void *getValueFromDynamicParameter( ext::oneapi::experimental::detail::dynamic_parameter_base &DynamicParamBase); -#if __SYCL_ID_QUERIES_FIT_IN_INT__ -template struct NotIntMsg; - -template struct NotIntMsg> { - constexpr static const char *Msg = - "Provided range is out of integer limits. Pass " - "`-fno-sycl-id-queries-fit-in-int' to disable range check."; -}; - -template struct NotIntMsg> { - constexpr static const char *Msg = - "Provided offset is out of integer limits. Pass " - "`-fno-sycl-id-queries-fit-in-int' to disable offset check."; -}; -#endif - // Helper for merging properties with ones defined in an optional kernel functor // getter. template @@ -265,70 +250,6 @@ struct GetMergedKernelProperties< PropertiesT, get_method_properties>; }; -#if __SYCL_ID_QUERIES_FIT_IN_INT__ -template -typename std::enable_if_t::value || - std::is_same::value> -checkValueRangeImpl(ValT V) { - static constexpr size_t Limit = - static_cast((std::numeric_limits::max)()); - if (V > Limit) - throw sycl::exception(make_error_code(errc::nd_range), NotIntMsg::Msg); -} -#endif - -template -typename std::enable_if_t> || - std::is_same_v>> -checkValueRange(const T &V) { -#if __SYCL_ID_QUERIES_FIT_IN_INT__ - for (size_t Dim = 0; Dim < Dims; ++Dim) - checkValueRangeImpl(V[Dim]); - - { - unsigned long long Product = 1; - for (size_t Dim = 0; Dim < Dims; ++Dim) { - Product *= V[Dim]; - // check value now to prevent product overflow in the end - checkValueRangeImpl(Product); - } - } -#else - (void)V; -#endif -} - -template -void checkValueRange(const range &R, const id &O) { -#if __SYCL_ID_QUERIES_FIT_IN_INT__ - checkValueRange(R); - checkValueRange(O); - - for (size_t Dim = 0; Dim < Dims; ++Dim) { - unsigned long long Sum = R[Dim] + O[Dim]; - - checkValueRangeImpl>(Sum); - } -#else - (void)R; - (void)O; -#endif -} - -template -typename std::enable_if_t>> -checkValueRange(const T &V) { -#if __SYCL_ID_QUERIES_FIT_IN_INT__ - checkValueRange(V.get_global_range()); - checkValueRange(V.get_local_range()); - checkValueRange(V.get_offset()); - - checkValueRange(V.get_global_range(), V.get_offset()); -#else - (void)V; -#endif -} - template class RoundedRangeIDGenerator { id Id; id InitId; @@ -1353,8 +1274,10 @@ class __SYCL_EXPORT handler { /// \param Kernel is a SYCL kernel function. /// \param Properties is the properties. template - void parallel_for_impl(range NumWorkItems, PropertiesT Props, - kernel Kernel) { + void parallel_for_impl([[maybe_unused]] range NumWorkItems, + [[maybe_unused]] PropertiesT Props, + [[maybe_unused]] kernel Kernel) { +#ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); MKernel = detail::getSyclObjImpl(std::move(Kernel)); detail::checkValueRange(NumWorkItems); @@ -1364,6 +1287,7 @@ class __SYCL_EXPORT handler { setNDRangeUsed(false); extractArgsAndReqs(); MKernelName = getKernelName(); +#endif } /// Defines and invokes a SYCL kernel function for the specified range and @@ -1376,8 +1300,10 @@ class __SYCL_EXPORT handler { /// \param Properties is the properties. /// \param Kernel is a SYCL kernel function. template - void parallel_for_impl(nd_range NDRange, PropertiesT Props, - kernel Kernel) { + void parallel_for_impl([[maybe_unused]] nd_range NDRange, + [[maybe_unused]] PropertiesT Props, + [[maybe_unused]] kernel Kernel) { +#ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); MKernel = detail::getSyclObjImpl(std::move(Kernel)); detail::checkValueRange(NDRange); @@ -1387,6 +1313,7 @@ class __SYCL_EXPORT handler { setNDRangeUsed(true); extractArgsAndReqs(); MKernelName = getKernelName(); +#endif } /// Hierarchical kernel invocation method of a kernel defined as a lambda @@ -2136,8 +2063,10 @@ class __SYCL_EXPORT handler { /// \param Kernel is a SYCL kernel function. template __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020") - void parallel_for(range NumWorkItems, id WorkItemOffset, - kernel Kernel) { + void parallel_for([[maybe_unused]] range NumWorkItems, + [[maybe_unused]] id WorkItemOffset, + [[maybe_unused]] kernel Kernel) { +#ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); MKernel = detail::getSyclObjImpl(std::move(Kernel)); detail::checkValueRange(NumWorkItems, WorkItemOffset); @@ -2146,6 +2075,7 @@ class __SYCL_EXPORT handler { setNDRangeUsed(false); extractArgsAndReqs(); MKernelName = getKernelName(); +#endif } /// Defines and invokes a SYCL kernel function for the specified range and diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index a9ca815bcfffd..ad859b47c6495 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -145,6 +145,7 @@ // CHECK-NEXT: CL/cl_version.h // CHECK-NEXT: CL/cl_platform.h // CHECK-NEXT: CL/cl_ext.h +// CHECK-NEXT: detail/id_queries_fit_in_int.hpp // CHECK-NEXT: detail/reduction_forward.hpp // CHECK-NEXT: detail/ur.hpp // CHECK-NEXT: ur_api_funcs.def From 3528cb4c4def57e91c7622f633ac18e46f08e78b Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 11 Oct 2024 10:23:19 -0700 Subject: [PATCH 2/2] Apply code review comments --- sycl/include/sycl/detail/id_queries_fit_in_int.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/include/sycl/detail/id_queries_fit_in_int.hpp b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp index 2c77b44d5e7a2..d3ce74dfdfc0a 100644 --- a/sycl/include/sycl/detail/id_queries_fit_in_int.hpp +++ b/sycl/include/sycl/detail/id_queries_fit_in_int.hpp @@ -45,9 +45,7 @@ template struct NotIntMsg> { "Provided offset is out of integer limits. Pass " "`-fno-sycl-id-queries-fit-in-int' to disable offset check."; }; -#endif -#if __SYCL_ID_QUERIES_FIT_IN_INT__ template typename std::enable_if_t::value || std::is_same::value>