@@ -37,7 +37,7 @@ https://github.com/intel/llvm/issues
37
37
38
38
== Dependencies
39
39
40
- This extension is written against the SYCL 2020 specification, Revision 4 and
40
+ This extension is written against the SYCL 2020 specification, Revision 9 and
41
41
the following extensions:
42
42
43
43
- link:sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties]
@@ -287,136 +287,6 @@ by the property, the implementation must throw a synchronous exception with the
287
287
288
288
|===
289
289
290
- === Adding a Property List to a Kernel Launch
291
-
292
- To enable properties to be associated with kernels, this extension adds
293
- new overloads to each of the variants of `single_task`, `parallel_for` and
294
- `parallel_for_work_group` defined in the `queue` and `handler` classes. These
295
- new overloads accept a `sycl::ext::oneapi::experimental::properties` argument. For
296
- variants accepting a parameter pack, the `sycl::ext::oneapi::experimental::properties`
297
- argument is inserted immediately prior to the parameter pack; for variants not
298
- accepting a parameter pack, the `sycl::ext::oneapi::experimental::properties` argument is
299
- inserted immediately prior to the kernel function.
300
-
301
- The overloads introduced by this extension are listed below:
302
- ```c++
303
- namespace sycl {
304
- class queue {
305
- public:
306
- template <typename KernelName, typename KernelType, typename PropertyList>
307
- event single_task(PropertyList properties, const KernelType &kernelFunc);
308
-
309
- template <typename KernelName, typename KernelType, typename PropertyList>
310
- event single_task(event depEvent, PropertyList properties,
311
- const KernelType &kernelFunc);
312
-
313
- template <typename KernelName, typename KernelType, typename PropertyList>
314
- event single_task(const std::vector<event> &depEvents,
315
- PropertyList properties,
316
- const KernelType &kernelFunc);
317
-
318
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
319
- event parallel_for(range<Dims> numWorkItems,
320
- Rest&&... rest);
321
-
322
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
323
- event parallel_for(range<Dims> numWorkItems, event depEvent,
324
- PropertyList properties,
325
- Rest&&... rest);
326
-
327
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
328
- event parallel_for(range<Dims> numWorkItems,
329
- const std::vector<event> &depEvents,
330
- PropertyList properties,
331
- Rest&&... rest);
332
-
333
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
334
- event parallel_for(nd_range<Dims> executionRange,
335
- PropertyList properties,
336
- Rest&&... rest);
337
-
338
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
339
- event parallel_for(nd_range<Dims> executionRange,
340
- event depEvent,
341
- PropertyList properties,
342
- Rest&&... rest);
343
-
344
- template <typename KernelName, int Dims, typename PropertyList, typename... Rest>
345
- event parallel_for(nd_range<Dims> executionRange,
346
- const std::vector<event> &depEvents,
347
- PropertyList properties,
348
- Rest&&... rest);
349
- }
350
- }
351
-
352
- namespace sycl {
353
- class handler {
354
- public:
355
- template <typename KernelName, typename KernelType, typename PropertyList>
356
- void single_task(PropertyList properties, const KernelType &kernelFunc);
357
-
358
- template <typename KernelName, int dimensions, typename PropertyList, typename... Rest>
359
- void parallel_for(range<dimensions> numWorkItems,
360
- PropertyList properties,
361
- Rest&&... rest);
362
-
363
- template <typename KernelName, int dimensions, typename PropertyList, typename... Rest>
364
- void parallel_for(nd_range<dimensions> executionRange,
365
- PropertyList properties,
366
- Rest&&... rest);
367
-
368
- template <typename KernelName, typename WorkgroupFunctionType, int dimensions, typename PropertyList>
369
- void parallel_for_work_group(range<dimensions> numWorkGroups,
370
- PropertyList properties,
371
- const WorkgroupFunctionType &kernelFunc);
372
-
373
- template <typename KernelName, typename WorkgroupFunctionType, int dimensions, typename PropertyList>
374
- void parallel_for_work_group(range<dimensions> numWorkGroups,
375
- range<dimensions> workGroupSize,
376
- PropertyList properties,
377
- const WorkgroupFunctionType &kernelFunc);
378
- }
379
- }
380
- ```
381
-
382
- Passing a property list as an argument in this way allows properties to be
383
- associated with a kernel function without modifying its type. This enables
384
- the same kernel function (e.g. a lambda) to be submitted multiple times with
385
- different properties, or for libraries building on SYCL to add properties
386
- (e.g. for performance reasons) to user-provided kernel functions.
387
-
388
- All the properties defined in this extension have compile-time values. However,
389
- an implementation may support additional properties which could have run-time
390
- values. When this occurs, the `properties` parameter may be a property list
391
- containing a mix of both run-time and compile-time values, and a SYCL
392
- implementation should respect both run-time and compile-time information when
393
- determining the correct way to launch a kernel. However, only compile-time
394
- information can modify the compilation of the kernel function itself.
395
-
396
- A simple example of using this extension to set a required work-group size
397
- and required sub-group size is given below:
398
-
399
- ```c++
400
- sycl::ext::oneapi::experimental::properties properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
401
- sycl::ext::oneapi::experimental::sub_group_size<8>};
402
- q.parallel_for(range<2>{16, 16}, properties, [=](id<2> i) {
403
- a[i] = b[i] + c[i];
404
- }).wait();
405
- ```
406
-
407
- NOTE: It is currently not possible to use the same kernel function in two
408
- commands with different properties. For example, the following will result in an
409
- error at compile-time:
410
-
411
- ```c++
412
- auto kernelFunc = [=](){};
413
- q.single_task(kernelFunc);
414
- q.single_task(
415
- sycl::ext::oneapi::experimental::properties{
416
- sycl::ext::oneapi::experimental::sub_group_size<8>},
417
- kernelFunc);
418
- ```
419
-
420
290
== Embedding Properties into a Kernel
421
291
422
292
In other situations it may be useful to embed a kernel's properties directly
0 commit comments