From 3c91ec66208c2cc78a155bf456fd88ccdf500d38 Mon Sep 17 00:00:00 2001 From: Charles Congdon Date: Thu, 28 May 2026 13:37:23 -0700 Subject: [PATCH] Updates for oneAPI 2026.0 tools and drivers Signed-off-by: Charles Congdon --- .../guided_matrix_mult_BadBuffers/README.md | 25 +- .../src/b1_matrix_mul_null_usm.cpp | 4 + .../src/b2_matrix_mul_usm.cpp | 4 + .../guided_matrix_mult_Exceptions/README.md | 50 +- .../src/1_matrix_mul_null_pointer.cpp | 4 + .../src/2_matrix_mul_multi_offload.cpp | 4 + .../src/3_matrix_mul.cpp | 4 + .../README.md | 433 ++++++++++++------ .../src/1_matrix_mul_invalid_contexts.cpp | 31 +- .../src/2_matrix_mul.cpp | 17 + .../README.md | 47 +- .../guided_matrix_mult_SLMSize/README.md | 78 ++-- .../src/1_matrix_mul_SLM_size.cpp | 4 + .../src/2_matrix_mul.cpp | 4 + 14 files changed, 479 insertions(+), 230 deletions(-) diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md index 9c8398b822..fe92dbbd9e 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md @@ -1,6 +1,6 @@ # `Guided Matrix Multiplication Bad Buffers` Sample -The `Guided Matrix Multiplication Bad Buffers` sample demonstrates how to use several tools in the Intel® oneAPI Base Toolkit (Base Kit) to triage incorrect use of the SYCL language. +The `Guided Matrix Multiplication Bad Buffers` sample demonstrates how to use several tools in Intel® oneAPI to triage incorrect use of the SYCL language. The sample is a simple program that multiplies together two large matrices and verifies the results. @@ -31,9 +31,9 @@ The sample includes different versions of a simple matrix multiplication program | Optimized for | Description |:--- |:--- | OS | Ubuntu* 24.04 LTS -| Hardware | GEN9 or newer -| Software | Intel® oneAPI DPC++/C++ Compiler 2025.3
Intel® Distribution for GDB* 2025.3
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). -| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.31 or later from https://dgpu-docs.intel.com/releases/releases.html +| Intel Graphics Hardware | GEN9 or newer +| Software | Intel® oneAPI DPC++/C++ Compiler 2026.0
Intel® Distribution for GDB* 2026.0
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). +| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.59 or later from https://dgpu-docs.intel.com/releases/releases.html ## Key Implementation Details @@ -46,7 +46,7 @@ The basic SYCL* standards implemented in the code include the use of the followi ## Set Environment Variables -When working with the command-line interface (CLI), configure the oneAPI toolkit environment variables. Set up your CLI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. +When working with the command-line interface (CLI), set up your oneAPI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. ## Build the `Guided Matrix Multiplication Bad Buffers` Programs @@ -121,7 +121,6 @@ the `VERBOSE=1` argument: ``` make VERBOSE=1 ``` -If you receive an error message, troubleshoot the problem using the **Diagnostics Utility for Intel® oneAPI Toolkits**. The diagnostic utility provides configuration and system checks to help find missing dependencies, permissions errors, and other issues. See the *[Diagnostics Utility for Intel® oneAPI Toolkits User Guide](https://www.intel.com/content/www/us/en/docs/oneapi/user-guide-diagnostic-utility/current/overview.html)* for more information on using the utility. ## Guided Debugging @@ -129,7 +128,7 @@ If you receive an error message, troubleshoot the problem using the **Diagnostic These instructions assume you have installed the Intel® Distribution for GDB* and have a basic working knowledge of GDB. ### Setting up to Debug on the GPU -To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GDB-PVC debugger](https://dgpu-docs.intel.com/system-user-guides/DNP-Max-1100-userguide/DNP-Max-1100-userguide.html#gdb-pvc-debugger)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. +To learn how setup and use the Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GPU Debugging](https://dgpu-docs.intel.com/driver/gpu-debugging.html)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. Documentation on using the debugger in a variety of situations can be found at *[Debug Examples in Linux](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/tutorial-debugging-dpcpp-linux/current/overview.html)* @@ -137,12 +136,12 @@ Documentation on using the debugger in a variety of situations can be found at * ### Getting the Tracing and Profiling Tool -In this tutorial, the instructions require a utility that was not installed with the Intel® oneAPI Base Toolkit (Base Kit). +In this tutorial, the instructions require a utility that was not installed with the Intel® oneAPI. To complete the steps in the following section, you must download the [Unified Tracing and Profiling Tool](https://github.com/intel/pti-gpu/tree/master/tools/unitrace) code from GitHub and build the utility. The build instructions are included in the README in the GitHub repository. This build will go much more smoothly if you first install the latest drivers from [the Intel GPU driver download site](https://dgpu-docs.intel.com/driver/overview.html), especially the development packages (only available in the Data Center GPU driver install). Once you have built the utility, you invoke it on the command line in front of your program (similar to using GDB). ### Guided Instructions for Zero Buffer using Address Sanitizer -A recent addition to the oneAPI compiler is that ability to use the "Address Sanitizer" you may have seen when using [GCC](https://gcc.gnu.org/onlinedocs/gcc/Instrumentation-Options.html) or [CLANG](https://clang.llvm.org/docs/AddressSanitizer.html) to catch invalid pointer addresses at runtime on the GPU rather than the host. This will require a special build of the application. +The oneAPI compiler has the ability to use the "Address Sanitizer" you may have seen when using [GCC](https://gcc.gnu.org/onlinedocs/gcc/Instrumentation-Options.html) or [CLANG](https://clang.llvm.org/docs/AddressSanitizer.html) to catch invalid pointer addresses at runtime on the GPU rather than the host. This will require a special build of the application. 1. Compile a version of the program with device-side Address Sanitizer (assuming that you are in the `build` directory) ``` @@ -198,7 +197,7 @@ A recent addition to the oneAPI compiler is that ability to use the "Address San ### Guided Instructions for Zero Buffer using gdb-oneapi and the OpenCL CPU device -In `a1_matrix_mul_zero_buff`, a zero-element buffer is passed to a SYCL submit `lambda` function. **This will cause the application to crash.** We saw in the previous section how we can catch this with the device-side Address Sanitizer. But what if the bad array allocation occured somewhere else deep in the program? How would we track the problem back to its source? Let's try one technique to locate the source of the error. +In `a1_matrix_mul_zero_buff`, a zero-element buffer is passed to a SYCL submit `lambda` function. **This will cause the application to crash.** We saw in the previous section how we can catch this with the device-side Address Sanitizer. But what if the bad array allocation occurred somewhere else deep in the program? How would we track the problem back to its source? Let's try one technique to locate the source of the error. 1. Run the program without the debugger. @@ -353,7 +352,7 @@ In `a1_matrix_mul_zero_buff`, a zero-element buffer is passed to a SYCL submit ` 0}}, static dimensions = }, MOffset = {> = {common_array = {0, 0}}, static dimensions = }}} ``` - Clearly there is a mismatch here! 'a' has no space reserved for it, yet we will be iterating over 150 by 300 elements (and updating element 131 by 0 in this thread), which is clearly an error. + Clearly there is a mismatch here: 'a' has no space reserved for it, yet we will be iterating over 150 by 300 elements (and updating element 131 by 0 in this thread), which is clearly an error. 9. To further root-cause the error, we will need to restart the program and look at the values of the buffers behind the accessors (`a_buf` and `b_buf`), which are not in scope in any of our stack frames. We'll set some breakpoints at the `parallel_for` statements where they are initialized. @@ -416,7 +415,7 @@ In `a1_matrix_mul_zero_buff`, a zero-element buffer is passed to a SYCL submit ` ``` ### Guided Instructions for Null Device Pointer using Address Sanitizer -Let us use the Address Sanitizer again to catch invalid pointer addresses at runtime, this time in code that makes use of explicit device memory allocations rather than using SYCL buffers. This will require a special build of the application. +Let us use the Address Sanitizer again to catch invalid pointer addresses at runtime, this time in code that makes use of explicit device-memory allocations rather than using SYCL buffers. This will require a special build of the application. 1. Compile a version of the program with device-side Address Sanitizer (assuming that you are in the `build` directory) ``` @@ -544,7 +543,7 @@ In `b1_matrix_mul_null_usm.cpp` a bad (in this case, null) pointer that is suppo #### Debugging the Problem -Why did we try with multiple backends? If one had shown correct or incorrect results, and one had crashed, we might be facing a race condition that only occasionally manifests as something that goes terribly wrong. Or one of the backbends might have a bug while the others do not. But here all three crash, so it's likely the program is doing something illegal to memory. The host CPU is a particularly good place to test for illegal memory accesses, because the CPU never allows pointers with an address within a few kilobytes of address `0x0`, while this may be legally allocated memory on the GPU. +Why did we try with multiple backends? If one had shown correct or incorrect results, and one had crashed, we might be facing a race condition that only occasionally manifests when something goes terribly wrong. Or one of the backbends might have a bug while the others do not. But here all three crash, so it's likely the program is doing something illegal to memory. The host CPU is a particularly good place to test for illegal memory accesses, because the CPU never allows pointers with an address within a few kilobytes of address `0x0`, while this may be legally allocated memory on the GPU. Another reason to try different backends is that debugging support may differ between different GPU drivers and/or different GPU models. Debugging the program using the OpenCL™ CPU driver gets around these issues. diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b1_matrix_mul_null_usm.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b1_matrix_mul_null_usm.cpp index e7841aec86..546e37050e 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b1_matrix_mul_null_usm.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b1_matrix_mul_null_usm.cpp @@ -132,6 +132,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b2_matrix_mul_usm.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b2_matrix_mul_usm.cpp index ac97b65b23..503265f223 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b2_matrix_mul_usm.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b2_matrix_mul_usm.cpp @@ -130,6 +130,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/README.md index dd8b72569f..b191c422a1 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/README.md @@ -1,6 +1,6 @@ # `Guided Matrix Multiplication Exception` Sample -The `Guided Matrix Multiplication Exception` sample demonstrates a guided approach to debugging SYCL exceptions from incorrect use of the SYCL* API. It uses the Intel® oneAPI Base Toolkit (Base Kit) and several tools included in the Base Kit. +The `Guided Matrix Multiplication Exception` sample demonstrates an approach to debugging SYCL exceptions from incorrect use of the SYCL* API using several tools in Intel® oneAPI. The sample code is a simple program that multiplies together two large matrices and verifies the results. @@ -34,9 +34,9 @@ The sample includes three different versions of some simple matrix multiplicatio | Optimized for | Description |:--- |:--- | OS | Ubuntu* 24.04 LTS -| Hardware | GEN9 or newer -| Software | Intel® oneAPI DPC++/C++ Compiler 2025.3
Intel® Distribution for GDB* 2025.3 -| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.31 or later from https://dgpu-docs.intel.com/releases/releases.html +| Intel Graphics Hardware | GEN9 or newer +| Software | Intel® oneAPI DPC++/C++ Compiler 2026.0
Intel® Distribution for GDB* 2026.0 +| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.59 or later from https://dgpu-docs.intel.com/releases/releases.html ## Key Implementation Details @@ -50,7 +50,7 @@ The basic SYCL* standards implemented in the code include the use of the followi ## Set Environment Variables -When working with the command-line interface (CLI), you should configure the oneAPI toolkits using environment variables. Set up your CLI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries, and tools are ready for development. +When working with the command-line interface (CLI), set up your oneAPI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. ## Build and Run the `Guided Matrix Multiplication Exception` Programs @@ -121,7 +121,6 @@ the `VERBOSE=1` argument: ``` make VERBOSE=1 ``` -If you receive an error message, troubleshoot the problem using the **Diagnostics Utility for Intel® oneAPI Toolkits**. The diagnostic utility provides configuration and system checks to help find missing dependencies, permissions errors, and other issues. See the *[Diagnostics Utility for Intel® oneAPI Toolkits User Guide](https://www.intel.com/content/www/us/en/docs/oneapi/user-guide-diagnostic-utility/current/overview.html)* for more information on using the utility. ## Guided Debugging @@ -129,7 +128,7 @@ If you receive an error message, troubleshoot the problem using the **Diagnostic These instructions assume you have installed the Intel® Distribution for GDB* and have a basic working knowledge of GDB. ### Setting up to Debug on the GPU -To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GDB-PVC debugger](https://dgpu-docs.intel.com/system-user-guides/DNP-Max-1100-userguide/DNP-Max-1100-userguide.html#gdb-pvc-debugger)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. +To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GPU Debugging](https://dgpu-docs.intel.com/driver/gpu-debugging.html)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. Documentation on using the debugger in a variety of situations can be found at *[Debug Examples in Linux](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/tutorial-debugging-dpcpp-linux/current/overview.html)* @@ -236,7 +235,7 @@ In the second version, the code attempts to execute more than one offload statem (gdb) backtrace ``` -4. Notice in the results (which should look something like the following) that the exception (frame 8) was triggered around line 98 (frame 17): +4. Notice in the results (which should look something like the following) that the exception (frame 8) was triggered around line 98 (frame 16): ``` #0 __pthread_kill_implementation (no_tid=0, signo=6, threadid=) at ./nptl/pthread_kill.c:44 #1 __pthread_kill_internal (signo=6, threadid=) at ./nptl/pthread_kill.c:78 @@ -247,30 +246,29 @@ In the second version, the code attempts to execute more than one offload statem #6 0x00007ffff78bb0da in ?? () from /lib/x86_64-linux-gnu/libstdc++.so.6 #7 0x00007ffff78a5a55 in std::terminate() () from /lib/x86_64-linux-gnu/libstdc++.so.6 #8 0x00007ffff78bb391 in __cxa_throw () from /lib/x86_64-linux-gnu/libstdc++.so.6 - #9 0x00007ffff7f11ed0 in sycl::_V1::handler::memcpy(void*, void const*, unsigned long) () - from /opt/intel/oneapi/compiler/2025.3/lib/libsycl.so.8 - #10 0x0000000000404812 in main::{lambda(auto:1&)#1}::operator()(sycl::_V1::handler&) const ( - this=0x7fffffffb098, h=sycl::handler& = {...}) + #9 0x00007ffff7f13f61 in sycl::_V1::handler::memcpy(void*, void const*, unsigned long) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #10 0x0000000000404762 in main::{lambda(auto:1&)#1}::operator()(sycl::_V1::handler&) const ( + this=0x7fffffffb318, h=sycl::handler& = {...}) at Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp:100 - #11 0x00000000004047ad in sycl::_V1::detail::type_erased_cgfo_ty::invoker::call(void const*, sycl::_V1::handler&) (object=0x7fffffffb098, cgh=sycl::handler& = {...}) - at /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/handler.hpp:190 - #12 0x00007ffff7e8a4a4 in sycl::_V1::detail::queue_impl::submit_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::queue_impl*, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::v1::SubmissionInfo const&) () from /opt/intel/oneapi/compiler/2025.3/lib/libsycl.so.8 - #13 0x00007ffff7e90022 in sycl::_V1::detail::queue_impl::submit_with_event(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::v1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) () - from /opt/intel/oneapi/compiler/2025.3/lib/libsycl.so.8 - #14 0x00007ffff7f58844 in sycl::_V1::queue::submit_with_event_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::v1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) const () - from /opt/intel/oneapi/compiler/2025.3/lib/libsycl.so.8 - #15 0x0000000000407392 in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb570, Props=..., CGF=..., CodeLoc=...) - at /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:3762 - #16 0x00000000004042a1 in sycl::_V1::queue::submit(main::{lambda(auto:1&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb570, CGF=..., CodeLoc=...) - at /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:429 - #17 0x0000000000403eac in main () + #11 0x00000000004046fd in sycl::_V1::detail::type_erased_cgfo_ty::invoker::call(void const*, sycl::_V1::handler&) (object=0x7fffffffb318, cgh=sycl::handler& = {...}) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/handler.hpp:191 + #12 0x00007ffff7e92c6d in sycl::_V1::detail::queue_impl::submit_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::_V1::SubmissionInfo const&) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #13 0x00007ffff7f58ac9 in sycl::_V1::queue::submit_with_event_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::_V1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) const () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #14 0x00000000004072be in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb780, Props=..., CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:3700 + #15 0x00000000004041f1 in sycl::_V1::queue::submit(main::{lambda(auto:1&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb780, CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:441 + #16 0x0000000000403e7c in main () at Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp:98 ``` 5. Examine the last frame (it may be different from the output above) using the following command: ``` - (gdb) frame 17 + (gdb) frame 16 ``` You may need to issue this command twice before you see output similar to the following example: ``` @@ -294,7 +292,7 @@ In the second version, the code attempts to execute more than one offload statem As the exception reported, we are trying to do two memory copies to the device within the `submit` statement, where only a single `parallel_for` or `memcpy` is allowed. -7. To fix the error, remove the extra `memcpy` from the code above. If this statement were actually issuing two different `memcpy` statements, you would update the code to break this up into two `submit` statements, each with a single `memcpy` . +7. To fix the error, remove the extra `memcpy` from the code above (as in `3_matrix_mul.cpp`). If this statement were actually issuing two different `memcpy` statements, you would update the code to break this up into two `submit` statements, each with a single `memcpy` . ``` q.submit([&](auto &h) { h.memcpy(); diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/1_matrix_mul_null_pointer.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/1_matrix_mul_null_pointer.cpp index 6d390a7a63..355b45a475 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/1_matrix_mul_null_pointer.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/1_matrix_mul_null_pointer.cpp @@ -130,6 +130,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp index fe8b9b85e4..98f38d6af3 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/2_matrix_mul_multi_offload.cpp @@ -131,6 +131,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/3_matrix_mul.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/3_matrix_mul.cpp index ac97b65b23..503265f223 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/3_matrix_mul.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_Exceptions/src/3_matrix_mul.cpp @@ -130,6 +130,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md index 7bca9f4b79..d2621a7a34 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md @@ -1,8 +1,6 @@ # `Guided Matrix Multiplication Invalid Contexts` Sample -The `Guided Matrix Multiplication Invalid Contexts` sample demonstrates how to -use the Intel® oneAPI Base Toolkit (Base Kit) and several tools found in it to -triage incorrect use of the SYCL language. +The `Guided Matrix Multiplication Invalid Contexts` sample demonstrates how to use several tools in Intel® oneAPI to triage incorrect use of the SYCL language. The sample is a simple program that multiplies together two large matrices and verifies the results. @@ -40,9 +38,9 @@ program. | Optimized for | Description |:--- |:--- | OS | Ubuntu* 24.04 LTS -| Hardware | GEN9 or newer -| Software | Intel® oneAPI DPC++/C++ Compiler 2025.3
Intel® Distribution for GDB* 2025.3
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). -| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.31 or later from https://dgpu-docs.intel.com/releases/releases.html +| Intel GraphicsHardware | GEN9 or newer +| Software | Intel® oneAPI DPC++/C++ Compiler 2026.0
Intel® Distribution for GDB* 2026.0
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). +| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.59 or later from https://dgpu-docs.intel.com/releases/releases.html ## Key Implementation Details @@ -59,11 +57,7 @@ following: ## Set Environment Variables -When working with the command-line interface (CLI), you should configure the -oneAPI toolkits using environment variables. Set up your CLI environment by -sourcing the `setvars` script every time you open a new terminal window. This -practice ensures that your compiler, libraries, and tools are ready for -development. +When working with the command-line interface (CLI), set up your oneAPI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. ## Build and Run the `Guided Matrix Multiplication Invalid Contexts` Programs @@ -146,8 +140,6 @@ the `VERBOSE=1` argument: make VERBOSE=1 ``` -If you receive an error message, troubleshoot the problem using the **Diagnostics Utility for Intel® oneAPI Toolkits**. The diagnostic utility provides configuration and system checks to help find missing dependencies, permissions errors, and other issues. See the *[Diagnostics Utility for Intel® oneAPI Toolkits User Guide](https://www.intel.com/content/www/us/en/docs/oneapi/user-guide-diagnostic-utility/current/overview.html)* for more information on using the utility. - ## Guided Debugging @@ -155,7 +147,7 @@ The following instructions assume you have installed Intel® Distribution for GDB* and have a basic working knowledge of GDB. ### Setting up to Debug on the GPU -To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GDB-PVC debugger](https://dgpu-docs.intel.com/system-user-guides/DNP-Max-1100-userguide/DNP-Max-1100-userguide.html#gdb-pvc-debugger)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. +To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GPU Debugging](https://dgpu-docs.intel.com/driver/gpu-debugging.html)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. Documentation on using the debugger in a variety of situations can be found at *[Debug Examples in Linux](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/tutorial-debugging-dpcpp-linux/current/overview.html)* @@ -163,16 +155,28 @@ Documentation on using the debugger in a variety of situations can be found at * ### Getting the Tracing and Profiling Tool -In this tutorial, the instructions require a utility that was not installed with the Intel® oneAPI Base Toolkit (Base Kit). +In this tutorial, the instructions require a utility that was not installed with Intel® oneAPI. To complete the steps in the following section, you must download the [Unified Tracing and Profiling Tool](https://github.com/intel/pti-gpu/tree/master/tools/unitrace) code from GitHub and build the utility. The build instructions are included in the README in the GitHub repository. This build will go much more smoothly if you first install the latest drivers from [the Intel GPU driver download site](https://dgpu-docs.intel.com/driver/overview.html), especially the development packages (only available in the Data Center GPU driver install). Once you have built the utility, you invoke it on the command line in front of your program (similar to using GDB). ### Check the Programs -1. Notice that both versions of the application run to completion and report - correct results. +1. Notice that both versions of the application (`1_matrix_mul_invalid_contexts` and `2_matrix_mul`) run to completion and report correct results. -2. SYCL applications use the Level Zero runtime by default with an Intel GPU. What happens if you use OpenCL™ software to run `1_matrix_mul_invalid_contexts` on the GPU? +2. Check the available compute devices available on your system using the `sycl-ls` command: + + ``` + $ sycl-ls + [level_zero:gpu][level_zero:0] Intel(R) oneAPI Unified Runtime over Level-Zero, Intel(R) Data Center GPU Max 1550 12.60.7 [1.6.33578+51] + [level_zero:gpu][level_zero:1] Intel(R) oneAPI Unified Runtime over Level-Zero, Intel(R) Data Center GPU Max 1550 12.60.7 [1.6.33578+51] + [opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz OpenCL 3.0 (Build 0) [2026.21.3.0.31_160000] + [opencl:gpu][opencl:1] Intel(R) OpenCL Graphics, Intel(R) Data Center GPU Max 1550 OpenCL 3.0 NEO [25.18.33578] + [opencl:gpu][opencl:2] Intel(R) OpenCL Graphics, Intel(R) Data Center GPU Max 1550 OpenCL 3.0 NEO [25.18.33578] + [opencl:cpu][opencl:3] Intel(R) OpenCL, Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz OpenCL 3.0 (Build 0) [2023.16.7.0.21_160000] + ``` + > **Note:** If you have only one `[level_zero:gpu]` device listed, or the order is different from the above, the the main example below may not work. Try to follow through anyway, and then try the bonus sample at the end of this document, which should work no matter what system configuration. + +3. SYCL applications use the Level Zero runtime by default with an Intel GPU. But what happens if you use OpenCL™ runtime to run `1_matrix_mul_invalid_contexts` on the GPU? ``` ONEAPI_DEVICE_SELECTOR=opencl:gpu ./1_matrix_mul_invalid_contexts @@ -194,21 +198,8 @@ To complete the steps in the following section, you must download the [Unified T Aborted (core dumped) ``` - > **Note:** this will only work if the `sycl-ls` command shows OpenCL - devices for the graphics card, such as like this: - - ``` - $ sycl-ls - [opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz OpenCL 3.0 (Build 0) [2024.18.6.0.02_160000] - [opencl:gpu][opencl:1] Intel(R) OpenCL Graphics, Intel(R) Data Center GPU Max 1550 OpenCL 3.0 NEO [24.22.29735.27] - [opencl:gpu][opencl:2] Intel(R) OpenCL Graphics, Intel(R) Data Center GPU Max 1550 OpenCL 3.0 NEO [24.22.29735.27] - [opencl:cpu][opencl:3] Intel(R) OpenCL, Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz OpenCL 3.0 (Build 0) [2023.16.7.0.21_160000] - [opencl:fpga][opencl:4] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.7.0.21_160000] - [level_zero:gpu][level_zero:0] Intel(R) Level-Zero, Intel(R) Data Center GPU Max 1550 1.3 [1.3.29735] - [level_zero:gpu][level_zero:1] Intel(R) Level-Zero, Intel(R) Data Center GPU Max 1550 1.3 [1.3.29735] - ``` - - > If you are missing `[opencl:gpu]` devices you may have to add the necessary libraries to your device path by setting the appropriate path in `DRIVERLOC` and then running the following four commands (for Ubuntu - adapt for other OSes): + > **Note:** you will only see the above if the `sycl-ls` command shows OpenCL + devices for the graphics card. If you are missing `[opencl:gpu]` devices you may have to add the necessary libraries to your device path by setting the appropriate path in `DRIVERLOC` and then running the following four commands (for Ubuntu - adapt for other OSes): ``` export DRIVERLOC=/usr/lib/x86_64-linux-gnu @@ -216,9 +207,11 @@ To complete the steps in the following section, you must download the [Unified T export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$DRIVERLOC export PATH=$PATH:/opt/intel/oneapi:$DRIVERLOC ``` + + Then issue the `ONEAPI_DEVICE_SELECTOR=opencl:gpu ./1_matrix_mul_invalid_contexts` command again to see the above error. -3. Check the results on the **CPU** using OpenCL. +4. Check the results on the **CPU** using OpenCL. ``` ONEAPI_DEVICE_SELECTOR=opencl:cpu ./1_matrix_mul_invalid_contexts @@ -226,7 +219,7 @@ To complete the steps in the following section, you must download the [Unified T Interestingly, this runs just fine. In the next sections we will try to explain the inconsistency. ### Guided Instructions for Zero Buffer using Address Sanitizer -A recent addition to the oneAPI compiler is that ability to use the "Address Sanitizer" you may have seen when using [GCC](https://gcc.gnu.org/onlinedocs/gcc/Instrumentation-Options.html) or [CLANG](https://clang.llvm.org/docs/AddressSanitizer.html) to catch invalid pointer addresses at runtime on the GPU rather than the host. This will require a special build of the application. +The oneAPI compiler has the ability to use the "Address Sanitizer" you may have seen when using [GCC](https://gcc.gnu.org/onlinedocs/gcc/Instrumentation-Options.html) or [CLANG](https://clang.llvm.org/docs/AddressSanitizer.html) to catch invalid pointer addresses at runtime *on the GPU rather* than the host. This will require a special build of the application. 1. Compile a version of the program with device-side address sanitizer (assuming that you are in the `build` directory) ``` @@ -239,30 +232,35 @@ A recent addition to the oneAPI compiler is that ability to use the "Address San ./1_matrix_mul_invalid_contexts_asan Initializing ==== DeviceSanitizer: ASAN + Devices: + [0] Intel(R) Data Center GPU Max 1550, 12.60.7 [1.6.33578+51] 1.6.33578+51, Intel(R) Corporation + [1] Intel(R) Data Center GPU Max 1550, 12.60.7 [1.6.33578+51] 1.6.33578+51, Intel(R) Corporation + [2] Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz, OpenCL 3.0 (Build 0) [2026.21.3.0.31_160000] 3.0, Intel(R) Corporation + [3] Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz, OpenCL 3.0 (Build 0) [2023.16.7.0.21_160000] 3.0, Intel(R) Corporation Computing - Device: Intel(R) Xeon(R) Platinum 8360Y CPU @ 2.40GHz - Device compute units: 144 - Device max work item size: 8192, 8192, 8192 - Device max work group size: 8192 + Device: Intel(R) Data Center GPU Max 1550 + Device compute units: 512 + Device max work item size: 1024, 1024, 1024 + Device max work group size: 1024 Problem size: c(150,600) = a(150,300) * b(300,600) ====ERROR: DeviceSanitizer: invalid-argument on kernel (auto&) const::'lambda'(auto)> - #0 in sycl::_V1::event sycl::_V1::queue::submit_with_event>>(sycl::_V1::ext::oneapi::experimental::properties>, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:3762:12 - #1 in std::enable_if, sycl::_V1::event>::type sycl::_V1::queue::submit(auto, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:429:12 - #2 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:106:7 - #3 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x7fddd2a5bd8f) - #4 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x7fddd2a5be3f) - #5 in _start (./1_matrix_mul_invalid_contexts_asan+0x403624) - - The 5th argument 0x7fddcf88a800 is allocated in other context - 0x7fddcf88a800 is located inside of Device USM region [0x7fddcf88a800, 0x7fddcf8e2640) + #0 in sycl::_V1::event sycl::_V1::queue::submit_with_event>>(sycl::_V1::ext::oneapi::experimental::properties>, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:3700:12 + #1 in std::enable_if, sycl::_V1::event>::type sycl::_V1::queue::submit(auto, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:441:12 + #2 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:126:7 + #3 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x70114642a1c9) + #4 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x70114642a28a) + #5 in _start (./1_matrix_mul_invalid_contexts_asan+0x403674) + + The 5th argument 0xff00ffffffe70800 is allocated in other device + 0xff00ffffffe70800 is located inside of Device USM region [0xff00ffffffe70800, 0xff00ffffffec8640) allocated here: - #0 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::device const&, sycl::_V1::context const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/usm.hpp:174:27 - #1 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::queue const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/usm.hpp:182:10 - #2 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:82:21 - #3 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x7fddd2a5bd8f) - #4 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x7fddd2a5be3f) - #5 in _start (./1_matrix_mul_invalid_contexts_asan+0x403624) + #0 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::device const&, sycl::_V1::context const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/usm.hpp:175:27 + #1 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::queue const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/usm.hpp:183:10 + #2 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:102:21 + #3 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x70114642a1c9) + #4 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x70114642a28a) + #5 in _start (./1_matrix_mul_invalid_contexts_asan+0x403674) Aborted (core dumped) ``` @@ -271,42 +269,44 @@ A recent addition to the oneAPI compiler is that ability to use the "Address San 3. Look at the reported source location - That Address Sanitizer threw an error at line 106 where we enter the main work kernel. + That Address Sanitizer threw an error at line 126 where we enter the main work kernel. ``` - 99 q.submit([&](auto &h) { - 100 h.memcpy(dev_c, &c_back[0], M*P * sizeof(float)); - 101 }); - 102 - 103 q.wait(); - 104 - 105 // Submit command group to queue to multiply matrices: c = a * b - 106 q.submit([&](auto &h) { // Something reported wrong here - 107 // Read from a and b, write to c - 108 int width_a = N; - 109 - 110 // Execute kernel. + 119 q.submit([&](auto &h) { + 120 h.memcpy(dev_c, &c_back[0], M*P * sizeof(float)); + 121 }); + 122 + 123 q.wait(); + 124 + 125 // Submit command group to queue to multiply matrices: c = a * b + 126 q.submit([&](auto &h) { + 127 // Read from a and b, write to c + 128 int width_a = N; + 129 + 130 // Execute kernel. ``` - And it complained that "the 5th argument" was allocated in another context, an element that was set up in line 82: + And it complained that "the 5th argument" was allocated in another context, an element that was set up in line 102: ``` - 77 float * dev_a = sycl::malloc_device(M*N, q); - 78 float * dev_b = sycl::malloc_device(N*P, q); - 79 device selected_device = device(default_selector_v); - 80 context devicecontext(selected_device); - 81 queue q2(devicecontext, selected_device); - 82 float * dev_c = sycl::malloc_device(M*P, q2); // Complaining about this - 83 - 84 cout << "Problem size: c(" << M << "," << P << ") = a(" << M << "," << N - 85 << ") * b(" << N << "," << P << ")\n"; + 93 float * dev_a = sycl::malloc_device(M*N, q); + 94 float * dev_b = sycl::malloc_device(N*P, q); + 95 #ifdef BAD_FREE + 96 device selected_device = devices[0]; + 97 #else + 98 device selected_device = devices[1]; + 99 #endif + 100 context devicecontext(selected_device); + 101 queue q2(devicecontext, selected_device); + 102 float * dev_c = sycl::malloc_device(M*P, q2); + 103 ``` - Unfortunately, figuring out which "5th argument" of which kernel is involved requires internal knowledge of Level Zero API calls and how SYCL collects the references used in the `q.submit` lambda function at line 106. We can get some of this by running `unitrace ` against `./1_matrix_mul_invalid_contexts_asan`, but the result is not satisfying. + Unfortunately, figuring out which "5th argument" of which kernel is involved requires internal knowledge of Level Zero API calls and how SYCL collects the references used in the `q.submit` lambda function at line 126. We can get some of this by running `unitrace ` against `./1_matrix_mul_invalid_contexts_asan`, but the result is not satisfying. - Also, it may not be useful to figure out which argument was involved, because this may not be where things first go bad. The SYCL syntax used in this program means that `q.submit` statements will immediately return control to the calling program even if the submitted kernel is still running. So the problem might have occurred before line 106, and we are only just now learning about it. + Also, it may not be useful to figure out which argument was involved, because this may not be where things first go bad. The SYCL syntax used in this program means that `q.submit` statements will immediately return control to the calling program even if the submitted kernel is still running. So the problem might have occurred before line 126, and we are only just now learning about it. - However, we don't need to go that deep into the internals. We are told that the invalid element encountered at line 106 was allocated to a different context in line 82. Here we see that `dev_c` is the variable, and it was allocated a size of `M*P` using queue `q2`. + However, we don't need to go that deep into the internals. We are told that the invalid element encountered at line 126 was allocated to a different device context in line 102. Here we see that `dev_c` is the variable, and it was allocated a size of `M*P` using queue `q2`. You've probably spotted the problem in this trivial example, but let's use the debugger to see if we can gather additional information to do a proper diagnosis. @@ -375,22 +375,21 @@ In case we need view code running on the GPU, we need to enable GPU debugging. #7 0x00007ffff78a5a55 in std::terminate() () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 #8 0x00007ffff78bb0c1 in std::rethrow_exception(std::__exception_ptr::exception_ptr) () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 - #9 0x00007ffff7ec9ef3 in sycl::_V1::detail::Scheduler::enqueueCommandForCG(std::shared_ptr, std::vector >&, sycl::_V1::detail::BlockingT) () from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #10 0x00007ffff7ec9695 in sycl::_V1::detail::Scheduler::addCG(std::unique_ptr >, std::shared_ptr const&, bool, ur_exp_command_buffer_handle_t_*, std::vector > const&) () - from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #11 0x00007ffff7efe432 in sycl::_V1::handler::finalize() () from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #12 0x00007ffff7e84277 in void sycl::_V1::detail::queue_impl::finalizeHandler(sycl::_V1::handler&, sycl::_V1::event&) () from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #13 0x00007ffff7e832b7 in sycl::_V1::detail::queue_impl::submit_impl(std::function const&, std::shared_ptr const&, std::shared_ptr const&, std::shared_ptr const&, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::SubmissionInfo const&) () from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #14 0x00007ffff7e895c8 in sycl::_V1::detail::queue_impl::submit_with_event(std::function const&, std::shared_ptr const&, sycl::_V1::detail::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) () from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #15 0x00007ffff7f33afa in sycl::_V1::queue::submit_with_event_impl(std::function, sycl::_V1::detail::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) () - from /opt/intel/oneapi/compiler/2025.1/lib/libsycl.so.8 - #16 0x00000000004049a3 in sycl::_V1::queue::submit_with_event(main::{lambda(auto:1&)#1}, sycl::_V1::queue*, sycl::_V1::detail::code_location const&) (this=0x7fffffffb778, CGF=..., SecondaryQueuePtr=0x0, - CodeLoc=...) at /opt/intel/oneapi/compiler/2025.1/bin/compiler/../../include/sycl/queue.hpp:2826 - #17 0x00000000004043bd in sycl::_V1::queue::submit(main::{lambda(auto:1&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb778, CGF=..., CodeLoc=...) - at /opt/intel/oneapi/compiler/2025.1/bin/compiler/../../include/sycl/queue.hpp:365 - #18 0x0000000000403f7b in main () - at 1_matrix_mul_invalid_contexts.cpp:99 - + #9 0x00007ffff7ee0f48 in sycl::_V1::detail::Scheduler::enqueueCommandForCG(sycl::_V1::detail::event_impl&, std::vector >&, sycl::_V1::detail::BlockingT) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #10 0x00007ffff7edff90 in sycl::_V1::detail::Scheduler::addCG(std::unique_ptr >, sycl::_V1::detail::queue_impl&, bool, ur_exp_command_buffer_handle_t_*, std::vector > const&) () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #11 0x00007ffff7f10267 in sycl::_V1::handler::finalize() () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #12 0x00007ffff7e93a38 in std::shared_ptr sycl::_V1::detail::queue_impl::finalizeHandlerOutOfOrder(sycl::_V1::handler&) () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #13 0x00007ffff7e92e04 in sycl::_V1::detail::queue_impl::submit_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::_V1::SubmissionInfo const&) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #14 0x00007ffff7f58ac9 in sycl::_V1::queue::submit_with_event_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::_V1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) const () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #15 0x0000000000407b4e in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb598, Props=..., CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:3700 + #16 0x00000000004048e1 in sycl::_V1::queue::submit(main::{lambda(auto:1&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb598, CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:441 + #17 0x00000000004043e8 in main () + at Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp:119 (gdb) ``` @@ -400,15 +399,15 @@ In case we need view code running on the GPU, we need to enable GPU debugging. 6. Switch the debugger focus to that frame. ``` - (gdb) frame 18 + (gdb) frame 17 ``` Your output will be similar to the following: ``` - #18 0x0000000000403f7b in main () - at Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp:99 - 99 q.submit([&](auto &h) { + #17 0x00000000004043e8 in main () + at Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp:119 + 119 q.submit([&](auto &h) { (gdb) ``` @@ -421,27 +420,27 @@ In case we need view code running on the GPU, we need to enable GPU debugging. You should see the code around the line reporting the problem. ``` - 94 - 95 // Submit command group to queue to initialize matrix b - 96 q.memcpy(dev_b, &b_back[0], N*P * sizeof(float)); - 97 - 98 // Submit command group to queue to initialize matrix c - 99 q.submit([&](auto &h) { - 100 h.memcpy(dev_c, &c_back[0], M*P * sizeof(float)); - 101 }); - 102 - 103 q.wait(); + 114 + 115 // Submit command group to queue to initialize matrix b + 116 q.memcpy(dev_b, &b_back[0], N*P * sizeof(float)); + 117 + 118 // Submit command group to queue to initialize matrix c + 119 q.submit([&](auto &h) { + 120 h.memcpy(dev_c, &c_back[0], M*P * sizeof(float)); + 121 }); + 122 + 123 q.wait(); ``` - Something wrong in line 99. Unfortunately, the + Something wrong in line 119. Unfortunately, the ` Enqueue process failed` message we saw when it crashed does not really tell us anything other than our attempt to submit the `memcpy` to the device failed - Fortunately, in this case the two variables, `dev_c` and `c_back`, are allocated only a few lines above line 99. In real code this might have happened in another source file or library, so hunting down this issue is going to be much harder. + Fortunately, in this case the two variables, `dev_c` and `c_back`, are allocated only a few lines above line 102. In real code this might have happened in another source file or library, so hunting down this issue is going to be much harder. Look at the source, and note that `dev_c` is defined as a pointer to device memory allocated on queue `q2` (we noticed this with Address Sanitizer): ``` - 82 float * dev_c = sycl::malloc_device(M*P, q2); + 102 float * dev_c = sycl::malloc_device(M*P, q2); ``` and `c_back` is defined as local memory @@ -450,15 +449,15 @@ In case we need view code running on the GPU, we need to enable GPU debugging. 49 float(*c_back)[P] = new float[M][P]; ``` -8. Look at line 99, and notice the discrepancy. +8. Look at line 119, and notice the discrepancy. ``` - 99 q.submit([&](auto &h) { + 119 q.submit([&](auto &h) { ``` - Variable `dev_c` was allocated on queue `q2` while the submit statement is being done on queue `q`. These queues are created using deferent devices contexts (the default one, and `devicecontext`) + Variable `dev_c` was allocated on queue `q2` while the submit statement is being done on queue `q`. These queues are created using different devices contexts (the default one, and `devicecontext`). - So unlike what Address Sanitizer suggested, the problem was first noticed in line 99, not 106. The debugger stopped immediately, while Address Sanitizer took a moment longer to spot the problem/respond to the exception. + So unlike what Address Sanitizer suggested, the problem was first noticed in line 119, not 126. The debugger stopped immediately, while Address Sanitizer took a moment longer to spot the problem/respond to the exception. ### Identify the Problem without Code Inspection @@ -472,20 +471,20 @@ is every low-level API call made to OpenCL™ or Level Zero. We will use it to a to enable call logging of API calls. > **Note**: You must modify the command shown below to include the path to - > where you installed the `unitrace` utility. + > where you installed the `unitrace` utility. We also recommend that you redirect the results to a file so that you can look at them in an editor via ` > results.txt 2>&1` ``` ONEAPI_DEVICE_SELECTOR=opencl:gpu [path]/unitrace -c --opencl ./1_matrix_mul_invalid_contexts ``` - The `unitrace` utility outputs extensive results, which will look similiar to the below. A few key excerpts with areas of interest are shown below (note - pulling these out is not easy without a program that can highlight multiple searches with different colors). + The `unitrace` utility outputs extensive results, which will look similar to the below. A few key excerpts with areas of interest are shown below (note - pulling these out is not easy without a program that can highlight multiple searches with different colors). ![alt text](./TraceOutput.png "unitrace -c output for OpenCL") - Let's work backwards from the error, starting with `clEnqueueMemcpyINTEL` (which corresponds to the `q.submit` at line 99 in the source). + Let's work backwards from the error, starting with `clEnqueueMemcpyINTEL` (which corresponds to the `q.submit` at line 119 in the source). The final `clEnqueueMemcpyINTEL` (line 22) uses `command_queue = 0x49d7130` and copies - `src_ptr` into device memory `dst_ptr = 0xff00ffffffeb0000` (NOTE: in some versions of `unitrace` these addresses may be returned in decimal rather than hexidecimal). Working back + `src_ptr` into device memory `dst_ptr = 0xff00ffffffeb0000` (NOTE: in some versions of `unitrace` these addresses may be returned in decimal rather than hexadecimal). Working back up the trace, you can see we allocated the destination device memory with the address `0xff00ffffffeb0000` using context `0x49dbff0` (line 16). However, the command queue (`0x49d7130`) being used in the `clEnqueueMemcpyINTEL` call @@ -497,19 +496,18 @@ is every low-level API call made to OpenCL™ or Level Zero. We will use it to a This *is legal* if both queues point to the same device context; however, in this example `q2` is actually defined pointing to another device context. You might do this in actual code if you have multiple offload compute - devices you are targeting. This code is sending work and data to each - device for processing. It is easy to accidentally send the wrong pointer to - the wrong queue in complex code. + devices you are targeting. It is easy to accidentally send the wrong pointer to + the wrong queue or device in complex code (as in this case). For comparison, an example of legal memory copy where the device context (`0x488d190`) used for the command queue (`0x49d7130`) is the same as that used for the memory allocation is shown as well (lines 4, 7, 19). -2. Let's also look at the output from Level Zero, and see if we could have +2. Let's also look at the output from Level Zero (the default GPU runtime), and see if we could have detected the issue there: ``` - ONEAPI_DEVICE_SELECTOR=level_zero:gpu [path]unitrace -c ./1_matrix_mul_invalid_contexts + [path]unitrace -c ./1_matrix_mul_invalid_contexts ``` Your output might be similar to the following: @@ -531,8 +529,10 @@ is every low-level API call made to OpenCL™ or Level Zero. We will use it to a timeline showing "what happened when" using the Chrome(tm) tracing browser tool. -### A final thought. -If you look at the source, `1_matrix_mul_invalid_contexts.cpp`, you will see that both queues where created with the `default_selector_v` device selector to make this tutorial work for most people (lines 66 and 79). If these were set to use different device selectors (say `gpu_selector_v` and `cpu_selector_v` or a custom device), or we specified exactly which device to use for each queue, our ability to use `ONEAPI_DEVICE_SELECTOR` to explore the problem might have been handicapped in this sample. But we hope it is clear that this code has an error in it in the case where two entirely different devices are selected. +### Cautionary thought. +Throughout this sample we made use of `ONEAPI_DEVICE_SELECTOR` to explore this problem. If it worked for you, it was because we were lucky - there were multiple instances of the same device, so we were able to explicitly create SYCL queues on two different devices of the same type. If not, we still we hope it is clear that this code has an error in it in the case where two entirely different devices are selected. + +If the above example did not work for you, we hope that you will try the bonus example below which shows another bug arising from using the wrong device context. For more on device discovery in SYCL, please see https://www.intel.com/content/www/us/en/developer/articles/technical/device-discovery-with-sycl.html. @@ -540,11 +540,184 @@ For more on device discovery in SYCL, please see https://www.intel.com/content/w To fix this problem, you must change the `malloc_device` allocating `dev_c` to use the same queue (and thus same device context) as the first two device -allocations or create queue `q` to use the same underlying device context. +allocations or create queue `q2` in such a way that it uses the same underlying device context. -If you really need to operate over multiple devices, this example will need to +If you really need to operate over multiple devices, exchanging data between devices, this example will need to be entirely re-written, which is beyond the scope of this tutorial. +### Bonus example + +1. Build some new binaries from our demo program: + ``` + icpx -fsycl -O0 -g -std=gnu++17 -Rno-debug-disables-optimization -D BAD_FREE -o 1_matrix_mul_invalid_contexts_badfree ../src/1_matrix_mul_invalid_contexts.cpp + + icpx -fsycl -O0 -g -Xarch_device -fsanitize=address -std=gnu++17 -Rno-debug-disables-optimization -D BAD_FREE -o 1_matrix_mul_invalid_contexts_asan_badfree ../src/1_matrix_mul_invalid_contexts.cpp + ``` + +2. If you run `1_matrix_mul_invalid_contexts_asan_badfree`, it should run without errors. + +3. Recall that SYCL applications use the Level Zero runtime by default with an Intel GPU. But what happens if you use OpenCL™ runtime to run 1_matrix_mul_invalid_contexts on the GPU? + + ``` + ONEAPI_DEVICE_SELECTOR=opencl:gpu ./1_matrix_mul_invalid_contexts_badfree + ``` + The output might look similar to the following: + ``` + ONEAPI_DEVICE_SELECTOR=opencl:gpu ./1_matrix_mul_invalid_contexts_badfree + Initializing + Devices: + [0] Intel(R) Data Center GPU Max 1100, OpenCL 3.0 NEO [25.18.33578] 3.0, Intel(R) Corporation + Computing + Device: Intel(R) Data Center GPU Max 1100 + Device compute units: 448 + Device max work item size: 1024, 1024, 1024 + Device max work group size: 1024 + Problem size: c(150,600) = a(150,300) * b(300,600) + terminate called after throwing an instance of 'sycl::_V1::exception' + what(): Enqueue process failed. + opencl backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) + Aborted (core dumped) + ``` + > **Note**: you will only see the above if the sycl-ls command shows OpenCL devices for the graphics card. If you are missing [opencl:gpu] devices you may have to add the necessary libraries to your device path by setting the appropriate path in DRIVERLOC and then running the following four commands (for Ubuntu - adapt for other OSes): + ``` + export DRIVERLOC=/usr/lib/x86_64-linux-gnu + export OCL_ICD_FILENAMES=$OCL_ICD_FILENAMES:$DRIVERLOC/intel-opencl/libigdrcl.so + export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$DRIVERLOC + export PATH=$PATH:/opt/intel/oneapi:$DRIVERLOC + ``` + Then issue the `ONEAPI_DEVICE_SELECTOR=opencl:gpu ./1_matrix_mul_invalid_contexts_badfree` command again to see the above error. + +4. However, if you run 1_matrix_mul_invalid_contexts_asan_badfree, you will see the following: + + ``` + ./1_matrix_mul_invalid_contexts_asan_badfree + Initializing + ==== DeviceSanitizer: ASAN + Devices: + [0] Intel(R) Data Center GPU Max 1100, 12.60.7 [1.6.33578+51] 1.6.33578+51, Intel(R) Corporation + [1] Intel(R) Xeon(R) Platinum 8480+, OpenCL 3.0 (Build 0) [2026.21.3.0.31_160000] 3.0, Intel(R) Corporation + [2] Intel(R) Xeon(R) Platinum 8480+, OpenCL 3.0 (Build 0) [2023.16.7.0.21_160000] 3.0, Intel(R) Corporation + Computing + Device: Intel(R) Data Center GPU Max 1100 + Device compute units: 448 + Device max work item size: 1024, 1024, 1024 + Device max work group size: 1024 + Problem size: c(150,600) = a(150,300) * b(300,600) + + ====ERROR: DeviceSanitizer: bad-context on address 0xff00ffffffe80800 + #0 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:157:5 + #1 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x795381e2a1c9) + #2 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x795381e2a28a) + #3 in _start (./1_matrix_mul_invalid_contexts_asan_badfree+0x403674) + + 0xff00ffffffe80800 is located inside of Device USM region [0xff00ffffffe80800, 0xff00ffffffed8640) + allocated here: + #0 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::device const&, sycl::_V1::context const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/usm.hpp:175:27 + #1 in float* sycl::_V1::malloc_device(unsigned long, sycl::_V1::queue const&, sycl::_V1::property_list const&, sycl::_V1::detail::code_location const&) /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/usm.hpp:183:10 + #2 in main Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/build/../src/1_matrix_mul_invalid_contexts.cpp:102:21 + #3 in ?? (/lib/x86_64-linux-gnu/libc.so.6+0x795381e2a1c9) + #4 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x795381e2a28a) + #5 in _start (./1_matrix_mul_invalid_contexts_asan_badfree+0x403674) + + Aborted (core dumped) + ``` + +5. Line 157 says this: + + ``` + 157 sycl::free(dev_c, q); + ``` + And here is where the memory was allocated: + ``` + 100 context devicecontext(selected_device); + 101 queue q2(devicecontext, selected_device); + 102 float * dev_c = sycl::malloc_device(M*P, q2); + ``` + As is hopefully obvious from the previous example, the problem is that we are trying to free memory allocated in SYCL queue `q2` that has a different device context fron SYCL queue `q`, even though under the covers they point to the same hardware device. + +6. While we are at it, what does the debugger show us? + ``` + ONEAPI_DEVICE_SELECTOR=opencl:gpu ZET_ENABLE_PROGRAM_DEBUGGING=1 gdb-oneapi ./1_matrix_mul_invalid_contexts_badfree + ``` + And then + ``` + (gdb) run + ``` + We crash like before: + ``` + Initializing + Devices: + [0] Intel(R) Data Center GPU Max 1100, OpenCL 3.0 NEO [25.18.33578] 3.0, Intel(R) Corporation + Computing + Device: Intel(R) Data Center GPU Max 1100 + Device compute units: 448 + Device max work item size: 1024, 1024, 1024 + Device max work group size: 1024 + Problem size: c(150,600) = a(150,300) * b(300,600) + [New Thread 0x7fffcae006c0 (LWP 2084018)] + terminate called after throwing an instance of 'sycl::_V1::exception' + what(): Enqueue process failed. + opencl backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) + + Thread 1 "1_matrix_mul_in" received signal SIGABRT, Aborted. + __pthread_kill_implementation (no_tid=0, signo=6, threadid=) at ./nptl/pthread_kill.c:44 + ⚠️ warning: 44 ./nptl/pthread_kill.c: No such file or directory + ``` +7. Looking at the backtrace + ``` + (gdb) backtrace + #0 __pthread_kill_implementation (no_tid=0, signo=6, threadid=) at ./nptl/pthread_kill.c:44 + #1 __pthread_kill_internal (signo=6, threadid=) at ./nptl/pthread_kill.c:78 + #2 __GI___pthread_kill (threadid=, signo=signo@entry=6) at ./nptl/pthread_kill.c:89 + #3 0x00007ffff764527e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26 + #4 0x00007ffff76288ff in __GI_abort () at ./stdlib/abort.c:79 + #5 0x00007ffff7ca5ff5 in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 + #6 0x00007ffff7cbb0da in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 + #7 0x00007ffff7ca5a55 in std::terminate() () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 + #8 0x00007ffff7cbb0c1 in std::rethrow_exception(std::__exception_ptr::exception_ptr) () + from /usr/lib/x86_64-linux-gnu/libstdc++.so.6 + #9 0x00007ffff7b23f48 in sycl::_V1::detail::Scheduler::enqueueCommandForCG(sycl::_V1::detail::event_impl&, std::vector >&, sycl::_V1::detail::BlockingT) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #10 0x00007ffff7b22f90 in sycl::_V1::detail::Scheduler::addCG(std::unique_ptr >, sycl::_V1::detail::queue_impl&, bool, ur_exp_command_buffer_handle_t_*, std::vector > const&) () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #11 0x00007ffff7b53267 in sycl::_V1::handler::finalize() () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #12 0x00007ffff7ad6a38 in std::shared_ptr sycl::_V1::detail::queue_impl::finalizeHandlerOutOfOrder(sycl::_V1::handler&) () from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #13 0x00007ffff7ad5e04 in sycl::_V1::detail::queue_impl::submit_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, bool, sycl::_V1::detail::code_location const&, bool, sycl::_V1::detail::_V1::SubmissionInfo const&) () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #14 0x00007ffff7b9bac9 in sycl::_V1::queue::submit_with_event_impl(sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::_V1::SubmissionInfo const&, sycl::_V1::detail::code_location const&, bool) const () + from /opt/intel/oneapi/compiler/2026.0/lib/libsycl.so.9 + #15 0x0000000000407b3e in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb588, Props=..., CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:3700 + #16 0x00000000004048d1 in sycl::_V1::queue::submit(main::{lambda(auto:1&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb588, CGF=..., CodeLoc=...) + at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:441 + #17 0x00000000004043de in main () at ../src/1_matrix_mul_invalid_contexts.cpp:119 + (gdb) + ``` + + We are back to the error above. Recall that line line 119 is doing this: + + ``` + 119 q.submit([&](auto &h) { + 120 h.memcpy(dev_c, &c_back[0], M*P * sizeof(float)); + 121 }); + ``` + + Variable `dev_c` was allocated on queue `q2` while the submit statement is being done on queue `q`. These queues are created using different devices contexts. Now, they point to the same device, but the device context is different. + + Interestingly, Address Sanitizer was fine with that. What it objected to (running under Level Zero rather than OpenCL here) was trying to free memory allocated in one SYCL device context in a different SYCL device context. + + Even more interestingly, if you run this program in the debugger under level zero like so: + ``` + ONEAPI_DEVICE_SELECTOR=level_zero:gpu ZET_ENABLE_PROGRAM_DEBUGGING=1 gdb-oneapi ./1_matrix_mul_invalid_contexts_badfree + ``` + The program will complete without errors. + + Both of these tools are correct that there is an error here. Level Zero gave the impression (including in the debugger) that everything was fine, while Address Sanitizer caught the bad free in Level Zero. Unfortunately Address Sanitizer does not run under OpenCL. Still, OpenCL itself helped identify the other issue (using queues belonging to different contexts interchangeably) because, unlike Level Zero, it does not allow this. + +### Final Thoughts +Both of these examples are cases of subtle coding errors that might work well in some environments, but not in others. When SYCL queues and devices start getting passed around between functions, or libraries, along with memory allocated using those queues, it can be easy to lose track of which queue owns a given buffer, and subtle errors can start to creep in. + +Moral of the story: check your program from time-time on different GPU runtimes (Level Zero and OpenCL), and with different tools, in the hopes of exposing these issues before they come too deeply imbedded and hard to track down. + ## License Code samples are licensed under the MIT license. See diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp index dc505e0c2f..9810af1bbd 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp @@ -63,7 +63,23 @@ int main() { // Initialize the device queue with the default selector. The device queue is // used to enqueue kernels. It encapsulates all states needed for execution. { - queue q(default_selector_v); + property_list propList = property_list{property::queue::enable_profiling()}; + + std::vector devices = sycl::device::get_devices(); + cout << "Devices:" << std::endl; + + for (size_t index = 0; index < devices.size(); index++){ + std::string device_name = devices[index].get_info(); + std::string device_driver = devices[index].get_info(); + std::string sycl_version = devices[index].get_info(); + std::string vendor = devices[index].get_info(); + std::string backend = devices[index].get_info(); + std::cout << " [" << index << "] " << device_name << ", " << sycl_version << " [" << device_driver + << "] " << backend << ", " << vendor << std::endl; + } + + // Be very specific about the device to use. + queue q(devices[0]); cout << "Computing" << "\n"; cout << "Device: " << q.get_device().get_info() << "\n"; @@ -76,7 +92,11 @@ int main() { // Create 2D buffers for matrices, buffer c is bound with host memory c_back float * dev_a = sycl::malloc_device(M*N, q); float * dev_b = sycl::malloc_device(N*P, q); - device selected_device = device(default_selector_v); +#ifdef BAD_FREE + device selected_device = devices[0]; +#else + device selected_device = devices[1]; +#endif context devicecontext(selected_device); queue q2(devicecontext, selected_device); float * dev_c = sycl::malloc_device(M*P, q2); @@ -131,6 +151,13 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + sycl::free(dev_a, q); + sycl::free(dev_b, q); +#ifdef BAD_FREE + sycl::free(dev_c, q); +#else + sycl::free(dev_c, q2); +#endif } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/2_matrix_mul.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/2_matrix_mul.cpp index ac97b65b23..0883604faa 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/2_matrix_mul.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/2_matrix_mul.cpp @@ -65,6 +65,19 @@ int main() { { property_list propList = property_list{property::queue::enable_profiling()}; + std::vector devices = sycl::device::get_devices(); + cout << "Devices:" << std::endl; + + for (size_t index = 0; index < devices.size(); index++){ + std::string device_name = devices[index].get_info(); + std::string device_driver = devices[index].get_info(); + std::string sycl_version = devices[index].get_info(); + std::string vendor = devices[index].get_info(); + std::string backend = devices[index].get_info(); + std::cout << " [" << index << "] " << device_name << ", " << sycl_version << " [" << device_driver + << "] " << backend << ", " << vendor << std::endl; + } + queue q(default_selector_v); cout << "Computing" << "\n"; @@ -130,6 +143,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md index 2bce82e253..5583300434 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md @@ -1,6 +1,6 @@ # `Guided Matrix Multiplication Race Condition` Sample -The `Guided Matrix Multiplication Race Condition` sample demonstrates a guided approach to debugging a race condition accessing data on the host before it has been fully copied back from the device. It uses the Intel® oneAPI Base Toolkit (Base Kit) and several tools included in the Base Kit. +The `Guided Matrix Multiplication Race Condition` sample demonstrates an approach to debugging a race condition using several tools in Intel® oneAPI. The race condition arises from accessing data on the host before it has been fully copied back from the device. The sample is a simple program that multiplies together two large matrices and verifies the results. @@ -22,7 +22,7 @@ The sample includes different versions of a simple matrix multiplication program | File name | Description |:--- |:--- | `1_matrix_mul_race_condition.cpp` |This example shows what happens when a developer tries to access data provided by the device before the copy to the host is complete. -| `2_matrix_mul.cpp` | A working version of the matrix multiply code that properly waits for the data to be copied back to the host. +| `2_matrix_mul.cpp` | A working version of the matrix multiply code that properly waits for the data to be copied back to the host but still has some issues. | `3_matrix_mul.cpp` | A working version of the application that corrects its errors using a host accessor and a `q.wait` command in place of parenthesis. ## Prerequisites @@ -30,9 +30,9 @@ The sample includes different versions of a simple matrix multiplication program | Optimized for | Description |:--- |:--- | OS | Ubuntu* 24.04 LTS -| Hardware | GEN9 or newer -| Software | Intel® oneAPI DPC++/C++ Compiler 2025.3
Intel® Distribution for GDB* 2025.3
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). -| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.31 or later from https://dgpu-docs.intel.com/releases/releases.html +| Intel GraphicsHardware | GEN9 or newer +| Software | Intel® oneAPI DPC++/C++ Compiler 2026.0
Intel® Distribution for GDB* 2026.0
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). +| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.59 or later from https://dgpu-docs.intel.com/releases/releases.html ## Key Implementation Details @@ -44,7 +44,7 @@ The basic SYCL* standards implemented in the code include the use of the followi ## Set Environment Variables -When working with the command-line interface (CLI), you should configure the oneAPI toolkits using environment variables. Set up your CLI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries, and tools are ready for development. +When working with the command-line interface (CLI), set up your oneAPI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. ## Build and Run the `Guided Matrix Multiply Race Condition` Programs @@ -115,8 +115,6 @@ the `VERBOSE=1` argument: make VERBOSE=1 ``` -If you receive an error message, troubleshoot the problem using the **Diagnostics Utility for Intel® oneAPI Toolkits**. The diagnostic utility provides configuration and system checks to help find missing dependencies, permissions errors, and other issues. See the *[Diagnostics Utility for Intel® oneAPI Toolkits User Guide](https://www.intel.com/content/www/us/en/docs/oneapi/user-guide-diagnostic-utility/current/overview.html)* for more information on using the utility. - ## Guided Debugging @@ -125,7 +123,7 @@ This example shows what happens when code tries to access data provided by the d These instructions assume you have installed the Intel® Distribution for GDB* and have a basic working knowledge of GDB. ### Setting up to Debug on the GPU -To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GDB-PVC debugger](https://dgpu-docs.intel.com/system-user-guides/DNP-Max-1100-userguide/DNP-Max-1100-userguide.html#gdb-pvc-debugger)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. +To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GPU Debugging](https://dgpu-docs.intel.com/driver/gpu-debugging.html)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. Documentation on using the debugger in a variety of situations can be found at *[Debug Examples in Linux](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/tutorial-debugging-dpcpp-linux/current/overview.html)* @@ -133,13 +131,13 @@ Documentation on using the debugger in a variety of situations can be found at * ### Getting the Tracing and Profiling Tool -In this tutorial, the instructions require a utility that was not installed with the Intel® oneAPI Base Toolkit (Base Kit). +In this tutorial, the instructions require a utility that was not installed with Intel® oneAPI. To complete the steps in the following section, you must download the [Unified Tracing and Profiling Tool](https://github.com/intel/pti-gpu/tree/master/tools/unitrace) code from GitHub and build the utility. The build instructions are included in the README in the GitHub repository. This build will go much more smoothly if you first install the latest drivers from [the Intel GPU driver download site](https://dgpu-docs.intel.com/driver/overview.html), especially the development packages (only available in the Data Center GPU driver install). Once you have built the utility, you invoke it on the command line in front of your program (similar to using GDB). ### Examine the Original Code -As you might have noticed, when you attempt to run `1_matrix_mul_race_condition.cpp` the code reports bad results and then exits. We can use the Intel® Distribution for GDB* to get a backtrace of the entire stack to understand the problem. +As you might have noticed, when you attempt to run `1_matrix_mul_race_condition` the code reports bad results and then exits. We can use the Intel® Distribution for GDB* to get a backtrace of the entire stack to understand the problem. In case we need view code running on the GPU, we need to enable GPU debugging. This will require [some setup on your system](#setting-up-to-debug-on-the-gpu) before you can see code running on the GPU. @@ -149,7 +147,7 @@ In case we need view code running on the GPU, we need to enable GPU debugging. ``` 2. Then run the application in the debugger. ``` - run + (gdb) run ``` 3. Examine the results. ``` @@ -276,13 +274,22 @@ In case we need view code running on the GPU, we need to enable GPU debugging. Fail - The result is incorrect for element: [0, 3], expected: 45150, but found: 0 Fail - The result is incorrect for element: [0, 4], expected: 45150, but found: 0 Fail - The results mismatch! - >>>> [211022976795378] zeEventCreate: hEventPool = 37983040 desc = 140734297023408 {ZE_STRUCTURE_TYPE_EVENT_DESC(0x11) 0 4 4 0} phEvent = 140734297023464 (hEvent = 15669694584003) - <<<< [211022976802066] zeEventCreate [1339 ns] hEvent = 37883384 -> ZE_RESULT_SUCCESS(0x0) - >>>> [211022976805342] zeCommandListAppendMemoryCopyRegion: hCommandList = 37277672 dstptr = 35936816 dstRegion = 140734297023744 dstPitch = 2400 dstSlicePitch = 360000 srcptr = 18374967954634571776 srcRegion = 140734297023768 srcPitch = 2400 srcSlicePitch = 360000 hSignalEvent = 37883384 numWaitEvents = 1 phWaitEvents = 37993232 (hWaitEvents = [37882840]) - <<<< [211022976856153] zeCommandListAppendMemoryCopyRegion [46127 ns] hWaitEvents = 37882840 -> ZE_RESULT_SUCCESS(0x0) - >>>> [211022976862571] zeEventHostSynchronize: hEvent = 37883384 timeout = 18446744073709551615 - <<<< [211022979801501] zeEventHostSynchronize [2937354 ns] -> ZE_RESULT_SUCCESS(0x0) - Segmentation fault (core dumped) + >>>> [670752705693142] zeEventCreate: hEventPool = 0x48187f8 desc = 0x7ffea023b540 {ZE_STRUCTURE_TYPE_EVENT_DESC(0x11) 0 4 4 0} phEvent = 0x7ffea023b578 (hEvent = 0x0) + <<<< [670752705701414] zeEventCreate [2045 ns] hEvent = 0x4935aa8 -> ZE_RESULT_SUCCESS(0x0) + >>>> [670752705704779] zeCommandListAppendMemoryCopyRegion: hCommandList = 0x4917708 dstptr = 0x5b103e38010 dstRegion = 0x7ffea023b740 dstPitch = 0x960 dstSlicePitch = 0x57e40 srcptr = 0xff00ffffff2e0000 srcRegion = 0x7ffea023b768 srcPitch = 0x960 srcSlicePitch = 0x57e40 hSignalEvent = 0x4935aa8 numWaitEvents = 0x1 phWaitEvents = 0x4935390 (hWaitEvents = [0x49353d8]) + <<<< [670752705716520] zeCommandListAppendMemoryCopyRegion [7285 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195) + >>>> [670752705801127] zeEventQueryStatus: hEvent = 0x49353d8 + <<<< [670752705806038] zeEventQueryStatus [2103 ns] -> ZE_RESULT_SUCCESS(0x0) + >>>> [670752705809927] zeEventHostSynchronize: hEvent = 0x49353d8 timeout = 0xffffffffffffffff + <<<< [670752706994834] zeEventHostSynchronize [1182409 ns] -> ZE_RESULT_SUCCESS(0x0) + >>>> [670752707005758] zeEventQueryStatus: hEvent = 0x4935aa8 + <<<< [670752707010312] zeEventQueryStatus [1742 ns] -> ZE_RESULT_NOT_READY(0x1) + >>>> [670752707014922] zeEventQueryStatus: hEvent = 0x4935aa8 + <<<< [670752707018218] zeEventQueryStatus [1273 ns] -> ZE_RESULT_NOT_READY(0x1) + >>>> [670752707038683] zeEventQueryStatus: hEvent = 0x4935aa8 + <<<< [670752707042673] zeEventQueryStatus [1348 ns] -> ZE_RESULT_NOT_READY(0x1) + >>>> [670752707045585] zeEventQueryStatus: hEvent = 0x4935aa8 + <<<< [670752707048684] zeEventQueryStatus [1206 ns] -> ZE_RESULT_NOT_READY(0x1) ``` ### Interpret the Results @@ -291,7 +298,7 @@ The first clue here is that the program throws an exception *after* it has compl Next, looking at the crash in the debugger, there are a couple of odd things that stand out. Look at stack `frame 9`. This frame shows us attempting to update the host memory from the device, while `frame 20` shows we are already at the end of the program and have started cleaning up the SYCL buffers (`frame 19`). The only variable containing data returned from the device is `c_back`. But the developer has already deleted `c_back` in line 126, so the *data the buffer being copied into (`c_back`) no longer exists*. -We see something like this in the `unitrace` output above. The kernel is executed, the results are immediately checked, we create and wait on some events, and then the last thing we try to do before crashing is to copy some memory from the device memory (`srcptr = 18374967954634571776`) to a host pointer (`dstptr = 35936816`) that previously was used to initialize this same device memory (around line 101). Since `c_buf` is the only accessor that is defined as writeable in the `q.submit` at line 97, it again is a likely suspect. +We see something like this in the `unitrace` output above. The kernel is executed, the results are immediately checked, we create and wait on some events, and then the last thing we try to do before crashing/timing out is to copy some memory from the device memory (`srcptr = 0xff00ffffff2e0000`) to a host pointer (`dstptr = 0x5b103e38010`) that previously was used to initialize this same device memory (around line 101). Since `c_buf` is the only accessor that is defined as writeable in the `q.submit` at line 97, it again is a likely suspect. But what if the developer didn't delete `c_back`, and let program termination clean it up? Try it! diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md index 79e5dc4c03..c86dbdbbf0 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md @@ -1,6 +1,6 @@ # `Guided Matrix Multiplication Illegal SLM Size` Sample -The `Guided Matrix Multiplication Illegal SLM Size` sample demonstrates a guided approach to debugging incorrect use of the SYCL language. The sample uses the Intel® oneAPI Base Toolkit (Base Kit) and several tools included in the Base Kit. +The `Guided Matrix Multiplication Illegal SLM Size` sample demonstrates an approach to debugging incorrect use of the SYCL language using several tools in Intel® oneAPI. The sample is a simple program that multiplies together two large matrices and verifies the results. @@ -30,9 +30,9 @@ The sample includes different versions of a simple matrix multiplication program | Optimized for | Description |:--- |:--- | OS | Ubuntu* 24.04 LTS -| Hardware | GEN9 or newer -| Software | Intel® oneAPI DPC++/C++ Compiler 2025.3
Intel® Distribution for GDB* 2025.3
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). -| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.31 or later from https://dgpu-docs.intel.com/releases/releases.html +| Intel Graphics Hardware | GEN9 or newer +| Software | Intel® oneAPI DPC++/C++ Compiler 2026.0
Intel® Distribution for GDB* 2026.0
Unified Tracing and Profiling Tool 2.3.0, which is available from the [following Github repository](https://github.com/intel/pti-gpu/tree/master/tools/unitrace). +| Intel GPU Driver | Intel® General-Purpose GPU Long-Term Support driver 2523.59 or later from https://dgpu-docs.intel.com/releases/releases.html ## Key Implementation Details @@ -49,7 +49,7 @@ This can be particularly painful. For example, you might experience this error i ## Set Environment Variables -When working with the command-line interface (CLI), you should configure the oneAPI toolkits using environment variables. Set up your CLI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries, and tools are ready for development. +When working with the command-line interface (CLI), set up your oneAPI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries and tools are ready for development. ## Build and Run the `Guided Matrix Multiplication Illegal SLM Size` Sample @@ -117,15 +117,13 @@ the `VERBOSE=1` argument: make VERBOSE=1 ``` -If you receive an error message, troubleshoot the problem using the **Diagnostics Utility for Intel® oneAPI Toolkits**. The diagnostic utility provides configuration and system checks to help find missing dependencies, permissions errors, and other issues. See the *[Diagnostics Utility for Intel® oneAPI Toolkits User Guide](https://www.intel.com/content/www/us/en/docs/oneapi/user-guide-diagnostic-utility/current/overview.html)* for more information on using the utility. - ## Guided Debugging These instructions assume you have installed the Intel® Distribution for GDB* and have a basic working knowledge of GDB. ### Setting up to Debug on the GPU -To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GDB-PVC debugger](https://dgpu-docs.intel.com/system-user-guides/DNP-Max-1100-userguide/DNP-Max-1100-userguide.html#gdb-pvc-debugger)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. +To learn how setup and use Intel® Distribution for GDB*, see the *[Get Started with Intel® Distribution for GDB* on Linux* OS Host](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/get-started-guide-linux/current/overview.html)*. Additional setup instructions you should follow are at *[GPU Debugging](https://dgpu-docs.intel.com/driver/gpu-debugging.html)* and *[Configuring Kernel Boot Parameters](https://dgpu-docs.intel.com/driver/configuring-kernel-boot-parameters.html)*. Documentation on using the debugger in a variety of situations can be found at *[Debug Examples in Linux](https://www.intel.com/content/www/us/en/docs/distribution-for-gdb/tutorial-debugging-dpcpp-linux/current/overview.html)* @@ -134,7 +132,7 @@ Documentation on using the debugger in a variety of situations can be found at * ### Getting the Tracing and Profiling Tool -In this tutorial, the instructions require a utility that was not installed with the Intel® oneAPI Base Toolkit (Base Kit). +In this tutorial, the instructions require a utility that was not installed with Intel® oneAPI. To complete the steps in the following section, you must download the [Unified Tracing and Profiling Tool](https://github.com/intel/pti-gpu/tree/master/tools/unitrace) code from GitHub and build the utility. The build instructions are included in the README in the GitHub repository. This build will go much more smoothly if you first install the latest drivers from [the Intel GPU driver download site](https://dgpu-docs.intel.com/driver/overview.html), especially the development packages (only available in the Data Center GPU driver install). Once you have built the utility, you invoke it on the command line in front of your program (similar to using GDB). @@ -180,11 +178,14 @@ In `1_matrix_mul_SLM_size`, the local_accessor class is used to reserve an illeg The application will fail and display the same message when we ran it outside of the debugger. ``` : - Problem size: c(150,600) = a(150,300) * b(300,600) + Continuing with GPU-debugging disabled. + terminate called after throwing an instance of 'sycl::_V1::exception' - what(): level_zero backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) + what(): level_zero backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) Thread 1.1 "1_matrix_mul_SL" received signal SIGABRT, Aborted. + __pthread_kill_implementation (no_tid=0, signo=6, threadid=) at ./nptl/pthread_kill.c:44 + ⚠️ warning: 44 ./nptl/pthread_kill.c: No such file or directory (gdb) ``` @@ -210,18 +211,17 @@ In `1_matrix_mul_SLM_size`, the local_accessor class is used to reserve an illeg #10 0x00007ffff7de164a in sycl::_V1::detail::enqueue_kernel_launch::handleErrorOrWarning(ur_result_t, sycl::_V1::detail::device_impl const&, ur_kernel_handle_t_*, sycl::_V1::detail::NDRDescT const&) () from /opt/intel/oneapi/compiler/2025.3/lib/libsycl.so.8 : - #19 0x0000000000407e62 in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb580, Props=..., CGF=..., CodeLoc=...) - at /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:3762 - #20 0x00000000004043e9 in sycl::_V1::queue::submit(main::{lambda(sycl::_V1::handler&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb580, CGF=..., CodeLoc=...) - at /opt/intel/oneapi/compiler/2025.3/bin/compiler/../../include/sycl/queue.hpp:429 - #21 0x0000000000403fa3 in main () + #18 0x0000000000407d5e in sycl::_V1::queue::submit_with_event > >(sycl::_V1::ext::oneapi::experimental::properties >, sycl::_V1::detail::type_erased_cgfo_ty const&, sycl::_V1::detail::code_location const&) const (this=0x7fffffffb680, Props=..., CGF=..., + CodeLoc=...) at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:3700 + #19 0x0000000000404419 in sycl::_V1::queue::submit(main::{lambda(sycl::_V1::handler&)#1}, sycl::_V1::detail::code_location const&) (this=0x7fffffffb680, CGF=..., CodeLoc=...) at /opt/intel/oneapi/compiler/2026.0/bin/compiler/../../include/sycl/queue.hpp:441 + #20 0x0000000000403f73 in main () at Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp:104 ``` 4. Look at the final frame. (Your frame number might differ, and you might have to repeat this command to get the frame to change). ``` - (gdb) frame 21 - #21 0x0000000000403fe3 in main () at Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp:104 + (gdb) frame 20 + #20 0x0000000000403fe3 in main () at Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp:104 104 q.submit([&](handler &h){ (gdb) ``` @@ -264,14 +264,14 @@ Among other things, the Tracing and Profiling utility can print every low-level 3. Let the output continue until the error occurs and the program stops. ``` : - >>>> [806881439215102] zeKernelSetGroupSize: hKernel = 0x2a54be8 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 - <<<< [806881439220367] zeKernelSetGroupSize [1631 ns] -> ZE_RESULT_SUCCESS(0x0) - >>>> [806881439225163] zeCommandListCreateImmediate: hContext = 0x29cd578 hDevice = 0x2936ee8 altdesc = 0x7ffcce631850 {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC(0xe) 0 0 0 0 2 0} phCommandList = 0x7ffcce631838 (hCommandList = 0x0) - <<<< [806881439391714] zeCommandListCreateImmediate [160772 ns] hCommandList = 0x2d3b218 -> ZE_RESULT_SUCCESS(0x0) - >>>> [806881439399881] zeEventHostReset: hEvent = 0x2cd8a88 - <<<< [806881439403570] zeEventHostReset [1577 ns] -> ZE_RESULT_SUCCESS(0x0) - >>>> [806881439411094] zeCommandListAppendLaunchKernel: hCommandList = 0x2d3b218 hKernel = 0x2a54be8 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7ffcce631d78 {16385, 1, 1} hSignalEvent = 0x2cd8a88 numWaitEvents = 0x0 phWaitEvents = 0x0 - <<<< [806881439467224] zeCommandListAppendLaunchKernel [47029 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195) + >>>> [672116802394124] zeKernelSetGroupSize: hKernel = 0x2689658 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 + <<<< [672116802398592] zeKernelSetGroupSize [1577 ns] -> ZE_RESULT_SUCCESS(0x0) + >>>> [672116802403366] zeCommandListCreateImmediate: hContext = 0x260be58 hDevice = 0x25752c8 altdesc = 0x7fffb216c300 {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC(0xe) 0 0 0 0 2 0} phCommandList = 0x7fffb216c2e8 (hCommandList = 0x0) + <<<< [672116802564843] zeCommandListCreateImmediate [156721 ns] hCommandList = 0x3d56268 -> ZE_RESULT_SUCCESS(0x0) + >>>> [672116802577317] zeEventHostReset: hEvent = 0x3b585a8 + <<<< [672116802581433] zeEventHostReset [1546 ns] -> ZE_RESULT_SUCCESS(0x0) + >>>> [672116802588167] zeCommandListAppendLaunchKernel: hCommandList = 0x3d56268 hKernel = 0x2689658 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7fffb216c840 {16385, 1, 1} hSignalEvent = 0x3b585a8 numWaitEvents = 0x0 phWaitEvents = 0x0 + <<<< [672116802646429] zeCommandListAppendLaunchKernel [49631 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195) terminate called after throwing an instance of 'sycl::_V1::exception' what(): level_zero backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) Aborted (core dumped) @@ -282,12 +282,12 @@ Among other things, the Tracing and Profiling utility can print every low-level A note about the output above. You will see that is has two lines that read: ``` - >>>> [806881439215102] zeKernelSetGroupSize: hKernel = 0x2a54be8 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 + >>>> [672116802394124] zeKernelSetGroupSize: hKernel = 0x2689658 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 : - >>>> [806881439411094] zeCommandListAppendLaunchKernel: hCommandList = 0x2d3b218 hKernel = 0x2a54be8 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7ffcce631d78 {16385, 1, 1} hSignalEvent = 0x2cd8a88 numWaitEvents = 0x0 phWaitEvents = 0x0 + >>>> [672116802588167] zeCommandListAppendLaunchKernel: hCommandList = 0x3d56268 hKernel = 0x2689658 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7fffb216c840 {16385, 1, 1} hSignalEvent = 0x3b585a8 numWaitEvents = 0x0 phWaitEvents = 0x0 ``` - We used the form of `parallel_for` that takes the `nd_range`, which specifies the global iteration range (163850) and the local work-group size (10) like so: `nd_range<1>{{163850}, {10}}`. The first line above shows the workgroup size (`groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1`), and the second shows how many total workgroups will be needed to process the global iteration range (`{16385, 1, 1}`). + At like 106 we used the form of `parallel_for` that takes the `nd_range`, which specifies the global iteration range (163850) and the local work-group size (10) like so: `nd_range<1>{{163850}, {10}}`. The first line above shows the workgroup size (`groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1`), and the second shows how many total workgroups will be needed to process the global iteration range (`{16385, 1, 1}`). ### Determine Device Limits @@ -322,21 +322,21 @@ If the `parallel_for` were operating over a multi-dimensional range (for example You should know that different devices will have different amounts of memory set aside as SLM. In SYCL, you can query this number by passing `info::device::local_mem_size` to the `get_info` member of the `device` class. -Finally, running under `unitrace -c` you see: +Finally, running under `unitrace -c` with these debug variables set you will see: ``` : ->>>> [807184674868230] zeKernelSetGroupSize: hKernel = 0x257bbe8 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 -<<<< [807184674879934] zeKernelSetGroupSize [1560 ns] -> ZE_RESULT_SUCCESS(0x0) ->>>> [807184674884930] zeCommandListCreateImmediate: hContext = 0x24f4578 hDevice = 0x245e2f8 altdesc = 0x7ffeb187b880 {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC(0xe) 0 0 0 0 2 0} phCommandList = 0x7ffeb187b868 (hCommandList = 0x0) +>>>> [672514990780902] zeKernelSetGroupSize: hKernel = 0x3fa2658 groupSizeX = 0xa groupSizeY = 0x1 groupSizeZ = 0x1 +<<<< [672514990785597] zeKernelSetGroupSize [1532 ns] -> ZE_RESULT_SUCCESS(0x0) +>>>> [672514990790934] zeCommandListCreateImmediate: hContext = 0x3f24e58 hDevice = 0x3e8e6d8 altdesc = 0x7fff875129a0 {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC(0xe) 0 0 0 0 2 0} phCommandList = 0x7fff87512988 (hCommandList = 0x0) Flush Task for Immediate command list : Enabled -Using PCI barrier ptr: 0xbbbf8206000 -<<<< [807184675059549] zeCommandListCreateImmediate [169347 ns] hCommandList = 0x2862218 -> ZE_RESULT_SUCCESS(0x0) ->>>> [807184675067615] zeEventHostReset: hEvent = 0x27ffa88 -<<<< [807184675071317] zeEventHostReset [1401 ns] -> ZE_RESULT_SUCCESS(0x0) ->>>> [807184675077771] zeCommandListAppendLaunchKernel: hCommandList = 0x2862218 hKernel = 0x257bbe8 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7ffeb187bda8 {16385, 1, 1} hSignalEvent = 0x27ffa88 numWaitEvents = 0x0 phWaitEvents = 0x0 +Using PCI barrier ptr: 0xc615206b000 +<<<< [672514990965418] zeCommandListCreateImmediate [169631 ns] hCommandList = 0x566f268 -> ZE_RESULT_SUCCESS(0x0) +>>>> [672514990977437] zeEventHostReset: hEvent = 0x54715a8 +<<<< [672514990981535] zeEventHostReset [1392 ns] -> ZE_RESULT_SUCCESS(0x0) +>>>> [672514990988038] zeCommandListAppendLaunchKernel: hCommandList = 0x566f268 hKernel = 0x3fa2658 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7fff87512ee0 {16385, 1, 1} hSignalEvent = 0x54715a8 numWaitEvents = 0x0 phWaitEvents = 0x0 Size of SLM (656384) larger than available (131072) -<<<< [807184675135249] zeCommandListAppendLaunchKernel [48600 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195) +<<<< [672514991042897] zeCommandListAppendLaunchKernel [46307 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195) terminate called after throwing an instance of 'sycl::_V1::exception' what(): level_zero backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES) Aborted (core dumped) diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp index ae77a0ebfe..0d506b8e23 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/1_matrix_mul_SLM_size.cpp @@ -138,6 +138,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/2_matrix_mul.cpp b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/2_matrix_mul.cpp index 8f79f07e70..8138f2fd46 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/2_matrix_mul.cpp +++ b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/src/2_matrix_mul.cpp @@ -130,6 +130,10 @@ int main() { q.memcpy(&c_back[0], dev_c, M*P * sizeof(float)); q.wait(); + + sycl::free(dev_a, q); + sycl::free(dev_b, q); + sycl::free(dev_c, q); } int result;