Skip to content

[SYCL] Unify compile/run-time properties more #12245

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

Closed
wants to merge 4 commits into from
Closed

[SYCL] Unify compile/run-time properties more #12245

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
7d22240
[SYCL] Unify compile/run-time properties more
rolandschulz Dec 25, 2023
e460012
Fix annotated_usm test
rolandschulz Dec 25, 2023
d187aa0
format
rolandschulz Dec 25, 2023
2c85623
Fix invalid runtime test
rolandschulz Dec 25, 2023
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
25 changes: 15 additions & 10 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_properties.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,8 +99,7 @@ The intention is to provide a robust mechanism with which to pass compile-time-c
property:: A property is represented by a key and value. Properties can be used to provide extra values to classes or functions.

property value:: An object of the property value class. A property value has zero or more property parameters.
For runtime properties the value type is the same as the key type.
For compile time properties the value type is given by the `value_t` type alias of the key type.
The value type is given by the `value_t` type alias of the key type.

property key:: A class representing the property key. It is used to query properties.

Expand Down Expand Up @@ -138,27 +137,33 @@ value to determine which of the extension's APIs the implementation supports.

Properties have a value and key type,
and by convention, these classes are declared in the root of the
`sycl::ext::oneapi::experimental` namespace. For a runtime property the key and value types are the same and the name of the property value
class has no suffix. A runtime property value typically has a constructor
`sycl::ext::oneapi::experimental` namespace.
The value type is a template specialization of `property_value`.
The property key class contains a `value_t` alias which is templated on the property parameters.

A runtime property value typically has a constructor
which takes the value(s) of the properties and member function(s) which return those values.

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

// The runtime property key
struct foo_key {
using value_t = property_value<foo_key>;
};
// This is a runtime property value with one integer parameter.
// The name of the property value class is the the name of the property without any suffix.
struct foo {
foo(int);
using foo = property_value<foo_key>;
template <> struct property_value<foo_key> {
using key_t = foo_key;
constexpr property_value(int v) : value(v) {}
int value;
};
// A runtime property key is an alias to the value type.
using foo_key = foo;

} // namespace experimental::oneapi::ext::sycl
```

For compile-time constant parameters the value type is a template specialization of `property_value`.
The property key class contains a `value_t` alias which is templated on the property parameters. The `property_value` class holds the
For compile-time constant parameters the `property_value` class holds the
values of the compile-time parameters as template arguments. The parameters to a compile-
time-constant property can be either types or non-type values.
The implementation provides a variable with the property value type. The variable has the name of the property without a suffix.
Expand Down
17 changes: 12 additions & 5 deletions sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,12 +26,19 @@ inline constexpr cache_config_enum large_slm =
inline constexpr cache_config_enum large_data =
cache_config_enum::large_data;

struct cache_config {
cache_config(cache_config_enum v) : value(v) {}
cache_config_enum value;
struct cache_config_key {
using value_t = oneapi::experimental::property_value<cache_config_key>;
};

using cache_config_key = cache_config;
using cache_config = cache_config_key::value_t;
} // namespace ext::intel::experimental
namespace ext::oneapi::experimental {
template <> struct property_value<intel::experimental::cache_config_key> {
using key_t = intel::experimental::cache_config_key;
property_value(intel::experimental::cache_config_enum v) : value(v) {}
intel::experimental::cache_config_enum value;
};
} // namespace ext::oneapi::experimental
namespace ext::intel::experimental {
Copy link
Contributor

Choose a reason for hiding this comment

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

You can see here that the definition of a typical runtime property expanded from 5 lines to about 12. However, the description of the PR says that a goal is to simplify the implementation. What is getting simplified? It seems like the definition of a runtime property is more complex, not less.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

What is getting simplified?

The main motivation is to simplify has_property and get_property. I'm worried that if/when properties are getting used a lot, that the current implementation will be bad for compilation speed. Currently we store the properties as a sorted list of values in the template argument of properties. And because not all have the same form (compile-time are of the form property_value<key, ...> and runtime are their own type) we have to iterate over the whole list every time.
With this change all property have the same form of property_value<key, ...>. This turns the sorted list into a sorted map. And with that efficient map lookup with mp_map_find (doc, impl) is possible.

If we don't like this change we could instead change the storage inside the template argument to a sorted set of runtime properties and a sorted map of compile-time properties (doesn't require any spec changes). Both would have fast lookup. We would just need to split the properties in the two group prior to sorting whenever we sort.

A 2nd motivation was to reduce where we need to distinguish between the two groups of properties in the implementation. With the alternative implementation this would become more not less.

For the spec we could allow both implementation strategies (runtime property values are of the form property_value<key, ...> or runtime property values/keys the same type). Then this becomes purely a quality of implementation question.

The property_value<key, ...> form allows a runtime property value to store some of its values as compile-time values. And those values can be queried at compile time. I'm not aware of use-cases for properties which have both runtime values (members) and also compile-time values (template args). So I don't know whether this is an advantage.

Copy link
Contributor Author

@rolandschulz rolandschulz Jan 5, 2024 •
edited
Loading

Choose a reason for hiding this comment

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

I created a prototype of what I think a better implementation should roughly look like: https://godbolt.org/z/PvaWzTnnW. For the PRs I was trying to break the changes into incremental steps to help with review. But hopefully this shows what is possible after all changes. In my view it makes the implementation of the properties class much shorter and also easier to define both runtime and compile-time properties. This implementation currently isn't conformant, because the spec contains several implementation details which aren't valid for this strategy. To make it conformant we would need to remove (or change):

  • Runtime properties key/value are the same type (this change)
  • Remove key_t (Remove key_t #12246, not necessary and makes defining runtime properties simpler )
  • Remove static from get_property (simplifies implementation, we don't have any use case of non-device copyable properties)
  • Remove value_t (not necessary and simplifies defining properties a lot)

None of the changes are necessary. We could do none or only some and still simplify the implementation. It just that each of them make it simpler to either implement the properties class or add new properties.

I would be interested in any feedback regarding:

  • Do we want the spec to mandate one implementation strategy or not?
  • Which implementation strategy do we want to use?
  • Do you prefer for review a set of incremental changes or one large PR?

Copy link
Contributor

Choose a reason for hiding this comment

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

I prefer multiple incremental PRs, personally, and I don't think we should specify anything about implementation that we don't have to.


inline bool operator==(const cache_config &lhs,
const cache_config &rhs) {
Expand Down
4 changes: 2 additions & 2 deletions sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -106,8 +106,8 @@ struct ValidAllocPropertyList<T, detail::properties_t<Prop, Props...>>
is_valid_property<T *, Prop>::value,
"Found invalid compile-time property in the property list.");
// check if a runtime property is valid for malloc
static_assert(!detail::IsRuntimeProperty<Prop>::value ||
IsRuntimePropertyValid<Prop>::value,
static_assert(!detail::IsRuntimePropertyValue<Prop>::value ||
IsRuntimePropertyValid<typename Prop::key_t>::value,
"Found invalid runtime property in the property list.");
};

Expand Down
17 changes: 8 additions & 9 deletions sycl/include/sycl/ext/oneapi/properties/properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,7 +77,7 @@ template <typename... Ts> struct RuntimePropertyStorage<std::tuple<Ts...>> {
};
template <typename T, typename... Ts>
struct RuntimePropertyStorage<std::tuple<T, Ts...>>
: std::conditional_t<IsRuntimeProperty<T>::value,
: std::conditional_t<IsRuntimePropertyValue<T>::value,
PrependTuple<T, typename RuntimePropertyStorage<
std::tuple<Ts...>>::type>,
RuntimePropertyStorage<std::tuple<Ts...>>> {};
Expand Down Expand Up @@ -149,9 +149,9 @@ template <typename PropertiesT> class properties {
template <typename PropertyT>
typename std::enable_if_t<detail::IsRuntimeProperty<PropertyT>::value &&
has_property<PropertyT>(),
PropertyT>
typename PropertyT::value_t>
get_property() const {
return std::get<PropertyT>(Storage);
return std::get<typename PropertyT::value_t>(Storage);
}

template <typename PropertyT>
Expand Down Expand Up @@ -254,16 +254,15 @@ struct all_props_are_keys_of<
template <typename SyclT, typename PropT>
struct all_props_are_keys_of<
SyclT, ext::oneapi::experimental::properties<std::tuple<PropT>>>
: std::bool_constant<
ext::oneapi::experimental::is_property_key_of<PropT, SyclT>::value> {
};
: std::bool_constant<ext::oneapi::experimental::is_property_key_of<
typename PropT::key_t, SyclT>::value> {};

template <typename SyclT, typename PropT, typename... PropTs>
struct all_props_are_keys_of<
SyclT, ext::oneapi::experimental::properties<std::tuple<PropT, PropTs...>>>
: std::bool_constant<
ext::oneapi::experimental::is_property_key_of<PropT, SyclT>::value &&
all_props_are_keys_of<SyclT, PropTs...>()> {};
: std::bool_constant<ext::oneapi::experimental::is_property_key_of<
typename PropT::key_t, SyclT>::value &&
all_props_are_keys_of<SyclT, PropTs...>()> {};

} // namespace detail
} // namespace ext::oneapi::experimental
Expand Down
4 changes: 3 additions & 1 deletion sycl/include/sycl/ext/oneapi/properties/property_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,8 @@ struct HasValue<T, decltype((void)T::value, 0)> : std::true_type {};
template <typename PropertyT>
struct IsCompileTimePropertyValue : std::false_type {};

template <typename PropertyT>
struct IsRuntimePropertyValue : std::false_type {};
// Checks if a type is either a runtime property or if it is a compile-time
// property
template <typename T> struct IsProperty {
Expand All @@ -73,7 +75,7 @@ template <typename T> struct IsProperty {
// property_value with a valid compile-time property
template <typename T> struct IsPropertyValue {
static constexpr bool value =
IsRuntimeProperty<T>::value || IsCompileTimePropertyValue<T>::value;
IsRuntimePropertyValue<T>::value || IsCompileTimePropertyValue<T>::value;
};

// Checks that all types in a tuple are valid properties.
Expand Down
18 changes: 9 additions & 9 deletions sycl/include/sycl/ext/oneapi/properties/property_value.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,15 +57,11 @@ operator!=(const property_value<PropertyT, A...> &,
return (!std::is_same<A, B>::value || ...);
}

template <typename V, typename = void> struct is_property_value {
static constexpr bool value =
detail::IsRuntimeProperty<V>::value && is_property_key<V>::value;
};
template <typename V, typename O, typename = void> struct is_property_value_of {
static constexpr bool value =
detail::IsRuntimeProperty<V>::value && is_property_key_of<V, O>::value;
};
// Specialization for compile-time-constant properties
template <typename V, typename = void>
struct is_property_value : std::false_type {};
template <typename V, typename O, typename = void>
struct is_property_value_of : std::false_type {};
// Specialization for properties
template <typename V>
struct is_property_value<V, std::void_t<typename V::key_t>>
: is_property_key<typename V::key_t> {};
Expand All @@ -85,6 +81,10 @@ template <typename PropertyT, typename... PropertyValueTs>
struct IsCompileTimePropertyValue<property_value<PropertyT, PropertyValueTs...>>
: IsCompileTimeProperty<PropertyT> {};

template <typename PropertyT, typename... PropertyValueTs>
struct IsRuntimePropertyValue<property_value<PropertyT, PropertyValueTs...>>
: IsRuntimeProperty<PropertyT> {};

} // namespace detail
} // namespace ext::oneapi::experimental
} // namespace _V1
Expand Down
31 changes: 19 additions & 12 deletions sycl/include/sycl/kernel_bundle.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -813,12 +813,16 @@ namespace ext::oneapi::experimental {
/////////////////////////
// PropertyT syclex::build_options
/////////////////////////
struct build_options {
struct build_options_key {
using value_t = property_value<build_options_key>;
};
using build_options = property_value<build_options_key>;
template <> struct property_value<build_options_key> {
using key_t = build_options_key;
std::vector<std::string> opts;
build_options(const std::string &optsArg) : opts{optsArg} {}
build_options(const std::vector<std::string> &optsArg) : opts(optsArg) {}
property_value(const std::string &optsArg) : opts{optsArg} {}
property_value(const std::vector<std::string> &optsArg) : opts(optsArg) {}
};
using build_options_key = build_options;

template <> struct is_property_key<build_options_key> : std::true_type {};

Expand Down Expand Up @@ -847,12 +851,15 @@ struct IsCompileTimeProperty<sycl::ext::oneapi::experimental::build_options_key>
/////////////////////////
// PropertyT syclex::save_log
/////////////////////////
struct save_log {
struct save_log_key {
using value_t = property_value<save_log_key>;
};
using save_log = property_value<save_log_key>;
template <> struct property_value<save_log_key> {
using key_t = save_log_key;
std::string *log;
save_log(std::string *logArg) : log(logArg) {}
property_value(std::string *logArg) : log(logArg) {}
};
using save_log_key = save_log;

template <> struct is_property_key<save_log_key> : std::true_type {};

template <>
Expand Down Expand Up @@ -917,11 +924,11 @@ build(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
const std::vector<device> &Devices, PropertyListT props = {}) {
std::vector<std::string> BuildOptionsVec;
std::string *LogPtr = nullptr;
if constexpr (props.template has_property<build_options>()) {
BuildOptionsVec = props.template get_property<build_options>().opts;
if constexpr (props.template has_property<build_options_key>()) {
BuildOptionsVec = props.template get_property<build_options_key>().opts;
}
if constexpr (props.template has_property<save_log>()) {
LogPtr = props.template get_property<save_log>().log;
if constexpr (props.template has_property<save_log_key>()) {
LogPtr = props.template get_property<save_log_key>().log;
}
return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr);
}
Expand Down
127 changes: 56 additions & 71 deletions sycl/test/extensions/annotated_usm/fake_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,85 +69,70 @@ template <> struct IsCompileTimeProperty<boo_key> : std::true_type {};

// Runtime properties
enum foo_enum : unsigned { a, b, c };
struct foo {
constexpr foo(foo_enum v) : value(v) {}
struct foo_key {
using value_t = property_value<foo_key>;
};
using foo = property_value<foo_key>;
template <> struct property_value<foo_key> {
using key_t = foo_key;
constexpr property_value(foo_enum v) : value(v) {}
foo_enum value;
};

struct foz {
struct foz_key {
using value_t = property_value<foz_key>;
};
using foz = property_value<foz_key>;
template <> struct property_value<foz_key> {
using key_t = foz_key;
float value1;
bool value2;
};

struct rt_prop1 {};
struct rt_prop2 {};
struct rt_prop3 {};
struct rt_prop4 {};
struct rt_prop5 {};
struct rt_prop6 {};
struct rt_prop7 {};
struct rt_prop8 {};
struct rt_prop9 {};
struct rt_prop10 {};
struct rt_prop11 {};
struct rt_prop12 {};
struct rt_prop13 {};
struct rt_prop14 {};
struct rt_prop15 {};
struct rt_prop16 {};
struct rt_prop17 {};
struct rt_prop18 {};
struct rt_prop19 {};
struct rt_prop20 {};
struct rt_prop21 {};
struct rt_prop22 {};
struct rt_prop23 {};
struct rt_prop24 {};
struct rt_prop25 {};
struct rt_prop26 {};
struct rt_prop27 {};
struct rt_prop28 {};
struct rt_prop29 {};
struct rt_prop30 {};
struct rt_prop31 {};
struct rt_prop32 {};
struct rt_prop33 {};
#define rt_prop(N) \
struct rt_prop##N##_key { \
using value_t = property_value<rt_prop##N##_key>; \
}; \
using rt_prop##N = property_value<rt_prop##N##_key>; \
template <> struct property_value<rt_prop##N##_key> { \
using key_t = rt_prop##N##_key; \
}

rt_prop(1);
rt_prop(2);
rt_prop(3);
rt_prop(4);
rt_prop(5);
rt_prop(6);
rt_prop(7);
rt_prop(8);
rt_prop(9);
rt_prop(10);
rt_prop(11);
rt_prop(12);
rt_prop(13);
rt_prop(14);
rt_prop(15);
rt_prop(16);
rt_prop(17);
rt_prop(18);
rt_prop(19);
rt_prop(20);
rt_prop(21);
rt_prop(22);
rt_prop(23);
rt_prop(24);
rt_prop(25);
rt_prop(26);
rt_prop(27);
rt_prop(28);
rt_prop(29);
rt_prop(30);
rt_prop(31);
rt_prop(32);
rt_prop(33);

using foo_key = foo;
using foz_key = foz;
using rt_prop1_key = rt_prop1;
using rt_prop2_key = rt_prop2;
using rt_prop3_key = rt_prop3;
using rt_prop4_key = rt_prop4;
using rt_prop5_key = rt_prop5;
using rt_prop6_key = rt_prop6;
using rt_prop7_key = rt_prop7;
using rt_prop8_key = rt_prop8;
using rt_prop9_key = rt_prop9;
using rt_prop10_key = rt_prop10;
using rt_prop11_key = rt_prop11;
using rt_prop12_key = rt_prop12;
using rt_prop13_key = rt_prop13;
using rt_prop14_key = rt_prop14;
using rt_prop15_key = rt_prop15;
using rt_prop16_key = rt_prop16;
using rt_prop17_key = rt_prop17;
using rt_prop18_key = rt_prop18;
using rt_prop19_key = rt_prop19;
using rt_prop20_key = rt_prop20;
using rt_prop21_key = rt_prop21;
using rt_prop22_key = rt_prop22;
using rt_prop23_key = rt_prop23;
using rt_prop24_key = rt_prop24;
using rt_prop25_key = rt_prop25;
using rt_prop26_key = rt_prop26;
using rt_prop27_key = rt_prop27;
using rt_prop28_key = rt_prop28;
using rt_prop29_key = rt_prop29;
using rt_prop30_key = rt_prop30;
using rt_prop31_key = rt_prop31;
using rt_prop32_key = rt_prop32;
using rt_prop33_key = rt_prop33;
#undef rt_prop

template <> struct is_property_key<foo_key> : std::true_type {};
template <> struct is_property_key<foz_key> : std::true_type {};
Expand Down
Loading