[SYCL][Doc] Add proposed sycl_ext_oneapi_spirv_queries spec

[bashbaug]

This is a proposed SYCL extension to query the SPIR-V extended instruction sets, SPIR-V extensions, and SPIR-V capabilities supported by a SYCL device. It essentially provides the same functionality as the OpenCL cl_khr_spirv_queries extension via SYCL.

[gmlueck]

I think we should not have this query. Instead, I think any implementation that supports this extension must support this query on all of its devices. I understand that this puts more of a burden on implementors, but I think it makes life easier for application developers.

Note that we do not have any similar query like this for OpenCL in sycl_ext_oneapi_kernel_compiler_opencl. Instead, we assume that any implementation that supports that extension also supports the queries.

[bashbaug]

Yeah, I agree the "do you support the queries query" is a little weird but I still think it's the best tradeoff. With the proposal above, if one device does not support this SPIR-V query functionality (say because it's using an older driver), it means that all of the devices do not support the SPIR-V query functionality, which seems unfortunate. Note also that these are runtime queries, not compile-time queries.

Other options we could consider are:

1. Return "false" if the device does not support SPIR-V queries or if the device does not support the particular SPIR-V feature being queried. Easy to implement, but also misleading, and I can already imagine applications querying for SPIR-V features that are known to be present for any SPIR-V implementation (e.g. the "Addresses" capability) to determine if the SPIR-V queries are supported.
2. Have a fallback implementation in the SYCL runtime that derives supported SPIR-V features based on other queries. I think this is doable, but it'd be error-prone, and it'd also need to be conservative so I'm not sure it's worth the effort.

[Pennycook]

I agree that there being two ways to get to "false" is bad design and potentially misleading.

Is there a reason we couldn't do this via an aspect, and turn these queries into optional features? Specifically, I am suggesting that:

• We add aspect::spirv to identify devices that support launching kernels written in SPIR-V.
• Remove ext_oneapi_supports_spirv_queries(), in favor of device.has(aspect::spirv).
• Say that the other queries throw an exception if the device doesn't have aspect::spirv.

[bashbaug]

I like where this is going, but I think we need to differentiate between devices that support launching kernels written in SPIR-V and devices that support the SPIR-V queries. Rationale: The SPIR-V queries extension is new and it will take time for drivers supporting it to become widespread, whereas devices and drivers that support launching kernels written in SPIR-V have been shipping for years.

So, what if:

1. We add aspect::spirv_queries to identify devices that support the SPIR-V queries. We could also add aspect::spirv, if desired, though that is outside of the scope of this extension.
2. We remove ext_oneapi_supports_spirv_queries() in favor of device.has(aspect::spirv_queries).
3. We say that the other queries thrown an exception if the device does not have aspect::spirv_queries.

Note that item (3) works regardless whether we keep the current functions or switch to the device info mechanism.

[gmlueck]

I have a slight preference for @Pennycook's proposal, but I'd name the aspect aspect::spirv_queries (well, really aspect::ext_oneapi_spirv_queries). This is because a backend / device might support SPIR-V kernels, but still not support the queries. I believe this is the case that @bashbaug is concerned about.

This design would allow the application to decide whether it wanted to test the aspect, or whether it wanted to assume that all devices it cares about will always support the queries. If the application's assumption is wrong, they will get an obvious feature_not_supported exception, and then they will know they need to test the aspect.

I would also define the queries to return false if the device does not support SPIR-V at all (e.g. a CUDA device). If a device does not support SPIR-V, then it clearly does not support any SPIR-V extensions or capabilities. We don't need any special backend support to know that.

[gmlueck]

I didn't see your comment @bashbaug until after I wrote mine. I think we are on the same page.

What do you think of my last paragraph?

[bashbaug]

I think we are on the same page.

Yes, I think we are aligned, great minds think alike!

What do you think of my last paragraph?

I think it's fine, though I also don't think we need to say anything specific about it. A device that doesn't support SPIR-V at all could return that it has aspect::spirv_queries, then unconditionally return false for each of the queries, and this could be a valid implementation.

FWIW, some applications might find it confusing for a device to have aspect::spirv_queries but not support SPIR-V at all, but we don't need to couple the two.

Interestingly, there is an open discussion in the CTS PR for the OpenCL extension whether it makes sense to support the extension without supporting SPIR-V: KhronosGroup/OpenCL-CTS#2409 (comment)

[Pennycook]

I'm not opposed to splitting the aspect in two, but I'd like to test something.

Because SYCL is a higher-level abstraction, I think it's nice sometimes to hide some of the complexity by bundling certain features and/or assuming a certain level of support. If/when the SPIR-V queries become supported widely, or perhaps even become a core/optional feature of OpenCL instead of an extension, would it then be acceptable for aspect::spirv to require an implementation of these queries?

I understand the motivation for keeping things separate while the extension is new -- because otherwise a lot of devices would be prevented from reporting aspect::spirv -- but I'd like to avoid death by a million extensions in the long-term.

[gmlueck]

Suggested change

	Returns `false` if the device does not support SPIR-V queries.
	Returns `false` if the device does not support kernels written in SPIR-V.

Consistent with my comment above, I'd change this last sentence like this. Same for the two queries below too.

[Pennycook]

This comment applies to the other APIs too, but I'll start the discussion here.

The OpenCL cl_khr_spirv_queries extension doesn't seem to work this way. I might have misunderstood something, but it looks like OpenCL exposes these queries via the regular device info mechanism. Is there a reason not to do that here?

If we followed the convention used by other device queries, I think this would become something like:

std::set<std::string> device.get_info<sycl::ext::oneapi::spirv_extended_instruction_sets>();

Similar to my suggestion above, we could define this in such a way that you get an exception when the query isn't supported at all, and otherwise the set contains all the supported instruction sets.

[bashbaug]

I'm very open to changing this, but the current interfaces are designed to be analogous to sycl_ext_oneapi_kernel_compiler_opencl, as we discussed. If we switch and use the device info mechanism for this extension, we should consider changing sycl_ext_oneapi_kernel_compiler_opencl to match, also.

[gmlueck]

I think we discussed this earlier when we added the OpenCL queries. My recollection is that we thought that SYCL was higher level then OpenCL, so it made sense to expose a more friendly API. I think ext_oneapi_supports_spirv_extended_instruction_set is more friendly than returning a list of instruction sets. If we returned a list, applications would need to search through that list in order to decide whether they can safely use a SPIR-V extended instruction.

[Pennycook]

But, conversely, if a developer wants to print all of the supported instruction sets, the proposed SYCL API requires them to maintain a list of all possible values, and then iterate through that list to check whether each value is supported. Building up the list that way seems especially silly if we expect the lower-level APIs will be returning a list for each query.

I think I would prefer if we had both, defined something like:

std::set<std::string> device::get_info<sycl::ext::oneapi::spirv_extended_instruction_sets>();

Returns: A set of all SPIR-V instruction sets supported by this device.
Throws: An exception with the feature_unsupported error code if this device does not have the spirv_queries aspect.

bool device::supports_spirv_extended_instruction_set(const std::string& name);
Returns: Equivalent to return device.get_info<sycl::ext::oneapi::spirv_extended_instruction_sets>().contains(name);

An added bonus of this is we don't have to explicitly explain what true and false mean, because that can be inferred from the behavior of contains.

---

Aside: Looking at this in the context of the OpenCL compiler extension, I think there might also be a way to combine the extensions. I'm not super keen to add a bunch of separate "supports_X", "supports_Y" functions for everything, when we could achieve the same thing with a single supports function, but I'm happy to defer that discussion or capture it as an issue.