[SYCL][Doc] Clarify InvokeSIMD + SYCL_EXTERNAL #5353
Conversation
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 <john.pennycook@intel.com>
Pennycook
added
the
spec extension
| 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. | ||
|
|
||
| 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. | ||
| This is a proposed extension specification, intended to gather community |
There was a problem hiding this comment.
This is just a formatting comment, but does asciidoc care about newlines in paragraphs? I ask because it seems like the previous version had long lines, while this change is re-flowed.
There was a problem hiding this comment.
asciidoc doesn't, but other reviewers have complained about long lines in extension documentation. Some of the boilerplate content from older specifications still have long lines; I saw this as an opportunity to update InvokeSIMD to use some of the newer boilerplate content (with short lines).
There was a problem hiding this comment.
I think this is a bit of a punt in the sense that we don't really solve the problem of wanting to call a properly compiled SYCL_EXTERNAL function, but I think it is fine for the time being. It would be good to think about the linkage issue and ways to control how external functions are compiled down the road.
Signed-off-by: John Pennycook <john.pennycook@intel.com>
You're right, but I think we need more implementation experience first. I suspect that the upcoming ESIMD-based implementation of InvokeSIMD is going to require an ESIMD-specific macro, and in the short-term we'll probably also need some CUDA-backend-specific control for overriding the default compiler behavior when targeting PTX. Long-term I hope we can adopt something portable and similar to the function properties design described in the latest revision of the KernelProperties extension (see here): SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::simd_property_name_tbd))
sycl::simd<float, 8> simd_function(sycl::simd<float, 8> x);...but we'll need KernelProperties to be implemented first, and to figure out exactly what the property describes. Some backends may be satisfied just by "this is for SIMD", but I'm worried that some will need more information (e.g. the sub-group size of the kernel expected to call the SIMD function, maybe?). |
|
@bader I think this one is fine to merge; @rolandschulz did review 7 days ago, and has marked the only issue raised as resolved. |
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 john.pennycook@intel.com