Optional Kernel Features

[hdelan]

Specification Version

SYCL 2020 (Revision 11)

Section Number(s)

Section 5.7. Optional Kernel Features

Issue Description

In SYCL 2020, fp16, fp64, and atomic64 are clearly optional device aspects. I’m less clear on the statement that reqd_work_group_size and reqd_sub_group_size are “optional kernel features.”

The term “optional” seems ambiguous:

1. The attribute mechanism itself is optional (an implementation may ignore it or fail at launch/runtime, provided kernels using the attribute can still be compiled),
   or
2. The mechanism is required, but specific requested sizes are optional per device (i.e., only some reqd_sub_group_size(N) / reqd_work_group_size(...) values are supported).

The wording “compiled in a way that might not run on all devices” suggests (2), but calling these attributes “optional kernel features” can be read as (1).

Could the spec clarify which interpretation is intended, and what conformance requires (e.g., must implementations enforce requested subgroup/workgroup sizes against device capabilities, or may they ignore the attributes)?

Context: implementations currently differ (honor vs ignore), e.g. intel/llvm#11687

[gmlueck]

Perhaps we should clarify this sentence in that section:

In order to guarantee source code portability of SYCL applications that use optional kernel features, all SYCL implementations must be able to compile device code that uses these optional features regardless of whether the implementation supports the features on any of its devices.

to:

In order to guarantee source code portability of SYCL applications that use optional kernel features, all SYCL implementations must be able to compile device code that uses these optional features regardless of whether the implementation supports the features on any of its devices. For the two C++ attributes listed above, all implementations must be able to compile kernels that are decorated with these attributes regardless of whether the implementation supports the specified work-group or sub-group size.

[hdelan]

Hi @gmlueck . Does this mean that the attribute can be ignored, provided that compilation doesn't fail? i.e. if my device has one possible max sg size of 16, when launching a kernel with reqd_sub_group_size(32) must the kernel launch fail or can the kernel still run with the sub group size of 16?

[gmlueck]

Users typically decorate a kernel with reqd_sub_group_size because the body of their kernel makes an assumption that the sub-group size has a specific value. If you run with a different sub-group size, the kernel's logic will probably be wrong, and it will produce a wrong result or otherwise run incorrectly.

[hdelan]

I agree that code would likely go badly wrong if this is ignored, but I am building on DPC++'s precedent of either honouring or ignoring the attribute (see above), as well as C++'s allowing nonstandard attributes to be ignored (at least this page says that only standardized attributes must not be ignored).

I think the specification needs to clarify whether this can be ignored or not. If the implementation must not ignore this attribute then I think library only implementations would have a hard time implementing this.

[gmlueck]

I think the next paragraph already clarifies this:

Of course, applications that make use of optional kernel features should ensure that a kernel using such a feature is submitted only to a device that supports the feature. [...] If an application fails to do this, the implementation throws an exception with the errc::kernel_not_supported error code from the kernel invocation command (e.g. parallel_for()).

That said, the WG has decided that the CTS should not contain negative tests which verify that an implementation diagnoses an error for invalid applications.

[hdelan]

a kernel using such a feature

Does feature mean reqd_sub_group_size in general? Or does it mean reqd_sub_group_size for a particular value?

[gmlueck]

The specification of reqd_sub_group_size in section 5.8.1. Kernel attributes seems even clearer:

Indicates that the kernel must be compiled and executed with the specified sub-group size.
[...]
If a kernel is decorated with this attribute and then submitted to a device that does not support the sub-group size or if the kernel uses a feature that the device does not support with this sub-group size, the implementation must throw a synchronous exception with the errc::kernel_not_supported error code.

I do not think this wording allows an implementation to ignore the attribute and compile the kernel with a different sub-group size.

[hdelan]

Thanks for clarification. So "Optional" in the title in section 5.7.1 means optional for some specific value, not in general.

I think this use of the word "Optional" is ambiguous, especially given that C++ allows non standard attributes to be ignored.

I propose reqd_work_group_size and reqd_sub_group_size are not introduced as optional kernel features but rather "Additional Kernel Features", which spells out that these attributes are considered standard SYCL and as such cannot be ignored, similar to standard C++ attributes. I think section 5.8.1 is clear and sufficient otherwise.

As an aside, what does this mean for library only implementations?

[gmlueck]

I propose reqd_work_group_size and reqd_sub_group_size are not introduced as optional kernel features but rather "Additional Kernel Features", which spells out that these attributes are considered standard SYCL and as such cannot be ignored

I'm not opposed to changing the name of the section, but I think whatever name we choose should apply to all of the features currently described in section 5.7. For example, the sycl::half type is currently specified as an "optional kernel feature". Implementation must still be able to compile device code that uses this type, so it's not optional in that sense. However, kernels using that type might not be able to run on all devices. In this sense, it is optional. The attributes have the same behavior: the compiler must consume them, but kernels decorated with them may not run on all devices.

The term "Additional kernel feature" seems more confusing to me because it doesn't convey the notion that the feature might not be supported on all devices.

Note that the compiler need not generate meaningful code for a kernel that uses an "optional" feature if the feature is not supported for your target device. For example, DPC++ just deletes the entire kernel body in this case. The important thing is that the compiler does not issue a diagnostic or otherwise fail, and that the runtime throws an exception if the application mistakenly submits the kernel to a device that does not support the feature.

As an aside, what does this mean for library only implementations?

We have discussed this in the past, and we recognize that it would be hard to implement with a library-only implementation. DPC++ has an extension that provides an alternate way to decorate a kernel with a required work-group or a required sub-group size. See the sycl_ext_oneapi_kernel_properties extension. The Khronos WG is actively discussing compile-time properties in general as a possible KHR extension. Once we agree on this basic property support, I expect we will proposed kernel properties also as a KHR. My gut feeling is that replacing the C++ attributes with kernel properties should make these easier to implement in a library-only implementation.

[TApplencourt]

• On the agenda for next week (2026/04/09, how to do the clarification.

[TApplencourt]

For the two C++ attributes listed above, all implementations must be able to compile kernels that are decorated with these attributes regardless of whether the implementation supports the specified work-group or sub-group size.

Sound good.

2nd clarification to do:

• They cannot be ignored, but the value possible are a runtime set. The value possible can be NULL, they are "Optional" in that sense.

Maybe Hugh will draft a clarificaiton PR :) Thanks!

[VerenaBeckham]

I concluded that, if you e.g. enumerate reqd_work_group_size(n) for all n then each of them counts as a "feature", and each of these "features" is optional. "Optional" still as "must compile" but "may give an error at run time".