Skip to content

Commit 834e66e

Browse files
committed
[SYCL][Doc] Clarify InvokeSIMD + SYCL_EXTERNAL
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 5746906 commit 834e66e

File tree

1 file changed

+27
-18
lines changed

1 file changed

+27
-18
lines changed

sycl/doc/extensions/InvokeSIMD/InvokeSIMD.asciidoc

Lines changed: 27 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,24 @@ 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+
a `SYCL_EXTERNAL` function that was compiled for a SPMD context. Behavior in
389+
this case is undefined.
390+
391+
NOTE: Implementations are encouraged to throw an error in cases where this
392+
behavior can be detected, and to do so as early as possible.
393+
394+
This extension does not define a mechanism for requesting a `SYCL_EXTERNAL`
395+
function to be compiled for a SIMD context. Any such mechanism and its
396+
interaction with `invoke_simd` is implementation-defined.
397+
398+
NOTE: An explicit mechanism is required because the presence or absence of a
399+
`sycl::ext::oneapi::experimental::simd` class is not sufficient to determine
400+
the execution model associated with a function. For example, the execution
401+
model for a function accepting only scalar parameters and performing only
402+
scalar arithmetic is ambiguous: it could be called by each work-item in SPMD
403+
code, or called as-if by a single work-item in a SIMD function.
404+
388405
== Issues
389406

390407
. Should we allow reference arguments?
@@ -428,13 +445,5 @@ SYCL 2020 specification states that this only applies to inter-device transfers.
428445
|1|2021-03-30|John Pennycook|*Initial public working draft*
429446
|2|2021-03-31|John Pennycook|*Rename extension and add feature test macro*
430447
|3|2021-04-23|John Pennycook|*Split uniform wrapper into separate extension*
448+
|4|2022-01-20|John Pennycook|*Clarify interaction with SYCL_EXTERNAL*
431449
|========================================
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)