Adding a test case for reqd_work_group_size using L0 backend. #2005

clevels · 2020-06-29T20:01:20Z

Resolved issues test case discovered.

Signed-off-by: Cory Levels cory.levels@intel.com

Resolved issues test case discovered. Signed-off-by: Cory Levels <cory.levels@intel.com>

smaslov-intel · 2020-06-30T15:06:19Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

@@ -21,6 +21,35 @@ namespace detail {

 namespace enqueue_kernel_launch {

+bool L0HandleInvalidWorkGroupSize(const device_impl &DeviceImpl,


There should be no L0-specifc in this checking, it should be covered by PI API. Can you combine this with oclHandleInvalidWorkGroupSize into just HandleInvalidWorkGroupSize?

Hi Sergey, could you please see this section here https://github.com/intel/llvm/pull/1998/files#diff-5b57d29ae1eb53bd221ca97692974846R48? I think that is valid for L0 and should be added here.

Hi @smaslov-intel , If you look at the current function L0HandleInvalidWorkGroupSize the code is doing the same thing as a couple parts of oclHandleInvalidWorkGroupSize. So agreed they should be combined. I spoke to Rehana offline and we discussed renaming oclHandleInvalidWorkGroupSize to the generic HandleInvalidWorkGroupSize and just adding conditional statements if (PlatformName.find("OpenCL") != std::string::npos) around all the checks that are openCL specific. Would this be the best approach?

Probably, yes. I'd still want to see if any backend specifics can be hidden under PI, so that those specifics are implemented in the relevant plugins not in in SYCL RT.

rbegam

LGTM

…roup_size2

…nd() method

smaslov-intel

LGTM

sycl/source/detail/error_handling/enqueue_kernel.cpp

rbegam

LGTM

sycl/source/detail/error_handling/enqueue_kernel.cpp

bader · 2020-07-03T09:10:14Z

sycl/source/detail/error_handling/enqueue_kernel.cpp


  size_t VerSize = 0;
  Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION, 0,
                                          nullptr, &VerSize);
-  assert(VerSize >= 10 &&
-         "Unexpected device version string"); // strlen("OpenCL X.Y")
+  if (Platform.get_backend() == cl::sycl::backend::opencl) {


I expected SYCL runtime to performs back-end agnostic checks only, so parsing OpenCL version in this file doesn't seem right to me.
Shouldn't we sink most of these error checks into corresponding plug-ins?

sycl/source/detail/error_handling/enqueue_kernel.cpp

bader · 2020-07-03T09:18:43Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

+        throw sycl::nd_range_error(
+            "Total number of work-items in a work-group cannot exceed "
+            "info::kernel_work_group::work_group_size which is equal to " +
+                std::to_string(KernelWGSize) + " for this kernel",


"for this kernel" doesn't help if I have more than one. I suggest adding a kernel name to the message.

Note, however, that:

Kernels do not always have a name that is meaningful to the user since the KernelName template parameter is optional as of SYCL 2020 provisional.

The exception is thrown synchronously when submitting a kernel, so the problematic kernel should be clear if the application uses try/catch around the code that submits the kernel.

Adding the name of the kernel to the exception string does seem reasonable in the case where the user specified a name. If we decide to do that, we should do it uniformly for all the exceptions diagnosed in this function, not just for this one exception.

sycl/source/detail/error_handling/enqueue_kernel.cpp

bader · 2020-07-03T09:25:21Z

sycl/test/basic_tests/reqd_work_group_size.cpp

@@ -0,0 +1,109 @@
+// XFAIL: cuda || opencl


There were changes in opencl specific code, so I suggest validate opencl back-end as well.

I think it better to use unittest framework with mocking API.

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

Co-authored-by: smaslov-intel <48694368+smaslov-intel@users.noreply.github.com>

bader · 2020-07-03T14:48:02Z

@clevels and I agreed to address the rest of the comments in follow-up PRs.

gmlueck · 2020-07-06T12:13:38Z

sycl/test/basic_tests/reqd_work_group_size.cpp

+    return 1;
+  }
+
+  // Positive test-cases that should pass on any underlying OpenCL runtime


This test case is not specific to OpenCL, so I would change this comment to something like: "... that should pass on any backend".

gmlueck · 2020-07-06T12:50:29Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

+  if (Platform.get_backend() == cl::sycl::backend::opencl) {
+    assert(VerSize >= 10 &&
+           "Unexpected device version string"); // strlen("OpenCL X.Y")
+  }
  string_class VerStr(VerSize, '\0');
  Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION,
                                          VerSize, &VerStr.front(), nullptr);


This comment applies to the entire block of code:

size_t VerSize = 0; Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION, 0, nullptr, &VerSize); if (Platform.get_backend() == cl::sycl::backend::opencl) { assert(VerSize >= 10 && "Unexpected device version string"); // strlen("OpenCL X.Y") } string_class VerStr(VerSize, '\0'); Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION, VerSize, &VerStr.front(), nullptr); const char *Ver = &VerStr[7]; // strlen("OpenCL ")

There are two problems here:

The Ver pointer could point to uninitialized memory when the backend is not OpenCL.

The name of the Ver variable should indicate that it is specific to OpenCL.

It would also be nice to include a comment, so readers understand why you are getting the version only for OpenCL.
I'd suggest rewriting the code to follow this form:

// Some of the error handling below is special for particular OpenCL versions. // If this is an OpenCL backend, get the version. const char *OpenClVer = nullptr; string_class OpenClVerStr; if (Platform.get_backend() == cl::sycl::backend::opencl) { size_t OclVerSize = 0; Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION, 0, nullptr, &OclVerSize); assert(VerSize >= 10 && "Unexpected device version string"); // strlen("OpenCL X.Y") OpenClVerStr.assign(OclVerSize, '\0'); Plugin.call<PiApiKind::piDeviceGetInfo>(Device, PI_DEVICE_INFO_VERSION, OclVerSize, &OpenClVerStr.front(), nullptr); OpenClVer = &OpenClVerStr[7]; // strlen("OpenCL ") }

gmlueck · 2020-07-06T12:55:30Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

+            "OpenCL 1.x and 2.0 requires to pass local size argument even if "
+            "required work-group size was specified in the program source",
+                          PI_INVALID_WORK_GROUP_SIZE);

            
+    }
    // Any OpenCL version:


