forked from codeplaysoftware/syclacademy
-
Notifications
You must be signed in to change notification settings - Fork 0
/
index.html
328 lines (311 loc) · 12.3 KB
/
index.html
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
<!DOCTYPE html>
<html>
<head>
<meta charset="utf-8">
<link rel="stylesheet" href="../common-revealjs/css/reveal.css">
<link rel="stylesheet" href="../common-revealjs/css/theme/white.css">
<link rel="stylesheet" href="../common-revealjs/css/custom.css">
</head>
<body>
<span class='overlay-image bottom-left'>SYCL and the SYCL logo are trademarks of the Khronos Group Inc.</span>
<span class='overlay-image bottom-right'><img src="../common-revealjs/images/codeplay.png" alt="Codeplay Software"></span>
<span class='overlay-image top-left'><img src="../common-revealjs/images/sycl_academy.png" alt="SYCL Academy"></span>
<span class='overlay-image top-right'><img src="../common-revealjs/images/sycl_logo.png" alt="SYCL"></span>
<div class="reveal">
<div class="slides">
<!--Slide 1-->
<section class="hbox" data-markdown>
## Handling SYCL Errors
</section>
<!--Slide 2-->
<section class="hbox" data-markdown>
## Learning Objectives
* Learn about how SYCL handles errors
* Learn about the difference between synchronous and asynchronous exceptions
* Learn how to handle exceptions and retrieve further information
* Learn about the different exception types
* Learn about the host device and how to use it
</section>
<!--Slide 3-->
<section class="hbox" data-markdown>
* In SYCL errors are handled by throwing exceptions
* It is crucial that these errors are handled otherwise your application may silently fail
* In SYCL there are two kinds of error
* Synchronous errors (thrown in user thread)
* Asynchronous errors (thrown by the SYCL scheduler)
</section>
<!--Slide 4-->
<section class="hbox">
<div class="hbox" data-markdown>
</div>
</section>
<!--Slide 5-->
<section class="hbox" data-markdown>
## Handling Errors
</section>
<!--Slide 6-->
<section>
<div class="hbox" >
<code class="code-60pc"><pre>
#include <CL/sycl.hpp>
using namespace cl::sycl;
class add;
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
queue gpuQueue(gpu_selector{});
buffer<float, 1> bufA(dA.data(), range<1>(dA.size()));
buffer<float, 1> bufB(dB.data(), range<1>(dB.size()));
buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
gpuQueue.submit([&](handler &cgh){
auto inA = bufA.get_access<access::mode::read>(cgh);
auto inB = bufB.get_access<access::mode::read>(cgh);
auto out = bufO.get_access<access::mode::write>(cgh);
cgh.parallel_for<add>(range<1>(dA.size()), [=](id<1> i){
out[i] = inA[i] + inB[i];
});
});
gpuQueue.wait();
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* If errors are not handled, the application can fail silently
</div>
</section>
<!--Slide 7-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
<mark>try{</mark>
queue gpuQueue(gpu_selector{});
buffer<float, 1> bufA(dA.data(), range<1>(dA.size()));
buffer<float, 1> bufB(dB.data(), range<1>(dB.size()));
buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
gpuQueue.submit([&](handler &cgh){
auto inA = bufA.get_access<access::mode::read>(cgh);
auto inB = bufB.get_access<access::mode::read>(cgh);
auto out = bufO.get_access<access::mode::write>(cgh);
cgh.parallel_for<add>(range<1>(dA.size()), [=](id<1> i){
out[i] = inA[i] + inB[i];
});
});
gpuQueue.wait();
<mark>} catch (...) { /* handle errors */ }</mark>
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* Synchronous errors are typically thrown by SYCL API functions
* In order to handle all SYCL errors you must wrap everything in a try-catch block
</div>
</section>
<!--Slide 8-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try{
queue gpuQueue(gpu_selector{}, <mark>async_handler{}</mark>);
buffer<float, 1> bufA(dA.data(), range<1>(dA.size()));
buffer<float, 1> bufB(dB.data(), range<1>(dB.size()));
buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
gpuQueue.submit([&](handler &cgh){
auto inA = bufA.get_access<access::mode::read>(cgh);
auto inB = bufB.get_access<access::mode::read>(cgh);
auto out = bufO.get_access<access::mode::write>(cgh);
cgh.parallel_for<add>(range<1>(dA.size()), [=](id<1> i){
out[i] = inA[i] + inB[i];
});
});
<mark>gpuQueue.wait_and_throw();</mark>
} catch (...) { /* handle errors */
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* Asynchronous errors errors that may have occurred will be thrown after a command group has been submitted to a queue
* To handle these errors you must provide an async handler when constructing the queue object
* Then you must also call the **throw_asynchronous or wait_and_throw** member functions of the queue class
* This will pass the exceptions to the async handler in the user thread so they can be thrown
</div>
</section>
<!--Slide 9-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try{
queue gpuQueue(gpu_selector{}, <mark>[=](sycl::exception_list eL) {
for (auto e : eL) { std::rethrow_exception(e); }
}</mark>);
buffer<float, 1> bufA(dA.data(), range<1>(dA.size()));
buffer<float, 1> bufB(dB.data(), range<1>(dB.size()));
buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
gpuQueue.submit([&](handler &cgh){ // Command group submitted to queue
auto inA = bufA.get_access<access::mode::read>(cgh);
auto inB = bufB.get_access<access::mode::read>(cgh);
auto out = bufO.get_access<access::mode::write>(cgh);
cgh.parallel_for<add>(range<1>(dA.size()), [=](id<1> i){
out[i] = inA[i] + inB[i];
});
});
gpuQueue.wait_and_throw(); } catch (...) { /* handle errors */ }
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* The async handler is a C++ lambda or function object that takes as a parameter an **exception_list**
* The exception_list class is a wrapper around a list of **exception_ptrs** which can be iterated over
* The exception_ptrs can be rethrown by passing them to **std::rethrow_exception**
</div>
</section>
<!--Slide 10-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try {
queue gpuQueue(gpu_selector{}, [=](sycl::exception_list eL) {
for (auto e : eL) { std::rethrow_exception(e); }
});
...
gpuQueue.wait_and_throw();
} catch (<mark>std::exception e</mark>) {
<mark>std::cout << “Exception caught: ” << e.what()
<< std::endl;</mark>
}
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* Once rethrown and caught, a SYCL exception can provide information about the error
* The **what** member function will return a string with more details
</div>
</section>
<!--Slide 11-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try {
queue gpuQueue(gpu_selector{}, [=](sycl::exception_list eL) {
for (auto e : eL) { std::rethrow_exception(e); }
});
...
gpuQueue.wait_and_throw();
} catch (std::exception e) {
std::cout << “Exception caught: ” << e.what();
<mark>std:: cout << “ With OpenCL error code: ”</mark>
<mark><< e.get_cl_code() << std::endl;</mark>
}
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* If the exception has an OpenCL error code associated with it this can be retrieved by calling the <mark>get_cl_code</mark> member function
* If there is no OpenCL error code this will return <mark>CL_SUCCESS</mark>
</div>
</section>
<!--Slide 12-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try {
queue gpuQueue(gpu_selector{}, [=](sycl::exception_list eL) {
for (auto e : eL) { std::rethrow_exception(e); }
});
...
gpuQueue.wait_and_throw();
} catch (std::exception e) {
<mark>if (e.has_context()) {</mark>
<mark>if (e.get_context() == gpuContext) {</mark>
<mark>/* handle error */</mark>
<mark>}</mark>
<mark>}</mark>
}
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* The **has_context** member function will tell you if there is a SYCL context associated with the error
* If that returns true then the **get_context** member function will return the associated SYCL context object
</div>
</section>
<!--Slide 13-->
<section>
<div class="hbox" data-markdown>
## Exception Types
</div>
</section>
<!--Slide 14-->
<section>
<div class="hbox" data-markdown>
* In SYCL there are a number of different exception types that inherit from **std::exception**
* E.g. runtime_error, kernel_error
* The [SYCL 1.2.1 specification](https://sycl.tech/) will detail cases where a specific error can be expected
</div>
</section>
<!--Slide 15-->
<section>
<div class="hbox" data-markdown>
## Debugging SYCL Kernel Functions
</div>
</section>
<!--Slide 16-->
<section>
<div class="hbox" data-markdown>
* Every SYCL implementation is required to provide a host device
* This device executes native C++ code but is guaranteed to emulate the SYCL execution and memory model
* This means you can debug a SYCL kernel function by switching to the host device and using a standard C++ debugger
* For example gdb
</div>
</section>
<!--Slide 17-->
<section>
<div class="hbox">
<code class="code-60pc"><pre>
int main(int argc, char *argv[]) {
std::vector<float> dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try{
queue <mark>hostQueue(host_selector{}</mark>, async_handler{});
buffer<float, 1> bufA(dA.data(), range<1>(dA.size()));
buffer<float, 1> bufB(dB.data(), range<1>(dB.size()));
buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
<mark>hostQueue</mark>.submit([&](handler &cgh){
auto inA = bufA.get_access<access::mode::read>(cgh);
auto inB = bufB.get_access<access::mode::read>(cgh);
auto out = bufO.get_access<access::mode::write>(cgh);
cgh.parallel_for<add>(range<1>(dA.size()),
[=](id<1> i){out[i] = inA[i] + inB[i];});
});
gpuQueue.wait_and_throw();
} catch (...) { /* handle errors */ }
}
</code></pre>
</div>
<div class="bottom-bullets" data-markdown>
* Any SYCL application can be debugged on the host device by switching the queue for a host queue
* By replacing the device selector for the host_selector will ensure that the queue submits all work to the host device
</div>
</section>
</div>
</div>
<script src="../common-revealjs/js/reveal.js"></script>
<script src="../common-revealjs/plugin/markdown/marked.js"></script>
<script src="../common-revealjs/plugin/markdown/markdown.js"></script>
<script src="../common-revealjs/plugin/notes/notes.js"></script>
<script>
Reveal.initialize();
Reveal.configure({ slideNumber: true });
</script>
</body>
</html>