From 08f8656280738c9e59b623ea51764d07fd844f9a Mon Sep 17 00:00:00 2001 From: sergei <57672082+s-kanaev@users.noreply.github.com> Date: Thu, 28 May 2020 12:22:13 +0300 Subject: [PATCH] [SYCL] Throw exception if range/offset of kernel execution exceeds INT_MAX (#1713) The exception will only be thrown if `__SYCL_ID_QUERIES_FIT_IN_INT__` macro is defined. Signed-off-by: Sergey Kanaev --- sycl/include/CL/sycl/handler.hpp | 76 +++++++- .../basic_tests/range_offset_fit_in_int.cpp | 177 ++++++++++++++++++ 2 files changed, 251 insertions(+), 2 deletions(-) create mode 100644 sycl/test/basic_tests/range_offset_fit_in_int.cpp diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index ddcdb1fe08188..7642fc5dc6f33 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -128,6 +129,39 @@ struct check_fn_signature { __SYCL_EXPORT device getDeviceFromHandler(handler &); +#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) +template struct NotIntMsg; + +// TODO reword for "`fsycl-id-queries-fit-in-int' optimization flag." when +// implemented +template struct NotIntMsg> { + constexpr static char *Msg = "Provided range is out of integer limits. " + "Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to " + "disable range check."; +}; + +template struct NotIntMsg> { + constexpr static char *Msg = "Provided offset is out of integer limits. " + "Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to " + "disable offset check."; +}; +#endif + +template +typename std::enable_if>::value || + std::is_same>::value>::type +checkValueRange(const T &V) { +#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) + static constexpr size_t Limit = + static_cast((std::numeric_limits::max)()); + for (size_t Dim = 0; Dim < Dims; ++Dim) + if (V[Dim] > Limit) + throw runtime_error(NotIntMsg::Msg, PI_INVALID_VALUE); +#else + (void)V; +#endif +} + } // namespace detail namespace intel { @@ -770,6 +804,8 @@ class __SYCL_EXPORT handler { #ifdef __SYCL_DEVICE_ONLY__ kernel_single_task(KernelFunc); #else + // No need to check if range is out of INT_MAX limits as it's compile-time + // known constant. MNDRDesc.set(range<1>{1}); StoreLambda(KernelFunc); @@ -798,6 +834,7 @@ class __SYCL_EXPORT handler { (void)NumWorkItems; kernel_parallel_for(KernelFunc); #else + detail::checkValueRange(NumWorkItems); MNDRDesc.set(std::move(NumWorkItems)); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; @@ -810,6 +847,8 @@ class __SYCL_EXPORT handler { /// named function object type. template void run_on_host_intel(FuncT Func) { throwIfActionIsCreated(); + // No need to check if range is out of INT_MAX limits as it's compile-time + // known constant MNDRDesc.set(range<1>{1}); MArgs = std::move(MAssociatedAccesors); @@ -856,6 +895,8 @@ class __SYCL_EXPORT handler { (void)WorkItemOffset; kernel_parallel_for(KernelFunc); #else + detail::checkValueRange(NumWorkItems); + detail::checkValueRange(WorkItemOffset); MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset)); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; @@ -884,6 +925,9 @@ class __SYCL_EXPORT handler { (void)ExecutionRange; kernel_parallel_for_nd_range(KernelFunc); #else + detail::checkValueRange(ExecutionRange.get_global_range()); + detail::checkValueRange(ExecutionRange.get_local_range()); + detail::checkValueRange(ExecutionRange.get_offset()); MNDRDesc.set(std::move(ExecutionRange)); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; @@ -1053,6 +1097,7 @@ class __SYCL_EXPORT handler { (void)NumWorkGroups; kernel_parallel_for_work_group(KernelFunc); #else + detail::checkValueRange(NumWorkGroups); MNDRDesc.setNumWorkGroups(NumWorkGroups); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; @@ -1084,7 +1129,12 @@ class __SYCL_EXPORT handler { (void)WorkGroupSize; kernel_parallel_for_work_group(KernelFunc); #else - MNDRDesc.set(nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize)); + nd_range ExecRange = + nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize); + detail::checkValueRange(ExecRange.get_global_range()); + detail::checkValueRange(ExecRange.get_local_range()); + detail::checkValueRange(ExecRange.get_offset()); + MNDRDesc.set(std::move(ExecRange)); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; #endif // __SYCL_DEVICE_ONLY__ @@ -1099,6 +1149,8 @@ class __SYCL_EXPORT handler { void single_task(kernel Kernel) { throwIfActionIsCreated(); verifyKernelInvoc(Kernel); + // No need to check if range is out of INT_MAX limits as it's compile-time + // known constant MNDRDesc.set(range<1>{1}); MKernel = detail::getSyclObjImpl(std::move(Kernel)); MCGType = detail::CG::KERNEL; @@ -1117,6 +1169,7 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); verifyKernelInvoc(Kenrel); MKernel = detail::getSyclObjImpl(std::move(Kenrel)); + detail::checkValueRange(NumWorkItems); MNDRDesc.set(std::move(NumWorkItems)); MCGType = detail::CG::KERNEL; extractArgsAndReqs(); @@ -1136,6 +1189,8 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); verifyKernelInvoc(Kernel); MKernel = detail::getSyclObjImpl(std::move(Kernel)); + detail::checkValueRange(NumWorkItems); + detail::checkValueRange(WorkItemOffset); MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset)); MCGType = detail::CG::KERNEL; extractArgsAndReqs(); @@ -1153,6 +1208,9 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); verifyKernelInvoc(Kernel); MKernel = detail::getSyclObjImpl(std::move(Kernel)); + detail::checkValueRange(NDRange.get_global_range()); + detail::checkValueRange(NDRange.get_local_range()); + detail::checkValueRange(NDRange.get_offset()); MNDRDesc.set(std::move(NDRange)); MCGType = detail::CG::KERNEL; extractArgsAndReqs(); @@ -1173,6 +1231,8 @@ class __SYCL_EXPORT handler { (void)Kernel; kernel_single_task(KernelFunc); #else + // No need to check if range is out of INT_MAX limits as it's compile-time + // known constant MNDRDesc.set(range<1>{1}); MKernel = detail::getSyclObjImpl(std::move(Kernel)); MCGType = detail::CG::KERNEL; @@ -1211,6 +1271,7 @@ class __SYCL_EXPORT handler { (void)NumWorkItems; kernel_parallel_for(KernelFunc); #else + detail::checkValueRange(NumWorkItems); MNDRDesc.set(std::move(NumWorkItems)); MKernel = detail::getSyclObjImpl(std::move(Kernel)); MCGType = detail::CG::KERNEL; @@ -1243,6 +1304,8 @@ class __SYCL_EXPORT handler { (void)WorkItemOffset; kernel_parallel_for(KernelFunc); #else + detail::checkValueRange(NumWorkItems); + detail::checkValueRange(WorkItemOffset); MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset)); MKernel = detail::getSyclObjImpl(std::move(Kernel)); MCGType = detail::CG::KERNEL; @@ -1274,6 +1337,9 @@ class __SYCL_EXPORT handler { (void)NDRange; kernel_parallel_for_nd_range(KernelFunc); #else + detail::checkValueRange(NDRange.get_global_range()); + detail::checkValueRange(NDRange.get_local_range()); + detail::checkValueRange(NDRange.get_offset()); MNDRDesc.set(std::move(NDRange)); MKernel = detail::getSyclObjImpl(std::move(Kernel)); MCGType = detail::CG::KERNEL; @@ -1309,6 +1375,7 @@ class __SYCL_EXPORT handler { (void)NumWorkGroups; kernel_parallel_for_work_group(KernelFunc); #else + detail::checkValueRange(NumWorkGroups); MNDRDesc.setNumWorkGroups(NumWorkGroups); MKernel = detail::getSyclObjImpl(std::move(Kernel)); StoreLambda(std::move(KernelFunc)); @@ -1345,7 +1412,12 @@ class __SYCL_EXPORT handler { (void)WorkGroupSize; kernel_parallel_for_work_group(KernelFunc); #else - MNDRDesc.set(nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize)); + nd_range ExecRange = + nd_range(NumWorkGroups * WorkGroupSize, WorkGroupSize); + detail::checkValueRange(ExecRange.get_global_range()); + detail::checkValueRange(ExecRange.get_local_range()); + detail::checkValueRange(ExecRange.get_offset()); + MNDRDesc.set(std::move(ExecRange)); MKernel = detail::getSyclObjImpl(std::move(Kernel)); StoreLambda(std::move(KernelFunc)); MCGType = detail::CG::KERNEL; diff --git a/sycl/test/basic_tests/range_offset_fit_in_int.cpp b/sycl/test/basic_tests/range_offset_fit_in_int.cpp new file mode 100644 index 0000000000000..d57446928983a --- /dev/null +++ b/sycl/test/basic_tests/range_offset_fit_in_int.cpp @@ -0,0 +1,177 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_ID_QUERIES_FIT_IN_INT__=1 %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out + +#include +#include + +namespace S = cl::sycl; + +void checkRangeException(S::runtime_error &E) { + constexpr char Msg[] = "Provided range is out of integer limits. " + "Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to " + "disable range check."; + + std::cerr << E.what() << std::endl; + + assert(std::string(E.what()).find(Msg) == 0 && "Unexpected message"); +} + +void checkOffsetException(S::runtime_error &E) { + constexpr char Msg[] = "Provided offset is out of integer limits. " + "Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to " + "disable offset check."; + + std::cerr << E.what() << std::endl; + + assert(std::string(E.what()).find(Msg) == 0 && "Unexpected message"); +} + +void test() { + auto EH = [](S::exception_list EL) { + for (const std::exception_ptr &E : EL) { + throw E; + } + }; + + S::queue Queue(EH); + + static constexpr size_t OutOfLimitsSize = static_cast(INT_MAX) + 1; + + S::range<1> RangeOutOfLimits{OutOfLimitsSize}; + S::range<1> RangeInLimits{1}; + S::id<1> OffsetOutOfLimits{OutOfLimitsSize}; + S::id<1> OffsetInLimits{1}; + S::nd_range<1> NDRange_ROL_LIL_OIL{RangeOutOfLimits, RangeInLimits, + OffsetInLimits}; + S::nd_range<1> NDRange_RIL_LOL_OIL{RangeInLimits, RangeOutOfLimits, + OffsetInLimits}; + S::nd_range<1> NDRange_RIL_LIL_OOL{RangeInLimits, RangeInLimits, + OffsetOutOfLimits}; + S::nd_range<1> NDRange_RIL_LIL_OIL(RangeInLimits, RangeInLimits, + OffsetInLimits); + + int Data = 0; + S::buffer Buf{&Data, 1}; + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for(RangeOutOfLimits, + [=](S::id<1> Id) { Acc[0] += 1; }); + }); + + assert(false && "Exception expected"); + } catch (S::runtime_error &E) { + checkRangeException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for(RangeInLimits, + [Acc](S::id<1> Id) { Acc[0] += 1; }); + }); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for(RangeOutOfLimits, OffsetInLimits, + [Acc](S::id<1> Id) { Acc[0] += 1; }); + }); + + assert(false && "Exception expected"); + } catch (S::runtime_error &E) { + checkRangeException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for(RangeInLimits, OffsetOutOfLimits, + [Acc](S::id<1> Id) { Acc[0] += 1; }); + }); + + assert(false && "Exception expected"); + } catch (S::runtime_error &E) { + checkOffsetException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for(RangeInLimits, OffsetInLimits, + [Acc](S::id<1> Id) { Acc[0] += 1; }); + }); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for( + NDRange_ROL_LIL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; }); + }); + } catch (S::runtime_error &E) { + checkRangeException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for( + NDRange_RIL_LOL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; }); + }); + } catch (S::runtime_error &E) { + checkRangeException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for( + NDRange_RIL_LIL_OOL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; }); + }); + } catch (S::runtime_error &E) { + checkOffsetException(E); + } catch (...) { + assert(false && "Unexpected exception catched"); + } + + try { + Queue.submit([&](S::handler &CGH) { + auto Acc = Buf.get_access(CGH); + + CGH.parallel_for( + NDRange_RIL_LIL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; }); + }); + } catch (...) { + assert(false && "Unexpected exception catched"); + } +} + +int main(void) { + test(); + return 0; +}