Skip to content

Commit b892f48

Browse files
KornevNikitaKornevNikita
authored
[SYCL] Check if reqd_work_group_size dim is valid (#12179)
The number of arguments must match the dimensionality of the work-group used to invoke the kernel.
1 parent 6baf1a3 commit b892f48

File tree

6 files changed

+86
-65
lines changed

6 files changed

+86
-65
lines changed

‎sycl/include/sycl/detail/cg_types.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,8 @@ class NDRDescT {
6464

6565
public:
6666
NDRDescT()
67-
: GlobalSize{0, 0, 0}, LocalSize{0, 0, 0}, NumWorkGroups{0, 0, 0} {}
67+
: GlobalSize{0, 0, 0}, LocalSize{0, 0, 0}, NumWorkGroups{0, 0, 0},
68+
Dims{0} {}
6869

6970
template <int Dims_> void set(sycl::range<Dims_> NumWorkItems) {
7071
for (int I = 0; I < Dims_; ++I) {

‎sycl/source/detail/program_manager/program_manager.cpp

Lines changed: 14 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -530,7 +530,8 @@ static void emitBuiltProgramInfo(const pi_program &Prog,
530530
// its ref count incremented.
531531
sycl::detail::pi::PiProgram ProgramManager::getBuiltPIProgram(
532532
const ContextImplPtr &ContextImpl, const DeviceImplPtr &DeviceImpl,
533-
const std::string &KernelName, bool JITCompilationIsRequired) {
533+
const std::string &KernelName, const NDRDescT &NDRDesc,
534+
bool JITCompilationIsRequired) {
534535
KernelProgramCache &Cache = ContextImpl->getKernelProgramCache();
535536

536537
std::string CompileOpts;
@@ -565,7 +566,7 @@ sycl::detail::pi::PiProgram ProgramManager::getBuiltPIProgram(
565566
getDeviceImage(KernelName, Context, Device, JITCompilationIsRequired);
566567

567568
// Check that device supports all aspects used by the kernel
568-
if (auto exception = checkDevSupportDeviceRequirements(Device, Img))
569+
if (auto exception = checkDevSupportDeviceRequirements(Device, Img, NDRDesc))
569570
throw *exception;
570571

571572
auto BuildF = [this, &Img, &Context, &ContextImpl, &Device, &CompileOpts,
@@ -649,7 +650,8 @@ std::tuple<sycl::detail::pi::PiKernel, std::mutex *, const KernelArgMask *,
649650
sycl::detail::pi::PiProgram>
650651
ProgramManager::getOrCreateKernel(const ContextImplPtr &ContextImpl,
651652
const DeviceImplPtr &DeviceImpl,
652-
const std::string &KernelName) {
653+
const std::string &KernelName,
654+
const NDRDescT &NDRDesc) {
653655
if (DbgProgMgr > 0) {
654656
std::cerr << ">>> ProgramManager::getOrCreateKernel(" << ContextImpl.get()
655657
<< ", " << DeviceImpl.get() << ", " << KernelName << ")\n";
@@ -685,7 +687,7 @@ ProgramManager::getOrCreateKernel(const ContextImplPtr &ContextImpl,
685687
}
686688

687689
sycl::detail::pi::PiProgram Program =
688-
getBuiltPIProgram(ContextImpl, DeviceImpl, KernelName);
690+
getBuiltPIProgram(ContextImpl, DeviceImpl, KernelName, NDRDesc);
689691

690692
auto BuildF = [this, &Program, &KernelName, &ContextImpl] {
691693
sycl::detail::pi::PiKernel Kernel = nullptr;
@@ -2679,7 +2681,8 @@ std::optional<sycl::exception> checkDevSupportJointMatrixMad(
26792681

26802682
std::optional<sycl::exception>
26812683
checkDevSupportDeviceRequirements(const device &Dev,
2682-
const RTDeviceBinaryImage &Img) {
2684+
const RTDeviceBinaryImage &Img,
2685+
const NDRDescT &NDRDesc) {
26832686
auto getPropIt = [&Img](const std::string &PropName) {
26842687
const RTDeviceBinaryImage::PropertyRange &PropRange =
26852688
Img.getDeviceRequirements();
@@ -2804,6 +2807,12 @@ checkDevSupportDeviceRequirements(const device &Dev,
28042807
Dims++;
28052808
}
28062809

2810+
if (NDRDesc.Dims != 0 && NDRDesc.Dims != static_cast<size_t>(Dims))
2811+
return sycl::exception(
2812+
sycl::errc::nd_range,
2813+
"The local size dimension of submitted nd_range doesn't match the "
2814+
"required work-group size dimension");
2815+
28072816
// The SingleDimSize was computed in an uint64_t; size_t does not
28082817
// necessarily have to be the same uint64_t (but should fit in an
28092818
// uint64_t).

‎sycl/source/detail/program_manager/program_manager.hpp

Lines changed: 10 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212
#include <detail/host_pipe_map_entry.hpp>
1313
#include <detail/kernel_arg_mask.hpp>
1414
#include <detail/spec_constant_impl.hpp>
15+
#include <sycl/detail/cg_types.hpp>
1516
#include <sycl/detail/common.hpp>
1617
#include <sycl/detail/device_global_map.hpp>
1718
#include <sycl/detail/export.hpp>
@@ -53,7 +54,8 @@ bool doesDevSupportDeviceRequirements(const device &Dev,
5354
const RTDeviceBinaryImage &BinImages);
5455
std::optional<sycl::exception>
5556
checkDevSupportDeviceRequirements(const device &Dev,
56-
const RTDeviceBinaryImage &BinImages);
57+
const RTDeviceBinaryImage &BinImages,
58+
const NDRDescT &NDRDesc = {});
5759

5860
// This value must be the same as in libdevice/device_itt.h.
5961
// See sycl/doc/design/ITTAnnotations.md for more info.
@@ -135,9 +137,11 @@ class ProgramManager {
135137
/// \param KernelName the kernel's name
136138
/// \param JITCompilationIsRequired If JITCompilationIsRequired is true
137139
/// add a check that kernel is compiled, otherwise don't add the check.
138-
sycl::detail::pi::PiProgram getBuiltPIProgram(
139-
const ContextImplPtr &ContextImpl, const DeviceImplPtr &DeviceImpl,
140-
const std::string &KernelName, bool JITCompilationIsRequired = false);
140+
sycl::detail::pi::PiProgram
141+
getBuiltPIProgram(const ContextImplPtr &ContextImpl,
142+
const DeviceImplPtr &DeviceImpl,
143+
const std::string &KernelName, const NDRDescT &NDRDesc = {},
144+
bool JITCompilationIsRequired = false);
141145

142146
sycl::detail::pi::PiProgram
143147
getBuiltPIProgram(const context &Context, const device &Device,
@@ -149,7 +153,8 @@ class ProgramManager {
149153
sycl::detail::pi::PiProgram>
150154
getOrCreateKernel(const ContextImplPtr &ContextImpl,
151155
const DeviceImplPtr &DeviceImpl,
152-
const std::string &KernelName);
156+
const std::string &KernelName,
157+
const NDRDescT &NDRDesc = {});
153158

154159
sycl::detail::pi::PiProgram
155160
getPiProgramFromPiKernel(sycl::detail::pi::PiKernel Kernel,

‎sycl/source/detail/scheduler/commands.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2549,7 +2549,7 @@ pi_int32 enqueueImpKernel(
25492549
} else {
25502550
std::tie(Kernel, KernelMutex, EliminatedArgMask, Program) =
25512551
detail::ProgramManager::getInstance().getOrCreateKernel(
2552-
ContextImpl, DeviceImpl, KernelName);
2552+
ContextImpl, DeviceImpl, KernelName, NDRDesc);
25532553
}
25542554

25552555
// We may need more events for the launch, so we make another reference.

‎sycl/test-e2e/Basic/reqd_work_group_size_matches_dimensions.cpp

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,59 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
// https://github.com/intel/llvm/issues/9353
5+
// UNSUPPORTED: hip
6+
7+
#include "sycl/sycl.hpp"
8+
9+
using namespace sycl;
10+
11+
#define CHECK_INVALID_REQD_WORK_GROUP_SIZE(Range, Item, ...) \
12+
{ \
13+
bool ExceptionThrown = false; \
14+
std::error_code Errc; \
15+
try { \
16+
q.submit([&](sycl::handler &h) { \
17+
h.parallel_for( \
18+
Range, [=](Item) [[sycl::reqd_work_group_size(__VA_ARGS__)]] {}); \
19+
}); \
20+
q.wait(); \
21+
} catch (sycl::exception & e) { \
22+
ExceptionThrown = true; \
23+
Errc = e.code(); \
24+
} \
25+
assert(ExceptionThrown && \
26+
"Invalid use of reqd_work_group_size should throw an exception."); \
27+
assert(Errc == sycl::errc::nd_range); \
28+
}
29+
30+
int main() {
31+
queue q;
32+
range<1> range1D(1);
33+
range<2> range2D(1, 1);
34+
range<3> range3D(1, 1, 1);
35+
36+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range1D, item<1> it, 1, 1)
37+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range1D, item<1> it, 1, 1, 1)
38+
39+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range2D, item<2> it, 1)
40+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range2D, item<2> it, 1, 1, 1)
41+
42+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range3D, item<3> it, 1)
43+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(range3D, item<3> it, 1, 1)
44+
45+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range1D, range1D), nd_item<1> it,
46+
1, 1)
47+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range1D, range1D), nd_item<1> it,
48+
1, 1, 1)
49+
50+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range2D, range2D), nd_item<2> it,
51+
1)
52+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range2D, range2D), nd_item<2> it,
53+
1, 1, 1)
54+
55+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range3D, range3D), nd_item<3> it,
56+
1)
57+
CHECK_INVALID_REQD_WORK_GROUP_SIZE(nd_range(range3D, range3D), nd_item<3> it,
58+
1, 1)
59+
}

‎sycl/test-e2e/Scheduler/HandleException.cpp

Lines changed: 0 additions & 53 deletions
This file was deleted.

0 commit comments

Comments
 (0)