Skip to content

Commit 6320c46

Browse files
Address concerns raised with previous pull request
Signed-off-by: Charles Congdon <charles.w.congdon@intel.com>
1 parent e81aa30 commit 6320c46

7 files changed

Lines changed: 14 additions & 10 deletions

File tree

Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/README.md

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -543,7 +543,7 @@ In `b1_matrix_mul_null_usm.cpp` a bad (in this case, null) pointer that is suppo
543543
544544
#### Debugging the Problem
545545
546-
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.
546+
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.
547547
548548
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.
549549

Tools/ApplicationDebugger/guided_matrix_mult_BadBuffers/src/b1_matrix_mul_null_usm.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -133,7 +133,11 @@ int main() {
133133

134134
q.wait();
135135

136-
sycl::free(dev_a, q);
136+
sycl::free(dev_a, q); // For the purposes of this demo, since we lost the original pointer value
137+
// and will crash before this point, this free is effectively a no-op and a
138+
// bug in its own right (we should be checking pointers best we can before
139+
// trying to free them).
140+
// All USM allocations should be released after use.
137141
sycl::free(dev_b, q);
138142
sycl::free(dev_c, q);
139143
}

Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/README.md

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -174,7 +174,7 @@ To complete the steps in the following section, you must download the [Unified T
174174
[opencl:gpu][opencl:2] Intel(R) OpenCL Graphics, Intel(R) Data Center GPU Max 1550 OpenCL 3.0 NEO [25.18.33578]
175175
[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]
176176
```
177-
> **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.
177+
> **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.
178178
179179
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?
180180

@@ -461,7 +461,7 @@ In case we need view code running on the GPU, we need to enable GPU debugging.
461461

462462
### Identify the Problem without Code Inspection
463463

464-
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).
464+
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).
465465

466466
One of the things that the Unified Tracing and Profiling utility can help us see
467467
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.
633633
101 queue q2(devicecontext, selected_device);
634634
102 float * dev_c = sycl::malloc_device<float>(M*P, q2);
635635
```
636-
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.
636+
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.
637637
638638
6. While we are at it, what does the debugger show us?
639639
```

Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/1_matrix_mul_invalid_contexts.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -79,7 +79,7 @@ int main() {
7979
}
8080

8181
// Be very specific about the device to use.
82-
queue q(devices[0]);
82+
queue q(devices[0], propList);
8383

8484
cout << "Computing" << "\n";
8585
cout << "Device: " << q.get_device().get_info<info::device::name>() << "\n";
@@ -95,7 +95,7 @@ int main() {
9595
#ifdef BAD_FREE
9696
device selected_device = devices[0];
9797
#else
98-
device selected_device = devices[1];
98+
device selected_device = devices.size() > 1 ? devices[1] : devices[0];
9999
#endif
100100
context devicecontext(selected_device);
101101
queue q2(devicecontext, selected_device);

Tools/ApplicationDebugger/guided_matrix_mult_InvalidContexts/src/2_matrix_mul.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -78,7 +78,7 @@ int main() {
7878
<< "] " << backend << ", " << vendor << std::endl;
7979
}
8080

81-
queue q(default_selector_v);
81+
queue q(default_selector_v, propList);
8282

8383
cout << "Computing" << "\n";
8484
cout << "Device: " << q.get_device().get_info<info::device::name>() << "\n";

Tools/ApplicationDebugger/guided_matrix_mult_RaceCondition/README.md

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -336,7 +336,7 @@ Unfortunately not; pretty much the same thing happens - they both produce incorr
336336
export PATH=$PATH:/opt/intel/oneapi:$DRIVERLOC
337337
```
338338
339-
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.
339+
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.
340340
```
341341
ONEAPI_DEVICE_SELECTOR=*:cpu ./1_matrix_mul_race_condition
342342
```

Tools/ApplicationDebugger/guided_matrix_mult_SLMSize/README.md

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -287,7 +287,7 @@ Among other things, the Tracing and Profiling utility can print every low-level
287287
>>>> [672116802588167] zeCommandListAppendLaunchKernel: hCommandList = 0x3d56268 hKernel = 0x2689658 (_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi1EEEE_) pLaunchFuncArgs = 0x7fffb216c840 {16385, 1, 1} hSignalEvent = 0x3b585a8 numWaitEvents = 0x0 phWaitEvents = 0x0
288288
```
289289

290-
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}`).
290+
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}`).
291291

292292
### Determine Device Limits
293293

0 commit comments

Comments
 (0)