|
160 | 160 | </div>
|
161 | 161 | <div class="container">
|
162 | 162 | <div class="col" data-markdown>
|
163 |
| - * Work-items are collected together into **work-groups** |
164 |
| - * The size of work-groups is generally relative to what is optimal on the device being targeted |
165 |
| - * It can also be affected by the resources used by each work-item |
| 163 | + * Work-items are launched in parallel in a `sycl::range`. |
| 164 | + * In order to maximize parallelism, the range should correspond to the problem size. |
166 | 165 | </div>
|
167 | 166 | <div class="col" data-markdown>
|
168 |
| -  |
| 167 | +  |
169 | 168 | </div>
|
170 | 169 | </div>
|
171 | 170 | </section>
|
172 |
| - <!--Slide 8--> |
173 |
| - <section> |
174 |
| - <div class="hbox" data-markdown> |
175 |
| - #### SYCL execution model |
176 |
| - </div> |
177 |
| - <div class="container"> |
178 |
| - <div class="col" data-markdown> |
179 |
| - * SYCL kernel functions are invoked within an **nd-range** |
180 |
| - * An nd-range has a number of work-groups and subsequently a number of work-items |
181 |
| - * Work-groups always have the same number of work-items |
182 |
| - </div> |
183 |
| - <div class="col" data-markdown> |
184 |
| -  |
185 |
| - </div> |
186 |
| - </div> |
187 |
| - </section> |
188 |
| - <!--Slide 9--> |
189 |
| - <section> |
190 |
| - <div class="hbox" data-markdown> |
191 |
| - #### SYCL execution model |
192 |
| - </div> |
193 |
| - <div class="container"> |
194 |
| - <div class="col" data-markdown> |
195 |
| - * The nd-range describes an **iteration space**; how the work-items and work-groups are composed |
196 |
| - * An nd-range can be 1, 2 or 3 dimensions |
197 |
| - * An nd-range has two components |
198 |
| - * The **global-range** describes the total number of workitems in each dimension |
199 |
| - * The **local-range** describes the number of work-items in a work-group in each dimension |
200 |
| - </div> |
201 |
| - <div class="col" data-markdown> |
202 |
| -  |
203 |
| - </div> |
204 |
| - </div> |
205 |
| - </section> |
206 |
| - <!--Slide 10--> |
207 |
| - <section> |
208 |
| - <div class="hbox" data-markdown> |
209 |
| - #### SYCL execution model |
210 |
| - </div> |
211 |
| - <div class="container"> |
212 |
| - <div class="col" data-markdown> |
213 |
| - * Multiple work-items will generally execute concurrently |
214 |
| - * On vector hardware this is often done in lock-step, which means the same hardware instructions |
215 |
| - * The number of work-items that will execute concurrently can vary from one device to another |
216 |
| - * Work-items will be batched along with other work-items in the same work-group |
217 |
| - * The order work-items and workgroups are executed in is implementation defined |
218 |
| - </div> |
219 |
| - <div class="col" data-markdown> |
220 |
| -  |
221 |
| - </div> |
222 |
| - </div> |
223 |
| - </section> |
224 |
| - <!--Slide 11--> |
225 |
| - <section> |
226 |
| - <div class="hbox" data-markdown> |
227 |
| - #### SYCL execution model |
228 |
| - </div> |
229 |
| - <div class="container"> |
230 |
| - <div class="col" data-markdown> |
231 |
| - * Work-items in a work-group can be synchronized using a work-group barrier |
232 |
| - * All work-items within a work-group must reach the barrier before any can continue on |
233 |
| - </div> |
234 |
| - <div class="col" data-markdown> |
235 |
| -  |
236 |
| - </div> |
237 |
| - </div> |
238 |
| - </section> |
239 |
| - <!--Slide 12--> |
240 |
| - <section> |
241 |
| - <div class="hbox" data-markdown> |
242 |
| - #### SYCL execution model |
243 |
| - </div> |
244 |
| - <div class="container"> |
245 |
| - <div class="col" data-markdown> |
246 |
| - * SYCL does not support synchronizing across all work-items in the nd-range |
247 |
| - * The only way to do this is to split the computation into separate SYCL kernel functions |
248 |
| - </div> |
249 |
| - <div class="col" data-markdown> |
250 |
| -  |
251 |
| - </div> |
252 |
| - </div> |
253 |
| - </section> |
254 |
| - <!--Slide 13--> |
255 |
| - <section> |
256 |
| - <div class="hbox" data-markdown> |
257 |
| - #### SYCL memory model |
258 |
| - </div> |
259 |
| - <div class="container"> |
260 |
| - <div class="col-left-3" data-markdown> |
261 |
| -  |
262 |
| - </div> |
263 |
| - <div class="col-right-2" data-markdown> |
264 |
| - Each work-item can access |
265 |
| - * a dedicated region of **private memory** |
266 |
| - * a dedicated region of **local memory** accessible to all work-items in a work-group |
267 |
| - * a single region of **global memory** that's accessible to all work-items in a ND-range |
268 |
| - * a region of global memory reserved as **constant memory**, which is read-only |
269 |
| - </div> |
270 |
| - |
271 |
| - </div> |
272 |
| - </section> |
273 | 171 |
|
274 | 172 | <!--Slide 14-->
|
275 | 173 | <section>
|
|
370 | 268 | #### Exercise
|
371 | 269 | </div>
|
372 | 270 | <div class="container" data-markdown>
|
373 |
| - Code_Exercises/Exercise_6_Vector_Add/source |
| 271 | + Code_Exercises/Exercise_06_Vector_Add/source.cpp |
374 | 272 | </div>
|
375 | 273 | <div class="container" data-markdown>
|
376 | 274 | Implement a SYCL application that adds two arrays of
|
|
0 commit comments