Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Clarify "is_compatible" #381

Merged
merged 1 commit into from
May 4, 2023
Merged

Clarify "is_compatible" #381

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 6 additions & 1 deletion adoc/chapters/programming_interface.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14999,7 +14999,12 @@ application is compatible with a device unless:
in <<sec:optional-kernel-features>>; or
* It is decorated with a [code]#[[sycl::device_has()]]# {cpp} attribute that
lists an aspect that is not supported by the device, as described in
<<sec:kernel.attributes>>.
<<sec:kernel.attributes>>; or
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

is_compatible doesn't really sound like "was this kernel compiled for this device?". With this change, users won't have an easy way to understand whether or not all optional features used in a kernel are supported by the device, because is_compatible could return false for other reasons.

Would it be better to instead expand definition of has_kernel_bundle to say that it may return false if compilation environment doesn't support the device?

Kernel definition/properties, i.e. list of aspects used by it should not change depending on compilation environment, because it is encoded in the source code. Content of kernel bundle and list of available device images, however, is not defined by the source code alone and can also be heavily dependent on compilation environment.

Tagging @rolandschulz

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the intent of is_compatible was that you can use it to determine whether a kernel can be submitted to a device. If we think that was the intent, then I believe this clarification makes sense.

I agree that an application could have some kernel K that uses features that are compatible with some device D, but due to the compilation environment K cannot be run on device D. This change to the spec makes it impossible for the application to test that K would be compatible with device D if the application was compiled with some other set of options. However, I'm not sure that such a query is useful. It seems like the thing the application cares about is whether K can be run on D, not some hypothetical question about whether K could be run on D if the application was compiled differently.

Note that it's less convenient to use has_kernel_bundle because the application need to create a kernel bundle in order to use that API. It's much more convenient to use is_compatible because you only need the kernel's type-name.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

However, I'm not sure that such a query is useful. It seems like the thing the application cares about is whether K can be run on D, not some hypothetical question about whether K could be run on D if the application was compiled differently.

Let's say you provide a library which is AOT-compiled for a certain list of devices. You may want to improve user experience of the library users by emitting a better error saying that they need to re-compile the library with different flags to make it work on that device; instead of having some unspecified exception at runtime.
However, I agree that it is kind of an artificial case, because as a library author you could somehow configure your compilation environment to ensure that there are some fallback paths for non-target devices; or simply have clear documentation about device requirements.

Note that it's less convenient to use has_kernel_bundle because the application need to create a kernel bundle in order to use that API.

You don't need a kernel bundle to query if a kernel bundle exists. has_kernel_bundle accepts context, vector of kernel_id and optionally list of devices.

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think both queries are useful. It is only hypothetical for run decision, but not for error reporting. An application might want to report "This kernel could run on the device you have available. But I was compiled without proper support. Please recompile me for best performance.".

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think I agree that since there are a few different reasons for why is_compatible may return false, it would be useful to have a way to expose the reason, I'm not sure how useful it would be as a mechanism for dispatching different kernels, but it could at least be a useful debug feature, to be able to identify exactly why the kernel couldn't be run on a certain device. Saying that, if we were to do this, I think that could be a different PR that doesn't need to hold up this one.

* The translation unit containing the kernel was compiled in a compilation
environment that does not support the device. Each implementation defines
AerialMantis marked this conversation as resolved.
Show resolved Hide resolved
the specific criteria for which devices are supported in its compilation
environment. For example, this might be dependent on options passed to the
compiler.

A device built-in kernel is only compatible with the device for which it is
built-in.
Expand Down