Work around crashes and incorrect results in scan-based algorithms when compiling with -O0

[mmichel11]

On certain integrated graphics architectures, sub-group sizes of 32 are not supported for kernels with certain properties when compiled with -O0 using the icpx compiler. The compiler is normally able to workaround this issue by compiling to a sub-group size of 16 instead. However, in cases in which an explicit sub-group size is required, then the compiler throws an exception at JIT time. This issue directly affects our reduce-then-scan implementation which has a required sub-group size of 32.

To properly work around this issue, several things must be done. Firstly, exception handling is implemented to catch this synchronous exception while re-throwing any other exceptions back to the user. Secondly, after discussion with compiler developers, kernel compilation must be separated from execution of the kernel to prevent corruption of the underlying sycl::queue that occurs when this exception is thrown after implicit buffer accessor dependencies around the kernel have been established. To do this, kernel bundles are used to first compile the kernel before executing.

[SergeyKopienko]

I never seen before in our code the examples of usage std::optional.
From my point of view we may simple write try / catch instead of __handle_sync_sycl_exception and don't use std::optional at all.

[mmichel11]

I never seen before in our code the examples of usage std::optional'. From my point of view we may simple write try / catchinstead of__handle_sync_sycl_exceptionand don't usestd::optional' at all.

I don't have a strong preference if we choose to go with the way I implemented or just directly add try ... catch throughout this header. I initially did this to avoiding having many try catch statements here.

Let me leave this open first for others' opinions to see if they prefer a more functional approach or just directly adding try...catch throughout.

[mmichel11]

I wanted to point this unique case out and ensure there are no issues with this approach. Benchmarking the single work-group sizes, I was seeing some overheads after switching from kernel provider to compiler.

Some of the overheads seem related to the kernel bundle which is unavoidable. However, __kernel_compiler creates a std::vector to call sycl::get_kernel_bundle. This allocation / deallocation overhead on each call was leading to measurable slowdowns with small input sizes. To fix this, I have made the variable static as __kernels should be the same for each call.

This is beneficial assuming the application makes multiple calls to the scan-based algorithm which I assume is the most common case.

[dmitriy-sobolev]

I wanted to point this unique case out and ensure there are no issues with this approach.

I am thinking about possible issues with uniqueness of the kernels in some corner cases. Is it possible to end up with different kernels with the same instantiation of __parallel_transform_reduce_then_scan? I can imagine that, for example:

// no kernel specified: unnamed lambda case
sycl::queue queue_a{selector_vendor_a{});
sycl::queue queue_b{selector_vendor_b{});

// policy_a and policy_b have the same type
dpl::execution::device_policy policy_a(queue_a);
dpl::execution::device_policy policy_b(queue_b);

// ... containers and predicates have the same types
dpl::copy_if(policy_a, ...);
dpl::copy_if(policy_b, ...); // will it use the kernels compiled for the queue_a (and thus device "a")?

Such cases are highly unlikely, and I cannot think of any others. I would rather document this as a known limitation with a workaround (name the kernel) than compromise performance (if it is large enough overhead).

[mmichel11]

Thanks, I agree the risk here is low, and I believe this change is worth it if we care about the small-to-medium input size cases. Will take a further look as well and see if we need to add a known limitation.

[akukanov]

We will need to have a proper fix.
Generally, caching kernels is not our business, but the one of the SYCL runtime. Caching memory is not our business, but the one of a memory manager.
If we do caching ourselves, we need to check if the cached value is valid for the call - for example, if it was compiled for the device for which the algorithm is called. Documenting it as a limitation is not a long-term substitute for a proper, reliable implementation.

[mmichel11]

Thank you for the response. Since this does not seem like a good option to keep, it might be worth incurring the performance hit and remove the static usage. The performance impact I've measured to be around a constant 3-4 us with this caching approach for small inputs after the first call. When incurring the additional overhead each call, it adds about another 4-5 us for for a total hit of 8-9 us on Intel GPU Series Max 1550 with inclusive_scan.

Once we scale our inputs past cache sizes the impact of this change is minimized with a lower threshold if we are loading data from global memory in the cold cache case. If this performance impact seems acceptable, then I can open a PR to make this single line change.

The overarching issue here is that the SYCL specification does not specify recovery behavior after handling this sort of exception which has led to the need for this workaround using kernel bundles. I've been told that it has been discussed at the Khronos working group before but there was no consensus.

[dmitriy-sobolev]

I agree with the approach, but I need to review the PR more thoroughly. Indeed, according to the SYCL 2020 spec, launching a kernel decorated with reqd_sub_group_size attribute may throw an exception depending on what inside that kernel...

I have a question regarding that part:

Secondly, after discussion with compiler developers, kernel compilation must be separated from execution of the kernel to prevent corruption of the underlying sycl::queue

Is it a limitation/bug of the DPC++ SYCL implementation?

[dmitriy-sobolev]

Let's keep the original variant with a forwarding reference, which is aligned with other patterns, unless the change is necessary.

[mmichel11]

There's a tricky scenario with aligning the future return type from the different potential paths when we use the forwarding reference. I will take another look and see if it can be implemented with the forwarding reference.

[mmichel11]

To clarify this issue, there are some instances where forwarding the execution policy in one path and not forwarding it in the other could result in differing future return types as the _ExecutionPolicy class template in __result_and_scratch_storage which is bundled in the future return may or may not be references of device_policy. Making this a const l-value reference aligns the return.

[danhoeflinger]

Let me preface this by saying I'm not arguing for any changes in the code here, but this should be considered in a refactor of __result_and_scratch_storage, #2003 :

This is seems like it reveals an issue in the design of __result_and_scratch_storage. We shouldn't have to tip toe around the type in our surrounding function definitions.

