Skip to content
72 changes: 70 additions & 2 deletions sycl/include/CL/sycl/handler.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -128,6 +128,36 @@ 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;

template <int Dims> struct NotIntMsg<range<Dims>> {
constexpr static char *Msg = "Provided range is out of integer limits. "
"Suggest disabling `id-queries-fit-in-int32' "
"optimizations flag.";
};

template <int Dims> struct NotIntMsg<id<Dims>> {
constexpr static char *Msg = "Provided offset is out of integer limits. "
"Suggest disabling `id-queries-fit-in-int32' "
"optimizations flag.";
};
#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
throwIfNotInt(const T &V) {
#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__)
static constexpr size_t Limit = static_cast<size_t>(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 {
Expand Down Expand Up @@ -764,6 +794,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);
Expand Down Expand Up @@ -792,6 +824,7 @@ class __SYCL_EXPORT handler {
(void)NumWorkItems;
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkItems);
MNDRDesc.set(std::move(NumWorkItems));
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
MCGType = detail::CG::KERNEL;
Expand All @@ -804,6 +837,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);
Expand Down Expand Up @@ -850,6 +885,8 @@ class __SYCL_EXPORT handler {
(void)WorkItemOffset;
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkItems);
detail::throwIfNotInt<Dims>(WorkItemOffset);
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -878,6 +915,9 @@ class __SYCL_EXPORT handler {
(void)ExecutionRange;
kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(ExecutionRange.get_global_range());
detail::throwIfNotInt<Dims>(ExecutionRange.get_local_range());
detail::throwIfNotInt<Dims>(ExecutionRange.get_offset());
MNDRDesc.set(std::move(ExecutionRange));
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -1047,6 +1087,7 @@ class __SYCL_EXPORT handler {
(void)NumWorkGroups;
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkGroups);
MNDRDesc.setNumWorkGroups(NumWorkGroups);
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -1078,7 +1119,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::throwIfNotInt<Dims>(ExecRange.get_global_range());
detail::throwIfNotInt<Dims>(ExecRange.get_local_range());
detail::throwIfNotInt<Dims>(ExecRange.get_offset());
MNDRDesc.set(std::move(ExecRange));
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
MCGType = detail::CG::KERNEL;
#endif // __SYCL_DEVICE_ONLY__
Expand All @@ -1093,6 +1139,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;
Expand All @@ -1111,6 +1159,7 @@ class __SYCL_EXPORT handler {
throwIfActionIsCreated();
verifyKernelInvoc(Kenrel);
MKernel = detail::getSyclObjImpl(std::move(Kenrel));
detail::throwIfNotInt<Dims>(NumWorkItems);
MNDRDesc.set(std::move(NumWorkItems));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
Expand All @@ -1130,6 +1179,8 @@ class __SYCL_EXPORT handler {
throwIfActionIsCreated();
verifyKernelInvoc(Kernel);
MKernel = detail::getSyclObjImpl(std::move(Kernel));
detail::throwIfNotInt<Dims>(NumWorkItems);
detail::throwIfNotInt<Dims>(WorkItemOffset);
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
Expand All @@ -1147,6 +1198,9 @@ class __SYCL_EXPORT handler {
throwIfActionIsCreated();
verifyKernelInvoc(Kernel);
MKernel = detail::getSyclObjImpl(std::move(Kernel));
detail::throwIfNotInt<Dims>(NDRange.get_global_range());
detail::throwIfNotInt<Dims>(NDRange.get_local_range());
detail::throwIfNotInt<Dims>(NDRange.get_offset());
MNDRDesc.set(std::move(NDRange));
MCGType = detail::CG::KERNEL;
extractArgsAndReqs();
Expand All @@ -1167,6 +1221,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;
Expand Down Expand Up @@ -1205,6 +1261,7 @@ class __SYCL_EXPORT handler {
(void)NumWorkItems;
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkItems);
MNDRDesc.set(std::move(NumWorkItems));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -1237,6 +1294,8 @@ class __SYCL_EXPORT handler {
(void)WorkItemOffset;
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkItems);
detail::throwIfNotInt<Dims>(WorkItemOffset);
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -1268,6 +1327,9 @@ class __SYCL_EXPORT handler {
(void)NDRange;
kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NDRange.get_global_range());
detail::throwIfNotInt<Dims>(NDRange.get_local_range());
detail::throwIfNotInt<Dims>(NDRange.get_offset());
MNDRDesc.set(std::move(NDRange));
MKernel = detail::getSyclObjImpl(std::move(Kernel));
MCGType = detail::CG::KERNEL;
Expand Down Expand Up @@ -1303,6 +1365,7 @@ class __SYCL_EXPORT handler {
(void)NumWorkGroups;
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
#else
detail::throwIfNotInt<Dims>(NumWorkGroups);
MNDRDesc.setNumWorkGroups(NumWorkGroups);
MKernel = detail::getSyclObjImpl(std::move(Kernel));
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
Expand Down Expand Up @@ -1339,7 +1402,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::throwIfNotInt<Dims>(ExecRange.get_global_range());
detail::throwIfNotInt<Dims>(ExecRange.get_local_range());
detail::throwIfNotInt<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;
Expand Down
177 changes: 177 additions & 0 deletions sycl/test/basic_tests/range_offset_fit_in_int.cpp
Original file line number Diff line number Diff line change
@@ -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. Suggest "
"disabling `id-queries-fit-in-int32' optimizations "
"flag.";

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. Suggest "
"disabling `id-queries-fit-in-int32' optimizations "
"flag.";

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;
}