Skip to content

Commit

Permalink
[SYCL][Doc] Update KernelProperties extension (#5343)
Browse files Browse the repository at this point in the history 
This commit makes several changes, to align with other extensions and
address feedback and concerns:

- To align with changes to the base property list extension, the _v
suffix has been dropped from property values.

- The "property_list" class is now called "properties". Some text in the
extension still refers to the concept of a "property list" because this
is easier to read.

- Using a member function in conjunction with a tag type prevents kernel
functors from accidentally implementing the properties interface, and
prevents naming conflicts with existing variables in user code.

Signed-off-by: John Pennycook <[email protected]>
  • Loading branch information
@Pennycook
Pennycook authored Jan 27, 2022
1 parent 33fdc58 commit b48be20
Showing 1 changed file with 91 additions and 74 deletions.
165 changes: 91 additions & 74 deletions sycl/doc/extensions/KernelProperties/KernelProperties.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,32 +23,26 @@ 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.

NOTE: This document is better viewed when rendered as html with asciidoctor.
GitHub does not render image icons.

This extension introduces a replacement for the kernel attributes defined in
Section 5.8.1 of the SYCL 2020 specification, in the form of a `property_list`
Section 5.8.1 of the SYCL 2020 specification, in the form of a property list
accepting properties with compile-time constant values.

== Notice

Copyright (c) 2021 Intel Corporation. All rights reserved.
Copyright (c) 2021-2022 Intel Corporation. All rights reserved.

== Status

Working Draft

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
feedback. Interfaces defined in this specification may not be implemented yet
or may be in a preliminary state. The specification itself may also change in
incompatible ways before it is finalized. Shipping software products should not
rely on APIs defined in this specification.

== Version

Built On: {docdate} +
Revision: 1

== Contributors
Expand All @@ -61,10 +55,10 @@ Roland Schulz, Intel

== Dependencies

This extension is written against the SYCL 2020 specification, Revision 3 and
This extension is written against the SYCL 2020 specification, Revision 4 and
the following extensions:

- https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/PropertyList/SYCL_EXT_ONEAPI_property_list.asciidoc[SYCL_EXT_ONEAPI_PROPERTY_LIST]
- https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/PropertyList/SYCL_EXT_ONEAPI_properties.asciidoc[SYCL_EXT_ONEAPI_PROPERTIES]

== Feature Test Macro

Expand Down Expand Up @@ -110,7 +104,7 @@ information, including:
performed at run-time and *after* compiling the kernel.

This extension proposes a replacement for these kernel attributes, in the form
of a `property_list` accepting properties with compile-time constant
of a property list accepting properties with compile-time constant
values, to address several of these issues.

== Kernel Properties
Expand All @@ -123,58 +117,60 @@ Section 5.8.1 of the SYCL 2020 specification. Note that deprecated attributes
namespace sycl {
namespace ext {
namespace oneapi {
namespace experimental {

// Corresponds to reqd_work_group_size
struct work_group_size {
struct work_group_size_key {
template <size_t... Dims>
using value_t = property_value<work_group_size, std::integral_constant<size_t, Dims>...>;
}; // work_group_size
using value_t = property_value<work_group_size_key, std::integral_constant<size_t, Dims>...>;
}; // work_group_size_key

// Corresponds to work_group_size_hint
struct work_group_size_hint {
struct work_group_size_hint_key {
template <size_t... Dims>
using value_t = property_value<work_group_size_hint, std::integral_constant<size_t, Dims>...>;
}; // work_group_size_hint
using value_t = property_value<work_group_size_hint_key, std::integral_constant<size_t, Dims>...>;
}; // work_group_size_hint_key

// Corresponds to reqd_sub_group_size
struct sub_group_size {
struct sub_group_size_key {
template <uint32_t Size>
using value_t = property_value<sub_group_size, std::integral_constant<uint32_t, Size>>;
}; // sub_group_size
using value_t = property_value<sub_group_size_key, std::integral_constant<uint32_t, Size>>;
}; // sub_group_size_key

// Corresponds to device_has
struct device_has {
struct device_has_key {
template <sycl::aspect... Aspects>
using value_t = property_value<device_has, std::integral_constant<sycl::aspect, Aspects>...>;
}; // device_has
using value_t = property_value<device_has_key, std::integral_constant<sycl::aspect, Aspects>...>;
}; // device_has_key

template <size_t... Dims>
struct property_value<work_group_size, std::integral_constant<size_t, Dims>...> {
struct property_value<work_group_size_key, std::integral_constant<size_t, Dims>...> {
constexpr size_t operator[](int dim);
};

template <size_t... Dims>
struct property_value<work_group_size_hint, std::integral_constant<size_t, Dims>...> {
struct property_value<work_group_size_hint_key, std::integral_constant<size_t, Dims>...> {
constexpr size_t operator[](int dim);
};

template <sycl::aspect... Aspects>
struct property_value<device_has, std::integral_constant<sycl::aspect, Aspects>...> {
struct property_value<device_has_key, std::integral_constant<sycl::aspect, Aspects>...> {
static constexpr std::array<sycl::aspect, sizeof...(Aspects)> value;
};

template <size_t... Dims>
inline constexpr work_group_size::value_t<Dims...> work_group_size_v;
inline constexpr work_group_size_key::value_t<Dims...> work_group_size;

template <size_t... Dims>
inline constexpr work_group_size_hint::value_t<Dims...> work_group_size_hint_v;
inline constexpr work_group_size_hint_key::value_t<Dims...> work_group_size_hint;

template <uint32_t Size>
inline constexpr sub_group_size::value_t<Size> sub_group_size_v;
inline constexpr sub_group_size_key::value_t<Size> sub_group_size;

template <sycl::aspect... Aspects>
inline constexpr device_has::value_t<Aspects...> device_has_v;
inline constexpr device_has_key::value_t<Aspects...> device_has;

} // namespace experimental
} // namespace oneapi
} // namespace ext
} // namespace sycl
Expand Down Expand Up @@ -221,15 +217,15 @@ SYCL implementations may introduce additional kernel properties. If any
combinations of kernel attributes are invalid, this must be clearly documented
as part of the new kernel property definition.

== Adding a `property_list` to a Kernel Launch
== Adding a Property List to a Kernel Launch

To enable properties to be associated with kernels, this extension adds
new overloads to each of the variants of `single_task`, `parallel_for` and
`parallel_for_work_group` defined in the `queue` and `handler` classes. These
new overloads accept a `sycl::ext::oneapi::property_list` argument. For
variants accepting a parameter pack, the `sycl::ext::oneapi::property_list`
new overloads accept a `sycl::ext::oneapi::experimental::properties` argument. For
variants accepting a parameter pack, the `sycl::ext::oneapi::experimental::properties`
argument is inserted immediately prior to the parameter pack; for variants not
accepting a parameter pack, the `sycl::ext::oneapi::property_list` argument is
accepting a parameter pack, the `sycl::ext::oneapi::experimental::properties` argument is
inserted immediately prior to the kernel function.

The overloads introduced by this extension are listed below:
Expand Down Expand Up @@ -313,15 +309,15 @@ class handler {
}
```

Passing properties as an argument in this way allows properties to be
Passing a property list as an argument in this way allows properties to be
associated with a kernel function without modifying its type. This enables
the same kernel function (e.g. a lambda) to be submitted multiple times with
different properties, or for libraries building on SYCL to add properties
(e.g. for performance reasons) to user-provided kernel functions.

All the properties defined in this extension have compile-time values. However,
an implementation may support additional properties which could have run-time
values. When this occurs, the `properties` parameter may be a `property_list`
values. When this occurs, the `properties` parameter may be a property list
containing a mix of both run-time and compile-time values, and a SYCL
implementation should respect both run-time and compile-time information when
determining the correct way to launch a kernel. However, only compile-time
Expand All @@ -331,28 +327,47 @@ A simple example of using this extension to set a required work-group size
and required sub-group size is given below:

```c++
sycl::ext::oneapi::property_list properties{sycl::ext::oneapi::work_group_size_v<8, 8>,
sycl::ext::oneapi::sub_group_size_v<8>};
sycl::ext::oneapi::experimental::properties properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
sycl::ext::oneapi::experimental::sub_group_size<8>};
q.parallel_for(range<2>{16, 16}, properties, [=](id<2> i) {
a[i] = b[i] + c[i];
}).wait();
```

== Encoding Properties into a Kernel
== Embedding Properties into a Kernel

In other situations it may be useful to encode a kernel's properties directly
In other situations it may be useful to embed a kernel's properties directly
into its type, to ensure that a kernel cannot be launched without a property
that it depends upon for correctness.

To enable this use-case, this extension adds a mechanism for implementations to
extract a property list from a kernel functor, if a kernel functor declares
a `property_list` member variable named `properties`. Note that this member
variable must be `static constexpr`, and kernel functors can therefore only
encode properties with compile-time values.
a member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
tag type and returning an instance of `sycl::ext::oneapi::experimental::properties`.

```c++
namespace sycl {
namespace ext {
namespace oneapi {
namespace experimental {

struct properties_tag {};

}
}
}
}
```

NOTE: https://wg21.link/p1895[P1895] proposes a function called `tag_invoke`
as a general mechanism for customization points that could be used as a
replacement for the `get` function proposed here. If `tag_invoke` becomes
a feature in a future version of {cpp}, a future version of this extension
may expose a new interface compatible with `tag_invoke`.

NOTE: The attribute mechanism in SYCL 2020 allows for different kernel
attributes to be applied to different call operators within the same
functor. The `property_list` member variable applies to all call operators in
functor. An embedded property list applies to all call operators in
the functor.

The example below shows how the kernel from the previous section could be
Expand All @@ -370,9 +385,10 @@ struct KernelFunctor {
a[i] = b[i] + c[i];
}

static constexpr auto properties =
sycl::ext::oneapi::property_list{sycl::ext::oneapi::work_group_size_v<8, 8>,
sycl::ext::oneapi::sub_group_size_v<8>};
auto get(sycl::ext::oneapi::experimental::properties_tag) {
return sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
sycl::ext::oneapi::experimental::sub_group_size<8>};
}

sycl::accessor<int, 2> a;
sycl::accessor<int, 2> b;
Expand All @@ -385,18 +401,18 @@ struct KernelFunctor {
q.parallel_for(range<2>{16, 16}, KernelFunctor(a, b, c)).wait();
```

If a kernel functor with a `property_list` member variable is enqueued for
execution using an invocation function with a `property_list` argument,
the kernel is launched as-if the member variable and argument were combined. If
the combined list contains any invalid combinations of properties, then this is
an error: invalid combinations that can be detected at compile-time should be
reported via a diagnostic; invalid combinations that can only be detected at
run-time should result in an implementation throwing an `exception` with the
`errc::invalid` error code.
If a kernel functor with embedded properties is enqueued for execution using an
invocation function with a property list argument, the kernel is launched as-if
the embedded properties and argument were combined. If the combined list
contains any invalid combinations of properties, then this is an error: invalid
combinations that can be detected at compile-time should be reported via a
diagnostic; invalid combinations that can only be detected at run-time should
result in an implementation throwing an `exception` with the `errc::invalid`
error code.

== Querying Properties in a Compiled Kernel

Any properties encoded into a kernel type via a property list are reflected
Any properties embedded into a kernel type via a property list are reflected
in the results of a call to `kernel::get_info` with the
`info::kernel::attributes` information descriptor, as if the corresponding
attribute from the SYCL 2020 specification had been applied to the kernel
Expand All @@ -408,9 +424,9 @@ The SYCL 2020 `sycl::device_has` attribute can be applied to the declaration
of a non-kernel device function, to assert that the device function uses a
specific set of optional features. This extension provides a mechanism exposing
similar behavior, allowing for kernel properties to be associated with
a function via the `SYCL_EXT_ONEAPI_PROPERTY` macro. Each instance of the
`SYCL_EXT_ONEAPI_PROPERTY` macro accepts one argument, corresponding to a
single property value.
a function via the `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY` macro. Each instance of
the `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY` macro accepts one argument,
corresponding to a single property value.

NOTE: Due to limitations of the C preprocessor, property value expressions
containing commas (e.g. due to template arguments) must be enclosed in
Expand All @@ -420,12 +436,13 @@ The example below shows a function that uses two optional features,
corresponding to the `fp16` and `atomic64` aspects.

```c++
SYCL_EXT_ONEAPI_PROPERTY((sycl::device_has_v<aspect::fp16, aspect::atomic64>))
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::device_has<aspect::fp16, aspect::atomic64>))
void foo();
```

The table below describes the effects of associating each kernel property
with a non-kernel device function via the `SYCL_EXT_ONEAPI_PROPERTY` macro.
with a non-kernel device function via the `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`
macro.

|===
|Property|Description
Expand All @@ -438,14 +455,14 @@ with a non-kernel device function via the `SYCL_EXT_ONEAPI_PROPERTY` macro.

|===

The `SYCL_EXT_ONEAPI_PROPERTY` macro can be used alongside the
The `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY` macro can be used alongside the
`SYCL_EXTERNAL` macro, and the macros may be specified in any order.
Whenever `SYCL_EXTERNAL` is used, there are two relevant translation units: the
translation unit that _defines_ the function and the translation unit that
_calls_ the function. If a given `SYCL_EXTERNAL` function _F_ is defined in
one translation unit with a set of properties _P_, then all other translation
units that declare that same function _F_ must list the same set of properties
_P_ via the `SYCL_EXT_ONEAPI_PROPERTY` macro. Programs which fail to do this
_P_ via the `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY` macro. Programs which fail to do this
are ill-formed, but no diagnostic is required.

== Issues
Expand All @@ -461,14 +478,14 @@ new properties, for example `device_has_all_of` and `device_has_any_of`:
device_has_any_of<device_has<aspect::fp16, device_has<aspect::fp64>>`.
--

. How should the `property_list` member variable behave with inheritance?
. How should an embedded property list behave with inheritance?
+
--
*UNRESOLVED*: The specification currently allows for a class to inspect the
`property_list` member variable from its base class(es) and construct a new
`property_list` member variable that applies to all call operators. Associating
different properties with different call operators via inheritance has the
potential to be confusing and would increase implementation complexity.
*RESOLVED*: The specification currently allows for a class to inspect the
property list embedded into its base class(es) and construct a new property
list that applies to all call operators. Associating different properties with
different call operators via inheritance has the potential to be confusing and
would increase implementation complexity.
--

//. asd
Expand Down

0 comments on commit b48be20

Please sign in to comment.