-
Notifications
You must be signed in to change notification settings - Fork 103
/
Copy pathindex.html
279 lines (271 loc) · 10.6 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
<!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">
<script>
// This is needed when printing the slides to pdf
var link = document.createElement( 'link' );
link.rel = 'stylesheet';
link.type = 'text/css';
link.href = window.location.search.match( /print-pdf/gi ) ? '../common-revealjs/css/print/pdf.css' : '../common-revealjs/css/print/paper.css';
document.getElementsByTagName( 'head' )[0].appendChild( link );
</script>
<script>
// This is used to display the static images on each slide,
// See global-images in this html file and custom.css
(function() {
if(window.addEventListener) {
window.addEventListener('load', () => {
let slides = document.getElementsByClassName("slide-background");
if (slides.length === 0) {
slides = document.getElementsByClassName("pdf-page")
}
// Insert global images on each slide
for(let i = 0, max = slides.length; i < max; i++) {
let cln = document.getElementById("global-images").cloneNode(true);
cln.removeAttribute("id");
slides[i].appendChild(cln);
}
// Remove top level global images
let elem = document.getElementById("global-images");
elem.parentElement.removeChild(elem);
}, false);
}
})();
</script>
</head>
<body>
<div class="reveal">
<div class="slides">
<div id="global-images" class="global-images">
<img src="../common-revealjs/images/sycl_academy.png" />
<img src="../common-revealjs/images/sycl_logo.png" />
<img src="../common-revealjs/images/trademarks.png" />
<img src="../common-revealjs/images/codeplay.png" />
</div>
<!--Slide 1-->
<section class="hbox">
<div class="hbox" data-markdown>
## Advanced data flow
</div>
</section>
<!--Slide 2-->
<section class="hbox" data-markdown>
## Learning Objectives
* Learn about ways to use initial data and pinned memory
* Learn about uninitialized buffers and when to use them
</section>
<!--Slide 3-->
<section>
<div class="hbox" data-markdown>
#### Precursor
</div>
<div class="container" data-markdown>
Parts of this lecture will focus primarily on the buffer/accessor model.
</div>
</section>
<!--Slide 4-->
<section>
<div class="hbox" data-markdown>
#### Initial data
</div>
<div class="container" data-markdown>
* Often when writing SYCL kernel functions there is initial data which has been allocated somewhere else in the application.
</div>
</section>
<!--Slide 5-->
<section>
<div class="hbox" data-markdown>
#### Initial data (USM)
</div>
<div class="container">
<code class="code-100pc"><pre>
auto devicePtr = sycl::malloc_device(sizeInBytes, gpuQueue);
gpuQueue.memcpy(devicePtr, initialData, sizeInBytes).wait();
</code></pre>
</div>
<div class="container" data-markdown>
* When using the USM model (unless using system USM) pointers passed to kernel functions must be allocated by the SYCL runtime.
* This means you have to copy the initial data to the USM memory allocation using `memcpy`.
</div>
</section>
<!--Slide 6-->
<section>
<div class="hbox" data-markdown>
#### Initial data (USM)
</div>
<div class="container">
<code class="code-100pc"><pre>
auto sharedData = sycl::malloc_shared(sizeInBytes, gpuQueue);
</code></pre>
</div>
<div class="container" data-markdown>
* Alternatively if your device supports shared USM allocations you can allocate memory which is shared across the host and device.
* Then there is no need to copy the data to the device.
</div>
</section>
<!--Slide 7-->
<section>
<div class="hbox" data-markdown>
#### Initial data (buffer/accessor)
</div>
<div class="container">
<code class="code-100pc"><pre>
auto buf = sycl::buffer{initialData, sycl::range{size}};
</code></pre>
</div>
<div class="container" data-markdown>
* When using the buffer/accessor model a `buffer` can manage already allocated memory or have the SYCL runtime allocate it.
* To do this simply provide an initial pointer when constructing a `buffer`.
* Note that the SYCL runtime is free to allocate memory and copy this into it, which can introduce an overhead.
</div>
</section>
<!--Slide 8-->
<section>
<div class="hbox" data-markdown>
#### Use_host_pointer property
</div>
<div class="container">
<code class="code-100pc"><pre>
auto buf = sycl::buffer{initialData, sycl::range{size},
{sycl::property::buffer::use_host_ptr{}}};
</code></pre>
</div>
<div class="container" data-markdown>
* To prevent the runtime allocation memory you can provide the `property::buffer::use_host_ptr` property when constructing the `buffer`.
* This instructs the SYCL runtime that it may not allocate any additional memory.
* Though note that the backend (such as OpenCL) may still allocate memory.
</div>
</section>
<!--Slide 9-->
<section>
<div class="hbox" data-markdown>
#### Copy back
</div>
<div class="container" data-markdown>
* A `buffer` will synchronize the latest modified copy of the data it manages back to the initial pointer on destruction.
</div>
</section>
<!--Slide 10-->
<section>
<div class="hbox" data-markdown>
#### Set_final_data
</div>
<div class="container">
<code class="code-100pc"><pre>
auto buf = sycl::buffer{initialData, sycl::range{size}};
buf.set_final_data(finalData);
</code></pre>
</div>
<div class="container" data-markdown>
* To change the destination that a `buffer` will synchronize to on destruction you can call `set_final_data` with another.
* The address provided must be capable of holding the size of the data the `buffer` manages.
</div>
</section>
<!--Slide 11-->
<section>
<div class="hbox" data-markdown>
#### Set_final_data
</div>
<div class="container">
<code class="code-100pc"><pre>
auto buf = sycl::buffer{initialData, sycl::range{size}};
buf.set_final_data(nullptr);
</code></pre>
</div>
<div class="container" data-markdown>
* Alternatively to prevent the `buffer` from synchronizing back to the initial data entirely you can call `set_final_data` with `nullptr`.
* A `buffer` with no final data address is useful because the data can left on a device and not copied back to the host from the device.
</div>
</section>
<!--Slide 12-->
<section>
<div class="hbox" data-markdown>
#### Uninitialized buffers (buffer/accessor)
</div>
<div class="container">
<code class="code-100pc"><pre>
auto buf = sycl::buffer<T>{sycl::range{size}};
</code></pre>
</div>
<div class="container" data-markdown>
* As we've seen in the USM model all memory is allocated initialized, but `buffer`s can be constructed without initial data.
* A `buffer` like this is called uninitialized.
* To do this simply construct a `buffer` without initial data. Just remember to explicitly specify the buffer's data type as it can't be inferred from the initial data any more.
* Uninitialized `buffer`s are useful for a couple of reasons because they can be allocated directly on a device and don't require moving data from the host.
</div>
</section>
<!--Slide 13-->
<section>
<div class="hbox" data-markdown>
#### Using initial data and uninitialized buffers
</div>
<div class="container">
<div class="col" data-markdown>
</div>
<div class="col" data-markdown>
* Here we have an example of using these techniques:
* **Input data** is initialized with initial data but doesn't need to be copied back so it can use `set_final_data(nullptr)`.
* **Temporary** is only used on the device so can be an uninitialized `buffer`.
* **Output data** is initialized on the device and needs to be copied back so it can be an uninitialized `buffer` and use `set_final_data` to provide the final data address.
</div>
</div>
</section>
<!--Slide 14-->
<section>
<div class="hbox" data-markdown>
#### Pinned memory (buffer/accessor)
</div>
<div class="container" data-markdown>
* Pinned memory is a feature supported by most SYCL backends and devices.
* It allows you to allocate memory which can be mapped between the host and device more efficiently, providing similar benefits to USM.
* Though the requirements can vary from one device to another.
* It's always best to check the vendor's programming guide.
</div>
</section>
<!--Slide 15-->
<section>
<div class="hbox" data-markdown>
#### Pinned memory (buffer/accessor)
</div>
<div class="container" data-markdown>
* The SYCL runtime will always aim to manage the memory for you in the most efficient way for the target device.
* Generally there are two approaches to facilitate pinned memory:
* Allocate memory according to the vendor's programming guide, usually involved allocating a size of a particular multiple and aligned to a particular size, and then use the `property::buffer::use_host_ptr` property.
* Create an uninitialized `buffer` and allow the runtime to allocate the memory the appropriate way.
</div>
</section>
<!--Slide 16-->
<section>
<div class="hbox" data-markdown>
## Questions
</div>
</section>
<!--Slide 17-->
<section>
<div class="hbox" data-markdown>
#### Exercise
</div>
<div class="container" data-markdown>
Code_Exercises/Advanced_Data_Flow/source
</div>
<div class="container" data-markdown>
Write a SYCL application which uses uninitialized `buffer`s and disabling write back.
</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({mouseWheel: true, defaultNotes: true});
Reveal.configure({ slideNumber: true });
</script>
</body>
</html>