This error check is not specific to OpenCL, so remove this part of the comment.

gmlueck · 2020-07-06T13:07:11Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

+                std::to_string(MaxWGSize),
+            PI_INVALID_WORK_GROUP_SIZE);
+    } else {
+      // RELEVENT // OpenCL 2.x:


I see you added the word "RELEVANT" to this comment. Was that a note to yourself indicating that this arm of the if statement was relevant to L0 (non-OpenCL backends)? If so, I think you should either change the code like this:

if (Platform.get_backend() == cl::sycl::backend::opencl && OpenClVer[0] == '1') { // Do the first check } else { // Do the second check }

Or, if you are planning to do this as part of a subsequent PR, add a more informative TODO comment indicating what needs to be done.

gmlueck · 2020-07-06T13:21:16Z

sycl/source/detail/error_handling/enqueue_kernel.cpp

      }
    }
+    throw sycl::nd_range_error(
+        "Non-uniform work-groups are not supported by the target device",
+        PI_INVALID_WORK_GROUP_SIZE);
  }


All of the checks above are done only for the OpenCL backend. I presume something is needed for the other backends too? If so, I'd recommend at least adding a TODO comment to that effect:

} else { // TODO: Similar checks should be done for the non-OpenCL backends. }

Here, the else should match the if (Platform.get_backend() == cl::sycl::backend::opencl) on line 143.

Currently only to DebugInfo/X86 Currently failing tests can be noticed by RUNx line Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com> Original commit: KhronosGroup/SPIRV-LLVM-Translator@772c7be

Adding a test case for reqd_work_size using L0 backend.

fc6914b

Resolved issues test case discovered. Signed-off-by: Cory Levels <cory.levels@intel.com>

clevels requested review from smaslov-intel and a team as code owners June 29, 2020 20:01

smaslov-intel reviewed Jun 30, 2020

View reviewed changes

rbegam reviewed Jun 30, 2020

View reviewed changes

clevels added 3 commits July 2, 2020 08:02

Merge remote-tracking branch 'upstream/sycl' into clevels/reqd_work_g…

53e2456

…roup_size2

Inline HandleInvalidWorkGroupSize function, formatting, use get_backe…

2a3296e

…nd() method

Remove old comment.

af50f43

smaslov-intel previously approved these changes Jul 2, 2020

View reviewed changes

sycl/source/detail/error_handling/enqueue_kernel.cpp Outdated Show resolved Hide resolved

rbegam previously approved these changes Jul 2, 2020

View reviewed changes

mark xfail with cuda.

14673ff

clevels dismissed stale reviews from rbegam and smaslov-intel via 14673ff July 2, 2020 18:39

smaslov-intel previously approved these changes Jul 2, 2020

View reviewed changes

rbegam previously approved these changes Jul 2, 2020

View reviewed changes

bader requested changes Jul 3, 2020

View reviewed changes

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

fbb0655

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

clevels dismissed stale reviews from rbegam and smaslov-intel via fbb0655 July 3, 2020 14:09

clevels and others added 7 commits July 3, 2020 09:10

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

9182228

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

12487ed

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

b9deaf3

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

33e41cf

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

Update sycl/source/detail/error_handling/enqueue_kernel.cpp

d5e7a22

Co-authored-by: smaslov-intel <48694368+smaslov-intel@users.noreply.github.com>

Address comments.

361fe85

Fix formatting.

c593556

bader previously approved these changes Jul 3, 2020

View reviewed changes

Update error message in parallel_for_range.cpp

e2c142f

clevels dismissed bader’s stale review via e2c142f July 3, 2020 16:33

bader approved these changes Jul 3, 2020

View reviewed changes

bader merged commit 03ef819 into intel:sycl Jul 4, 2020

gmlueck reviewed Jul 6, 2020

View reviewed changes

rbegam mentioned this pull request Jul 9, 2020

[SYCL] Implements error handling for invalid work group size for enqueue NDRange kernel while using level0. #1998

Closed

jsji mentioned this pull request May 16, 2023

llvmspirv pulldown #9479

Closed

jsji mentioned this pull request May 16, 2023

llvmspirv pulldown sys-ce-bb/llvm#2

Closed

This was referenced May 16, 2023

llvmspirv pulldown sys-ce-bb/llvm#3

Closed

test1 sys-ce-bb/llvm#4

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Adding a test case for reqd_work_group_size using L0 backend. #2005

Adding a test case for reqd_work_group_size using L0 backend. #2005

		@@ -21,6 +21,35 @@ namespace detail {

		namespace enqueue_kernel_launch {

		bool L0HandleInvalidWorkGroupSize(const device_impl &DeviceImpl,

Adding a test case for reqd_work_group_size using L0 backend. #2005

Adding a test case for reqd_work_group_size using L0 backend. #2005

Conversation

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment