diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md index fe92dbbd9e..9c1ff0b6b0 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md @@ -543,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 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. +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 backends 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 546e37050e..dac2d17c99 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 @@ -133,7 +133,11 @@ int main() { q.wait(); - sycl::free(dev_a, q); + sycl::free(dev_a, q); // For the purposes of this demo, since we lost the original pointer value + // and will crash before this point, this free is effectively a no-op and a + // bug in its own right (we should be checking pointers best we can before + // trying to free them). + // All USM allocations should be released after use. sycl::free(dev_b, q); sycl::free(dev_c, q); } diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md index d2621a7a34..3e87a0e43e 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md @@ -174,7 +174,7 @@ To complete the steps in the following section, you must download the [Unified T [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. + > **Note:** If you have only one `[level_zero:gpu]` device listed, or the order is different from the above, 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? @@ -461,7 +461,7 @@ In case we need view code running on the GPU, we need to enable GPU debugging. ### Identify the Problem without Code Inspection -You need to build the [Unified Tracing and Profiling Tool](#getting-the-tracing-and-profiling-tool) before completing this section. Once you have built the utility, you can start it before your program (similar to using GBD). +You need to build the [Unified Tracing and Profiling Tool](#getting-the-tracing-and-profiling-tool) before completing this section. Once you have built the utility, you can start it before your program (similar to using GDB). One of the things that the Unified Tracing and Profiling utility can help us see is every low-level API call made to OpenCL™ or Level Zero. We will use it to attempt to match the source to the events. @@ -633,7 +633,7 @@ be entirely re-written, which is beyond the scope of this tutorial. 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. + 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 from 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? ``` 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 9810af1bbd..3c97ca7510 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 @@ -79,7 +79,7 @@ int main() { } // Be very specific about the device to use. - queue q(devices[0]); + queue q(devices[0], propList); cout << "Computing" << "\n"; cout << "Device: " << q.get_device().get_info() << "\n"; @@ -95,7 +95,7 @@ int main() { #ifdef BAD_FREE device selected_device = devices[0]; #else - device selected_device = devices[1]; + device selected_device = devices.size() > 1 ? devices[1] : devices[0]; #endif context devicecontext(selected_device); queue q2(devicecontext, selected_device); 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 0883604faa..c9e6ea99f0 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 @@ -78,7 +78,7 @@ int main() { << "] " << backend << ", " << vendor << std::endl; } - queue q(default_selector_v); + queue q(default_selector_v, propList); cout << "Computing" << "\n"; cout << "Device: " << q.get_device().get_info() << "\n"; diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md index 5583300434..6f5ca12fe7 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md @@ -336,7 +336,7 @@ Unfortunately not; pretty much the same thing happens - they both produce incorr export PATH=$PATH:/opt/intel/oneapi:$DRIVERLOC ``` -Similarly, we an force the program to run on the CPU, which sometimes can avoid problems in your code that are specific to offloading to the GPU. +Similarly, we can force the program to run on the CPU, which sometimes can avoid problems in your code that are specific to offloading to the GPU. ``` ONEAPI_DEVICE_SELECTOR=*:cpu ./1_matrix_mul_race_condition ``` diff --git a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md index c86dbdbbf0..ff97c337fc 100644 --- a/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md +++ b/Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md @@ -287,7 +287,7 @@ Among other things, the Tracing and Profiling utility can print every low-level >>>> [672116802588167] zeCommandListAppendLaunchKernel: hCommandList = 0x3d56268 hKernel = 0x2689658 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7fffb216c840 {16385, 1, 1} hSignalEvent = 0x3b585a8 numWaitEvents = 0x0 phWaitEvents = 0x0 ``` - 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}`). + At line 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