The constructor of __result_and_scratch_storage takes a const _ExecutionPolicy&, but we also explicitly use _ExecutionPolicy frequently in defining the type of __result_and_scratch_storage types without std::decay_t. This seems ripe for problems to arise, and we should try to figure out a way to consistently arrive at the type of an appropriate __result_and_scratch_storage.

[akukanov]

The ExecutionPolicy type should not be used in templates etc. without first clearing from cv qualifiers and references. According to the standard, even the check of the is_execution_policy trait when the overload set is determined should use decay/remove_cvref. Sounds like __result_and_scratch_storage implementation or usage has a bug that needs to be fixed, instead of changing the pattern signatures, and it is orthogonal to any refactoring.

[danhoeflinger]

Yes, I can make a fix for this to be included for this release. My thought was to figure out some way to specify get the proper type in a way that didn't require the decay to make usage easier, but we should first fix the bug.

[danhoeflinger]

I believe #2031 should resolve any issues.

[akukanov]

Does it also revert the changes made by this patch? That should be a good test :)

[danhoeflinger]

No, but I can make that change.

[dmitriy-sobolev]

I wanted to point this unique case out and ensure there are no issues with this approach.

I am thinking about possible issues with uniqueness of the kernels in some corner cases. Is it possible to end up with different kernels with the same instantiation of __parallel_transform_reduce_then_scan? I can imagine that, for example:

// no kernel specified: unnamed lambda case
sycl::queue queue_a{selector_vendor_a{});
sycl::queue queue_b{selector_vendor_b{});

// policy_a and policy_b have the same type
dpl::execution::device_policy policy_a(queue_a);
dpl::execution::device_policy policy_b(queue_b);

// ... containers and predicates have the same types
dpl::copy_if(policy_a, ...);
dpl::copy_if(policy_b, ...); // will it use the kernels compiled for the queue_a (and thus device "a")?

Such cases are highly unlikely, and I cannot think of any others. I would rather document this as a known limitation with a workaround (name the kernel) than compromise performance (if it is large enough overhead).

[mmichel11]

I agree with the approach, but I need to review the PR more thoroughly. Indeed, according to the SYCL 2020 spec, launching a kernel decorated with reqd_sub_group_size attribute may throw an exception depending on what inside that kernel...

I have a question regarding that part:

Secondly, after discussion with compiler developers, kernel compilation must be separated from execution of the kernel to prevent corruption of the underlying sycl::queue

Is it a limitation/bug of the DPC++ SYCL implementation?

I think it is more of a limitation with the existing DPC++ implementation. From my discussion, it seems as if there is no easy fix which is why this workaround was recommended. From their perspective, it is not a bug as the SYCL specification does not specify any behavior regarding resubmissions to a queue after an exception such as this.

[mmichel11]

I never seen before in our code the examples of usage std::optional'. From my point of view we may simple write try / catchinstead of__handle_sync_sycl_exceptionand don't usestd::optional' at all.

I don't have a strong preference if we choose to go with the way I implemented or just directly add try ... catch throughout this header. I initially did this to avoiding having many try catch statements here.

Let me leave this open first for others' opinions to see if they prefer a more functional approach or just directly adding try...catch throughout.

@dmitriy-sobolev After looking through the PR, do you have any thoughts on Sergey's question here? Is the functional exception handler approach preferred in your opinion or just directly adding try-catch throughout parallel_backend_sycl.h?

[dmitriy-sobolev]

I never seen before in our code the examples of usage std::optional'. From my point of view we may simple write try / catchinstead of__handle_sync_sycl_exceptionand don't usestd::optional' at all.

I don't have a strong preference if we choose to go with the way I implemented or just directly add try ... catch throughout this header. I initially did this to avoiding having many try catch statements here.
Let me leave this open first for others' opinions to see if they prefer a more functional approach or just directly adding try...catch throughout.

@dmitriy-sobolev After looking through the PR, do you have any thoughts on Sergey's question here? Is the functional exception handler approach preferred in your opinion or just directly adding try-catch throughout parallel_backend_sycl.h?

I slightly prefer Sergey's option. __handle_sync_sycl_exception does not reuse much as an abstraction, but adds complexity. In contrast, __bypass_sycl_kernel_not_supported is helpful here.

[mmichel11]

@SergeyKopienko @dmitriy-sobolev I have removed __handle_sync_sycl_exception, turned __bypass_sycl_kernel_not_supported into a function, and made some other necessary changes to directly use try...catch.

I will run some internal testing to fully verify the change.

[mmichel11]

This change was necessary in a few places. When the range forwarding was removed in some places, make_zip_view resulted in a compiler error. It seemed to be caused by being unable to process l-value references to ranges. With the previous approach, we captured by value in a lambda and moved it in so this was a non-issue.

It looks like it is an instance of this issue: #1805.

[mmichel11]

clang-format wants to add an additional 4 space indentation throughout both the scan kernel implementations due to some of the changes. I disagree with this as we already have a lot of indentation here.

[SergeyKopienko]

I think better to move typedefs at this place and in other try / catch blocks outside of try / catch.

[mmichel11]

Either way is not a big deal to me, but my thought is constraining the scope of the alias as much as possible makes the most sense. Since it's only needed in this try...catch, it makes sense to limit it to this scope to me.

[SergeyKopienko]

The opposite opinion - to include into try block only the code which really may generate exception.
Writing aliases inside it you make me read what was written here and understand that it's only aliases.
But up to you.

[mmichel11]

I see your point here. I have made these changes.

[dmitriy-sobolev]

A minor comment. I am ready to approve the PR once it has been addressed.

[dmitriy-sobolev]

LGTM

[dmitriy-sobolev]

Reapproving. Please, wait for the CI to finish execution before merging.

[SergeyKopienko]

LGTM