DPC++ Runtime
Runtime libraries for oneAPI DPC++
memory_properties.hpp
Go to the documentation of this file.
1 //==-- memory_properties.hpp - ESIMD memory properties ---------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
15 
16 #define SYCL_EXT_INTEL_ESIMD_MEMORY_PROPERTIES 1
17 
18 namespace sycl {
19 inline namespace _V1 {
20 namespace ext::intel::esimd {
21 
22 template <typename PropertiesT>
24  : public sycl::ext::oneapi::experimental::properties<PropertiesT> {
25 public:
26  template <typename... PropertyValueTs>
27  constexpr properties(PropertyValueTs... props)
28  : sycl::ext::oneapi::experimental::properties<PropertiesT>(props...) {}
29 };
30 
31 #ifdef __cpp_deduction_guides
32 // Deduction guides
33 template <typename... PropertyValueTs>
34 properties(PropertyValueTs... props)
36  PropertyValueTs...>::type>;
37 #endif
38 
41 //
45 //
50 
52 
53 template <int K> inline constexpr alignment_key::value_t<K> alignment;
54 
63  template <cache_hint Hint>
65  cache_hint_L1_key, std::integral_constant<cache_hint, Hint>>;
66 };
68  template <cache_hint Hint>
70  cache_hint_L2_key, std::integral_constant<cache_hint, Hint>>;
71 };
73  template <cache_hint Hint>
75  cache_hint_L3_key, std::integral_constant<cache_hint, Hint>>;
76 };
77 
78 template <cache_hint Hint>
80 template <cache_hint Hint>
82 template <cache_hint Hint>
84 
85 #if 0
86 // TODO: Introduce the 2-parameter cache_hint property after looking
87 // for a better name for it. It cannot be 'esimd::cache_hint' because that
88 // may conflict with the enum name 'experimental::esimd::cache_hint' when
89 // both namespaces (esimd and experimental::esimd) are imported with 'using'
90 // statement.
91 // Naming alternatives: 'esimd::esimd_cache_hint, esimd::cache_hint_L,
92 // esimd::cache_hint_property'.
93 template <cache_level Level, cache_hint Hint>
94 inline constexpr std::conditional_t<
96  std::conditional_t<Level == cache_level::L2,
99  cache_hint; // Get a non-conflicting name
100 #endif
101 
105 
106 namespace detail {
110 template <typename PropertiesT, typename KeyT, typename KeyValueT,
111  typename = std::enable_if_t<
112  ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
113 constexpr auto getPropertyValue(KeyValueT DefaultValue) {
114  if constexpr (!PropertiesT::template has_property<KeyT>()) {
115  return DefaultValue;
116  } else if constexpr (std::is_same_v<KeyT, cache_hint_L1_key> ||
117  std::is_same_v<KeyT, cache_hint_L2_key> ||
118  std::is_same_v<KeyT, cache_hint_L3_key>) {
119  constexpr auto ValueT = PropertiesT::template get_property<KeyT>();
120  return ValueT.hint;
121  } else {
122  constexpr auto ValueT = PropertiesT::template get_property<KeyT>();
123  return ValueT.value;
124  }
125 }
126 } // namespace detail
127 
128 } // namespace ext::intel::esimd
129 
130 namespace ext::oneapi::experimental {
131 
132 template <__ESIMD_NS::cache_hint Hint>
133 struct property_value<__ESIMD_NS::cache_hint_L1_key,
134  std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
135  using key_t = __ESIMD_NS::cache_hint_L1_key;
136  static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L1;
137  static constexpr __ESIMD_NS::cache_hint hint = Hint;
138 };
139 template <__ESIMD_NS::cache_hint Hint>
140 struct property_value<__ESIMD_NS::cache_hint_L2_key,
141  std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
142  using key_t = __ESIMD_NS::cache_hint_L2_key;
143  static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L2;
144  static constexpr __ESIMD_NS::cache_hint hint = Hint;
145 };
146 template <__ESIMD_NS::cache_hint Hint>
147 struct property_value<__ESIMD_NS::cache_hint_L3_key,
148  std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
149  using key_t = __ESIMD_NS::cache_hint_L3_key;
150  static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L3;
151  static constexpr __ESIMD_NS::cache_hint hint = Hint;
152 };
153 
154 template <>
156  : std::true_type {};
157 template <>
159  : std::true_type {};
160 template <>
162  : std::true_type {};
163 
164 // Declare that esimd::properties is a property_list.
165 template <typename... PropertyValueTs>
166 struct is_property_list<__ESIMD_NS::properties<std::tuple<PropertyValueTs...>>>
167  : is_property_list<properties<std::tuple<PropertyValueTs...>>> {};
168 
169 namespace detail {
171  static constexpr PropKind Kind = PropKind::ESIMDL1CacheHint;
172 };
174  static constexpr PropKind Kind = PropKind::ESIMDL2CacheHint;
175 };
177  static constexpr PropKind Kind = PropKind::ESIMDL3CacheHint;
178 };
179 
180 template <>
181 struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L1_key> : std::true_type {};
182 template <>
183 struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L2_key> : std::true_type {};
184 template <>
185 struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L3_key> : std::true_type {};
186 
187 // We do not override the class ConflictingProperties for cache_hint properties
188 // because that mechanism would only allow to verify few obvious restrictions
189 // without the knowledge of the context in which the cache_hint properties are
190 // used (load, store, prefetch, atomic). Thus the function
191 // __ESIMD_DNS::check_cache_hint() is used to verify correctness of properties.
192 
193 } // namespace detail
194 } // namespace ext::oneapi::experimental
195 } // namespace _V1
196 } // namespace sycl
constexpr properties(PropertyValueTs... props)
constexpr auto getPropertyValue(KeyValueT DefaultValue)
Helper-function that returns the value of the compile time property KeyT if PropertiesT includes it.
constexpr alignment_key::value_t< K > alignment
cache_level
L1, L2 or L3 cache hint levels. L3 is reserved for future use.
Definition: common.hpp:444
cache_hint
L1, L2 or L3 cache hints.
Definition: common.hpp:348
constexpr cache_hint_L3_key::value_t< Hint > cache_hint_L3
constexpr cache_hint_L2_key::value_t< Hint > cache_hint_L2
constexpr cache_hint_L1_key::value_t< Hint > cache_hint_L1
sycl::ext::oneapi::experimental::alignment_key alignment_key
The 'alignment' property is used to specify the alignment of memory accessed in ESIMD memory operatio...
Definition: access.hpp:18
The 'cache_hint_L1', 'cache_hint_L2' and 'cache_hint_L3' properties are used to specify L1,...