Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -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.

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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?

Expand Down Expand Up @@ -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.
Expand Down Expand Up @@ -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<float>(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?
```
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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<info::device::name>() << "\n";
Expand All @@ -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);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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<info::device::name>() << "\n";
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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
```
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down