[SYCL] Throw exception if range/offset of kernel execution exceeds INT_MAX

sycl/include/CL/sycl/handler.hpp

@@ -25,6 +25,7 @@
 
 #include <algorithm>
 #include <functional>
+#include <limits>
 #include <memory>
 #include <type_traits>
 
@@ -128,6 +129,39 @@ struct check_fn_signature<F, RetT(Args...)> {
 
 __SYCL_EXPORT device getDeviceFromHandler(handler &);
 
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__)
+template <typename T> struct NotIntMsg;
+
+// TODO reword for "`fsycl-id-queries-fit-in-int' optimization flag." when
+// implemented
+template <int Dims> struct NotIntMsg<range<Dims>> {
+  constexpr static char *Msg = "Provided range is out of integer limits. "
+                               "Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to "
+                               "disable range check.";
+};
+
+template <int Dims> struct NotIntMsg<id<Dims>> {
+  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 <int Dims, typename T>
+typename std::enable_if<std::is_same<T, range<Dims>>::value ||
+                        std::is_same<T, id<Dims>>::value>::type
+checkValueRange(const T &V) {
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__)
+  static constexpr size_t Limit =
+      static_cast<size_t>((std::numeric_limits<int>::max)());
+  for (size_t Dim = 0; Dim < Dims; ++Dim)
+    if (V[Dim] > Limit)
+      throw runtime_error(NotIntMsg<T>::Msg, PI_INVALID_VALUE);
+#else
+  (void)V;
+#endif
+}
+
 } // namespace detail
 
 namespace intel {
@@ -764,6 +798,8 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     kernel_single_task<NameT>(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<NameT, KernelType, /*Dims*/ 0, void>(KernelFunc);
@@ -792,6 +828,7 @@ class __SYCL_EXPORT handler {
     (void)NumWorkItems;
     kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkItems);
     MNDRDesc.set(std::move(NumWorkItems));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;
@@ -804,6 +841,8 @@ class __SYCL_EXPORT handler {
   /// named function object type.
   template <typename FuncT> 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);
@@ -850,6 +889,8 @@ class __SYCL_EXPORT handler {
     (void)WorkItemOffset;
     kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkItems);
+    detail::checkValueRange<Dims>(WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;
@@ -878,6 +919,9 @@ class __SYCL_EXPORT handler {
     (void)ExecutionRange;
     kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(ExecutionRange.get_global_range());
+    detail::checkValueRange<Dims>(ExecutionRange.get_local_range());
+    detail::checkValueRange<Dims>(ExecutionRange.get_offset());
     MNDRDesc.set(std::move(ExecutionRange));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;
@@ -1047,6 +1091,7 @@ class __SYCL_EXPORT handler {
     (void)NumWorkGroups;
     kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;
@@ -1078,7 +1123,12 @@ class __SYCL_EXPORT handler {
     (void)WorkGroupSize;
     kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
 #else
-    MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
+    nd_range<Dims> ExecRange =
+        nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
+    detail::checkValueRange<Dims>(ExecRange.get_global_range());
+    detail::checkValueRange<Dims>(ExecRange.get_local_range());
+    detail::checkValueRange<Dims>(ExecRange.get_offset());
+    MNDRDesc.set(std::move(ExecRange));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;
 #endif // __SYCL_DEVICE_ONLY__
@@ -1093,6 +1143,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;
@@ -1111,6 +1163,7 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     verifyKernelInvoc(Kenrel);
     MKernel = detail::getSyclObjImpl(std::move(Kenrel));
+    detail::checkValueRange<Dims>(NumWorkItems);
     MNDRDesc.set(std::move(NumWorkItems));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
@@ -1130,6 +1183,8 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     verifyKernelInvoc(Kernel);
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    detail::checkValueRange<Dims>(NumWorkItems);
+    detail::checkValueRange<Dims>(WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
@@ -1147,6 +1202,9 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     verifyKernelInvoc(Kernel);
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    detail::checkValueRange<Dims>(NDRange.get_global_range());
+    detail::checkValueRange<Dims>(NDRange.get_local_range());
+    detail::checkValueRange<Dims>(NDRange.get_offset());
     MNDRDesc.set(std::move(NDRange));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
@@ -1167,6 +1225,8 @@ class __SYCL_EXPORT handler {
     (void)Kernel;
     kernel_single_task<NameT>(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;
@@ -1205,6 +1265,7 @@ class __SYCL_EXPORT handler {
     (void)NumWorkItems;
     kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkItems);
     MNDRDesc.set(std::move(NumWorkItems));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
@@ -1237,6 +1298,8 @@ class __SYCL_EXPORT handler {
     (void)WorkItemOffset;
     kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkItems);
+    detail::checkValueRange<Dims>(WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
@@ -1268,6 +1331,9 @@ class __SYCL_EXPORT handler {
     (void)NDRange;
     kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NDRange.get_global_range());
+    detail::checkValueRange<Dims>(NDRange.get_local_range());
+    detail::checkValueRange<Dims>(NDRange.get_offset());
     MNDRDesc.set(std::move(NDRange));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
@@ -1303,6 +1369,7 @@ class __SYCL_EXPORT handler {
     (void)NumWorkGroups;
     kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
 #else
+    detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
@@ -1339,7 +1406,12 @@ class __SYCL_EXPORT handler {
     (void)WorkGroupSize;
     kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
 #else
-    MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
+    nd_range<Dims> ExecRange =
+        nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
+    detail::checkValueRange<Dims>(ExecRange.get_global_range());
+    detail::checkValueRange<Dims>(ExecRange.get_local_range());
+    detail::checkValueRange<Dims>(ExecRange.get_offset());
+    MNDRDesc.set(std::move(ExecRange));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
     MCGType = detail::CG::KERNEL;

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 <CL/sycl.hpp>
+#include <climits>
+
+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<size_t>(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<int, 1> Buf{&Data, 1};
+
+  try {
+    Queue.submit([&](S::handler &CGH) {
+      auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ROL>(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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_RIL>(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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ROL_OIL>(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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_RIL_OOL>(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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_RIL_OIL>(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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ND_GOL_LIL_OIL>(
+          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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ND_GIL_LOL_OIL>(
+          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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ND_GIL_LIL_OOL>(
+          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<cl::sycl::access::mode::read_write>(CGH);
+
+      CGH.parallel_for<class PF_ND_GIL_LIL_OIL>(
+          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;
+}