Skip to content

Commit ac3e816

Browse files
authored
[SYCL][Doc] Clarify InvokeSIMD + SYCL_EXTERNAL (#5353)
The interaction between the InvokeSIMD extension and SYCL_EXTERNAL functions was previously undefined. This commit clarifies when such interactions are expected to fail and why. Signed-off-by: John Pennycook <[email protected]>
1 parent 12d7c1f commit ac3e816

File tree

1 file changed

+29
-18
lines changed

1 file changed

+29
-18
lines changed

sycl/doc/extensions/InvokeSIMD/InvokeSIMD.asciidoc

Lines changed: 29 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -21,27 +21,26 @@ IMPORTANT: This specification is a draft.
2121

2222
NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos.
2323

24-
NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons.
25-
2624
This extension introduces a mechanism to invoke explicitly vectorized functions
2725
written using SIMD types from a SYCL kernel.
2826

2927
== Notice
3028

31-
Copyright (c) 2021 Intel Corporation. All rights reserved.
29+
Copyright (c) 2021-2022 Intel Corporation. All rights reserved.
3230

3331
== Status
3432

3533
Working Draft
3634

37-
This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension.
38-
39-
Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products.
35+
This is a proposed extension specification, intended to gather community
36+
feedback. Interfaces defined in this specification may not be implemented yet
37+
or may be in a preliminary state. The specification itself may also change in
38+
incompatible ways before it is finalized. Shipping software products should not
39+
rely on APIs defined in this specification.
4040

4141
== Version
4242

43-
Built On: {docdate} +
44-
Revision: 3
43+
Revision: 4
4544

4645
== Contacts
4746

@@ -51,7 +50,7 @@ Jason Sewall, Intel (jason 'dot' sewall 'at' intel 'dot' com) +
5150

5251
== Dependencies
5352

54-
This extension is written against the SYCL 2020 specification, Revision 2 and the following extensions:
53+
This extension is written against the SYCL 2020 specification, Revision 4 and the following extensions:
5554

5655
- SPV_INTEL_function_pointers
5756
- SYCL_EXT_ONEAPI_UNIFORM
@@ -385,6 +384,26 @@ Since this execution model guarantees SIMD-like behavior, there is no need for
385384
the user to insert any form of explicit synchronization functions to ensure
386385
memory consistency across SIMD lanes.
387386

387+
Due to potential differences in execution model, a SIMD function cannot call
388+
any function that was compiled for a SPMD context. `SYCL_EXTERNAL` functions
389+
defined in a different translation unit to the SIMD function are not required
390+
to be compiled for SIMD contexts by default, and as a result such functions
391+
are incompatible with SIMD functions by default.
392+
393+
NOTE: Implementations are encouraged to throw an error in cases where this
394+
behavior can be detected, and to do so as early as possible.
395+
396+
This extension does not define a mechanism for requesting a `SYCL_EXTERNAL`
397+
function to be compiled for a SIMD context. Any such mechanism and its
398+
interaction with `invoke_simd` is implementation-defined.
399+
400+
NOTE: An explicit mechanism is required because the presence or absence of a
401+
`sycl::ext::oneapi::experimental::simd` class is not sufficient to determine
402+
the execution model associated with a function. For example, the execution
403+
model for a function accepting only scalar parameters and performing only
404+
scalar arithmetic is ambiguous: it could be called by each work-item in SPMD
405+
code, or called as-if by a single work-item in a SIMD function.
406+
388407
== Issues
389408

390409
. Should we allow reference arguments?
@@ -428,13 +447,5 @@ SYCL 2020 specification states that this only applies to inter-device transfers.
428447
|1|2021-03-30|John Pennycook|*Initial public working draft*
429448
|2|2021-03-31|John Pennycook|*Rename extension and add feature test macro*
430449
|3|2021-04-23|John Pennycook|*Split uniform wrapper into separate extension*
450+
|4|2022-01-20|John Pennycook|*Clarify interaction with SYCL_EXTERNAL*
431451
|========================================
432-
433-
//************************************************************************
434-
//Other formatting suggestions:
435-
//
436-
//* Use *bold* text for host APIs, or [source] syntax highlighting.
437-
//* Use +mono+ text for device APIs, or [source] syntax highlighting.
438-
//* Use +mono+ text for extension names, types, or enum values.
439-
//* Use _italics_ for parameters.
440-
//************************************************************************

0 commit comments

Comments
 (0)