Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
wants to merge 18 commits into
base: main
Choose a base branch
from

Conversation

mmichel11
Copy link
Contributor

@mmichel11 mmichel11 commented Jan 10, 2025

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.

IGC intentionally forces a sub-group size of 16 on certain iGPUs to
workaround a known issue. We have to determine this by first compiling
the kernels to see if the required sub-group size is respected.

Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
…xception

Signed-off-by: Matthew Michel <matthew.michel@intel.com>
…erence

Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
@mmichel11 mmichel11 added this to the 2022.8.0 milestone Jan 10, 2025
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
@mmichel11 mmichel11 changed the title [Draft] Workaround crashes and incorrect results in scan-based algorithms with integrated graphics when compiling with -O0 [Draft] Work around crashes and incorrect results in scan-based algorithms when compiling with -O0 Jan 10, 2025
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
@SergeyKopienko
Copy link
Contributor

SergeyKopienko commented Jan 14, 2025

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 mmichel11 changed the title [Draft] Work around crashes and incorrect results in scan-based algorithms when compiling with -O0 Work around crashes and incorrect results in scan-based algorithms when compiling with -O0 Jan 14, 2025
@mmichel11 mmichel11 marked this pull request as ready for review January 14, 2025 14:29
@mmichel11
Copy link
Contributor Author

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants