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

Jump to bottom

[NFCI][SYCL] Simplify property_key creation #12831

Merged
merged 5 commits into from
Feb 28, 2024
Merged

[NFCI][SYCL] Simplify property_key creation #12831

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
dc06285
[NFCI][SYCL] Simplify property_key creation
aelovikov-intel Feb 26, 2024
5ab3e4f
Update the example
aelovikov-intel Feb 27, 2024
1aa7447
Make Kind protected
aelovikov-intel Feb 27, 2024
65c6e14
Update the last remaining test
aelovikov-intel Feb 27, 2024
1995fc5
Simplify mock_compile_time_properties.hpp by fakePropKind helper
aelovikov-intel Feb 27, 2024
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
16 changes: 3 additions & 13 deletions sycl/include/sycl/detail/kernel_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,9 @@ enum class register_alloc_mode_enum : uint32_t {
large = 2,
};

struct register_alloc_mode_key {
struct register_alloc_mode_key
: ext::oneapi::experimental::detail::compile_time_property_key<
ext::oneapi::experimental::detail::PropKind::RegisterAllocMode> {
template <register_alloc_mode_enum Mode>
using value_t = sycl::ext::oneapi::experimental::property_value<
register_alloc_mode_key,
Expand All @@ -36,19 +38,7 @@ inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode
} // namespace detail

namespace ext::oneapi::experimental {
template <>
struct is_property_key<sycl::detail::register_alloc_mode_key> : std::true_type {
};

namespace detail {
template <> struct PropertyToKind<sycl::detail::register_alloc_mode_key> {
static constexpr PropKind Kind = PropKind::RegisterAllocMode;
};

template <>
struct IsCompileTimeProperty<sycl::detail::register_alloc_mode_key>
: std::true_type {};

template <sycl::detail::register_alloc_mode_enum Mode>
struct PropertyMetaInfo<sycl::detail::register_alloc_mode_key::value_t<Mode>> {
static constexpr const char *name = "sycl-register-alloc-mode";
Expand Down
40 changes: 9 additions & 31 deletions sycl/include/sycl/ext/intel/esimd/memory_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,17 +59,23 @@ template <int K> inline constexpr alignment_key::value_t<K> alignment;
/// L2 cache hint property must be used for the old/experimental LSC L3 cache
/// hints.
/// L3 cache property is reserved for future devices.
struct cache_hint_L1_key {
struct cache_hint_L1_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::ESIMDL1CacheHint> {
template <cache_hint Hint>
using value_t = ext::oneapi::experimental::property_value<
cache_hint_L1_key, std::integral_constant<cache_hint, Hint>>;
};
struct cache_hint_L2_key {
struct cache_hint_L2_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::ESIMDL2CacheHint> {
template <cache_hint Hint>
using value_t = ext::oneapi::experimental::property_value<
cache_hint_L2_key, std::integral_constant<cache_hint, Hint>>;
};
struct cache_hint_L3_key {
struct cache_hint_L3_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::ESIMDL3CacheHint> {
template <cache_hint Hint>
using value_t = ext::oneapi::experimental::property_value<
cache_hint_L3_key, std::integral_constant<cache_hint, Hint>>;
Expand Down Expand Up @@ -151,45 +157,17 @@ struct property_value<__ESIMD_NS::cache_hint_L3_key,
static constexpr __ESIMD_NS::cache_hint hint = Hint;
};

template <>
struct is_property_key<sycl::ext::intel::esimd::cache_hint_L1_key>
: std::true_type {};
template <>
struct is_property_key<sycl::ext::intel::esimd::cache_hint_L2_key>
: std::true_type {};
template <>
struct is_property_key<sycl::ext::intel::esimd::cache_hint_L3_key>
: std::true_type {};

// Declare that esimd::properties is a property_list.
template <typename... PropertyValueTs>
struct is_property_list<__ESIMD_NS::properties<std::tuple<PropertyValueTs...>>>
: is_property_list<properties<std::tuple<PropertyValueTs...>>> {};

namespace detail {
template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L1_key> {
static constexpr PropKind Kind = PropKind::ESIMDL1CacheHint;
};
template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L2_key> {
static constexpr PropKind Kind = PropKind::ESIMDL2CacheHint;
};
template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L3_key> {
static constexpr PropKind Kind = PropKind::ESIMDL3CacheHint;
};

template <>
struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L1_key> : std::true_type {};
template <>
struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L2_key> : std::true_type {};
template <>
struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L3_key> : std::true_type {};

// We do not override the class ConflictingProperties for cache_hint properties
// because that mechanism would only allow to verify few obvious restrictions
// without the knowledge of the context in which the cache_hint properties are
// used (load, store, prefetch, atomic). Thus the function
// __ESIMD_DNS::check_cache_hint() is used to verify correctness of properties.

} // namespace detail
} // namespace ext::oneapi::experimental
} // namespace _V1
Expand Down
37 changes: 9 additions & 28 deletions sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,17 +72,23 @@ template <typename PropertyT, typename... Ts>
using property_value =
sycl::ext::oneapi::experimental::property_value<PropertyT, Ts...>;

struct read_hint_key {
struct read_hint_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::CacheControlReadHint> {
template <typename... Cs>
using value_t = property_value<read_hint_key, Cs...>;
};

struct read_assertion_key {
struct read_assertion_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::CacheControlReadAssertion> {
template <typename... Cs>
using value_t = property_value<read_assertion_key, Cs...>;
};

struct write_hint_key {
struct write_hint_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::CacheControlWrite> {
template <typename... Cs>
using value_t = property_value<write_hint_key, Cs...>;
};
Expand All @@ -104,21 +110,14 @@ namespace experimental {

template <typename T, typename PropertyListT> class annotated_ptr;

template <>
struct is_property_key<intel::experimental::read_hint_key> : std::true_type {};
template <typename T, typename PropertyListT>
struct is_property_key_of<intel::experimental::read_hint_key,
annotated_ptr<T, PropertyListT>> : std::true_type {};

template <>
struct is_property_key<intel::experimental::read_assertion_key>
: std::true_type {};
template <typename T, typename PropertyListT>
struct is_property_key_of<intel::experimental::read_assertion_key,
annotated_ptr<T, PropertyListT>> : std::true_type {};

template <>
struct is_property_key<intel::experimental::write_hint_key> : std::true_type {};
template <typename T, typename PropertyListT>
struct is_property_key_of<intel::experimental::write_hint_key,
annotated_ptr<T, PropertyListT>> : std::true_type {};
Expand Down Expand Up @@ -178,12 +177,6 @@ template <cache_mode M> static constexpr int checkWriteHint() {
return 0;
}

template <> struct PropertyToKind<intel::experimental::read_hint_key> {
static constexpr PropKind Kind = PropKind::CacheControlReadHint;
};
template <>
struct IsCompileTimeProperty<intel::experimental::read_hint_key>
: std::true_type {};
template <typename... Cs>
struct PropertyMetaInfo<intel::experimental::read_hint_key::value_t<Cs...>> {
static constexpr const char *name = "sycl-cache-read-hint";
Expand All @@ -196,12 +189,6 @@ struct PropertyMetaInfo<intel::experimental::read_hint_key::value_t<Cs...>> {
((Cs::encoding) | ...));
};

template <> struct PropertyToKind<intel::experimental::read_assertion_key> {
static constexpr PropKind Kind = PropKind::CacheControlReadAssertion;
};
template <>
struct IsCompileTimeProperty<intel::experimental::read_assertion_key>
: std::true_type {};
template <typename... Cs>
struct PropertyMetaInfo<
intel::experimental::read_assertion_key::value_t<Cs...>> {
Expand All @@ -215,12 +202,6 @@ struct PropertyMetaInfo<
((Cs::encoding) | ...));
};

template <> struct PropertyToKind<intel::experimental::write_hint_key> {
static constexpr PropKind Kind = PropKind::CacheControlWrite;
};
template <>
struct IsCompileTimeProperty<intel::experimental::write_hint_key>
: std::true_type {};
template <typename... Cs>
struct PropertyMetaInfo<intel::experimental::write_hint_key::value_t<Cs...>> {
static constexpr const char *name = "sycl-cache-write-hint";
Expand Down
15 changes: 3 additions & 12 deletions sycl/include/sycl/ext/intel/experimental/fp_control_kernel_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,7 +85,9 @@ constexpr fp_mode setDefaultValuesIfNeeded(fp_mode mode) {
}
} // namespace detail

struct fp_control_key {
struct fp_control_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::FloatingPointControls> {
template <fp_mode option>
using value_t = ext::oneapi::experimental::property_value<
fp_control_key, std::integral_constant<fp_mode, option>>;
Expand All @@ -97,24 +99,13 @@ inline constexpr fp_control_key::value_t<option> fp_control;
} // namespace ext::intel::experimental

namespace ext::oneapi::experimental {
template <>
struct is_property_key<intel::experimental::fp_control_key> : std::true_type {};

template <typename T, typename PropertyListT>
struct is_property_key_of<
intel::experimental::fp_control_key,
intel::experimental::kernel_attribute<T, PropertyListT>> : std::true_type {
};

namespace detail {
template <> struct PropertyToKind<intel::experimental::fp_control_key> {
static constexpr PropKind Kind = FloatingPointControls;
};

template <>
struct IsCompileTimeProperty<intel::experimental::fp_control_key>
: std::true_type {};

template <intel::experimental::fp_mode FPMode>
struct PropertyMetaInfo<intel::experimental::fp_control_key::value_t<FPMode>> {
static_assert(intel::experimental::detail::checkMutuallyExclusive(FPMode),
Expand Down
91 changes: 24 additions & 67 deletions sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,54 +31,68 @@ using property_value =
//===----------------------------------------------------------------------===//
// FPGA properties of annotated_arg/annotated_ptr
//===----------------------------------------------------------------------===//
struct register_map_key {
struct register_map_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::RegisterMap> {
using value_t = property_value<register_map_key>;
};

struct conduit_key {
struct conduit_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::Conduit> {
using value_t = property_value<conduit_key>;
};

struct stable_key {
struct stable_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::Stable> {
using value_t = property_value<stable_key>;
};

struct buffer_location_key {
struct buffer_location_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::BufferLocation> {
Comment on lines +50 to +52
Copy link
Contributor

Choose a reason for hiding this comment

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

Would that be always ABI compatible taking into account that some properties are used in sycl/source? E.g.: https://github.com/intel/llvm/blob/sycl/sycl/source/detail/usm/usm_impl.cpp#L98-L99

Copy link
Contributor

Choose a reason for hiding this comment

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

The ones in that example are from SYCL 2020 properties, while these properties are all from the sycl_ext_oneapi_properties extension. Since that property-list is templated I don't believe it crosses the library boundaries, so it would only be if we pass the properties around by themselves.

template <int K>
using value_t =
property_value<buffer_location_key, std::integral_constant<int, K>>;
};

struct awidth_key {
struct awidth_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::AddrWidth> {
template <int K>
using value_t = property_value<awidth_key, std::integral_constant<int, K>>;
};

struct dwidth_key {
struct dwidth_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::DataWidth> {
template <int K>
using value_t = property_value<dwidth_key, std::integral_constant<int, K>>;
};

struct latency_key {
struct latency_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::Latency> {
template <int K>
using value_t = property_value<latency_key, std::integral_constant<int, K>>;
};

enum class read_write_mode_enum : std::uint16_t { read, write, read_write };

struct read_write_mode_key {
struct read_write_mode_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::RWMode> {
template <read_write_mode_enum Mode>
using value_t =
property_value<read_write_mode_key,
std::integral_constant<read_write_mode_enum, Mode>>;
};

struct maxburst_key {
struct maxburst_key : oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::MaxBurst> {
template <int K>
using value_t = property_value<maxburst_key, std::integral_constant<int, K>>;
};

struct wait_request_key {
struct wait_request_key
: oneapi::experimental::detail::compile_time_property_key<
oneapi::experimental::detail::PropKind::WaitRequest> {
template <int K>
using value_t =
property_value<wait_request_key, std::integral_constant<int, K>>;
Expand Down Expand Up @@ -131,18 +145,6 @@ using maxburst_key = intel::experimental::maxburst_key;
using wait_request_key = intel::experimental::wait_request_key;
using read_write_mode_enum = intel::experimental::read_write_mode_enum;

template <> struct is_property_key<register_map_key> : std::true_type {};
template <> struct is_property_key<conduit_key> : std::true_type {};
template <> struct is_property_key<stable_key> : std::true_type {};

template <> struct is_property_key<buffer_location_key> : std::true_type {};
template <> struct is_property_key<awidth_key> : std::true_type {};
template <> struct is_property_key<dwidth_key> : std::true_type {};
template <> struct is_property_key<latency_key> : std::true_type {};
template <> struct is_property_key<read_write_mode_key> : std::true_type {};
template <> struct is_property_key<maxburst_key> : std::true_type {};
template <> struct is_property_key<wait_request_key> : std::true_type {};

template <typename T, typename PropertyListT>
struct is_property_key_of<register_map_key, annotated_arg<T, PropertyListT>>
: std::true_type {};
Expand Down Expand Up @@ -224,51 +226,6 @@ struct is_property_key_of<wait_request_key, annotated_ptr<T, PropertyListT>>
: std::true_type {};

namespace detail {
template <> struct PropertyToKind<register_map_key> {
static constexpr PropKind Kind = PropKind::RegisterMap;
};
template <> struct PropertyToKind<conduit_key> {
static constexpr PropKind Kind = PropKind::Conduit;
};
template <> struct PropertyToKind<stable_key> {
static constexpr PropKind Kind = PropKind::Stable;
};
template <> struct PropertyToKind<buffer_location_key> {
static constexpr PropKind Kind = PropKind::BufferLocation;
};
template <> struct PropertyToKind<awidth_key> {
static constexpr PropKind Kind = PropKind::AddrWidth;
};
template <> struct PropertyToKind<dwidth_key> {
static constexpr PropKind Kind = PropKind::DataWidth;
};
template <> struct PropertyToKind<latency_key> {
static constexpr PropKind Kind = PropKind::Latency;
};
template <> struct PropertyToKind<read_write_mode_key> {
static constexpr PropKind Kind = PropKind::RWMode;
};
template <> struct PropertyToKind<maxburst_key> {
static constexpr PropKind Kind = PropKind::MaxBurst;
};
template <> struct PropertyToKind<wait_request_key> {
static constexpr PropKind Kind = PropKind::WaitRequest;
};

template <> struct IsCompileTimeProperty<register_map_key> : std::true_type {};
template <> struct IsCompileTimeProperty<conduit_key> : std::true_type {};
template <> struct IsCompileTimeProperty<stable_key> : std::true_type {};

template <>
struct IsCompileTimeProperty<buffer_location_key> : std::true_type {};
template <> struct IsCompileTimeProperty<awidth_key> : std::true_type {};
template <> struct IsCompileTimeProperty<dwidth_key> : std::true_type {};
template <>
struct IsCompileTimeProperty<read_write_mode_key> : std::true_type {};
template <> struct IsCompileTimeProperty<latency_key> : std::true_type {};
template <> struct IsCompileTimeProperty<maxburst_key> : std::true_type {};
template <> struct IsCompileTimeProperty<wait_request_key> : std::true_type {};

template <> struct PropertyMetaInfo<register_map_key::value_t> {
static constexpr const char *name = "sycl-register-map";
static constexpr std::nullptr_t value = nullptr;
Expand Down
Loading
Loading