DPC++ Runtime
Runtime libraries for oneAPI DPC++
properties.hpp
Go to the documentation of this file.
1 //==------- properties.hpp - SYCL properties associated with kernels -------==//
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 
11 #include <sycl/aspects.hpp> // for aspect
12 #include <sycl/ext/oneapi/properties/property.hpp> // for PropKind
13 #include <sycl/ext/oneapi/properties/property_utils.hpp> // for SizeListToStr
14 #include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
15 
16 #include <array> // for array
17 #include <stddef.h> // for size_t
18 #include <stdint.h> // for uint32_t
19 #include <type_traits> // for true_type
20 #include <utility> // for declval
21 
22 namespace sycl {
23 inline namespace _V1 {
24 namespace ext::oneapi::experimental {
25 namespace detail {
26 // Trait for checking that all size_t values are non-zero.
27 template <size_t... Xs> struct AllNonZero {
28  static constexpr bool value = true;
29 };
30 template <size_t X, size_t... Xs> struct AllNonZero<X, Xs...> {
31  static constexpr bool value = X > 0 && AllNonZero<Xs...>::value;
32 };
33 } // namespace detail
34 
35 struct properties_tag {};
36 
38  template <size_t... Dims>
40  std::integral_constant<size_t, Dims>...>;
41 };
42 
44  template <size_t... Dims>
46  std::integral_constant<size_t, Dims>...>;
47 };
48 
50  template <uint32_t Size>
52  std::integral_constant<uint32_t, Size>>;
53 };
54 
56  template <aspect... Aspects>
58  std::integral_constant<aspect, Aspects>...>;
59 };
60 
61 template <size_t Dim0, size_t... Dims>
62 struct property_value<work_group_size_key, std::integral_constant<size_t, Dim0>,
63  std::integral_constant<size_t, Dims>...> {
64  static_assert(
65  sizeof...(Dims) + 1 <= 3,
66  "work_group_size property currently only supports up to three values.");
68  "work_group_size property must only contain non-zero values.");
69 
71 
72  constexpr size_t operator[](int Dim) const {
73  return std::array<size_t, sizeof...(Dims) + 1>{Dim0, Dims...}[Dim];
74  }
75 };
76 
77 template <size_t Dim0, size_t... Dims>
79  std::integral_constant<size_t, Dim0>,
80  std::integral_constant<size_t, Dims>...> {
81  static_assert(sizeof...(Dims) + 1 <= 3,
82  "work_group_size_hint property currently "
83  "only supports up to three values.");
84  static_assert(
86  "work_group_size_hint property must only contain non-zero values.");
87 
89 
90  constexpr size_t operator[](int Dim) const {
91  return std::array<size_t, sizeof...(Dims) + 1>{Dim0, Dims...}[Dim];
92  }
93 };
94 
95 template <uint32_t Size>
97  std::integral_constant<uint32_t, Size>> {
98  static_assert(Size != 0,
99  "sub_group_size_key property must contain a non-zero value.");
100 
102  using value_t = std::integral_constant<uint32_t, Size>;
103  static constexpr uint32_t value = Size;
104 };
105 
106 template <aspect... Aspects>
108  std::integral_constant<aspect, Aspects>...> {
110  static constexpr std::array<aspect, sizeof...(Aspects)> value{Aspects...};
111 };
112 
113 template <size_t Dim0, size_t... Dims>
114 inline constexpr work_group_size_key::value_t<Dim0, Dims...> work_group_size;
115 
116 template <size_t Dim0, size_t... Dims>
117 inline constexpr work_group_size_hint_key::value_t<Dim0, Dims...>
119 
120 template <uint32_t Size>
122 
123 template <aspect... Aspects>
124 inline constexpr device_has_key::value_t<Aspects...> device_has;
125 
126 template <> struct is_property_key<work_group_size_key> : std::true_type {};
127 template <>
128 struct is_property_key<work_group_size_hint_key> : std::true_type {};
129 template <> struct is_property_key<sub_group_size_key> : std::true_type {};
130 template <> struct is_property_key<device_has_key> : std::true_type {};
131 
132 namespace detail {
133 template <> struct PropertyToKind<work_group_size_key> {
134  static constexpr PropKind Kind = PropKind::WorkGroupSize;
135 };
137  static constexpr PropKind Kind = PropKind::WorkGroupSizeHint;
138 };
139 template <> struct PropertyToKind<sub_group_size_key> {
140  static constexpr PropKind Kind = PropKind::SubGroupSize;
141 };
142 template <> struct PropertyToKind<device_has_key> {
143  static constexpr PropKind Kind = PropKind::DeviceHas;
144 };
145 
146 template <>
147 struct IsCompileTimeProperty<work_group_size_key> : std::true_type {};
148 template <>
150 template <>
151 struct IsCompileTimeProperty<sub_group_size_key> : std::true_type {};
152 template <> struct IsCompileTimeProperty<device_has_key> : std::true_type {};
153 
154 template <size_t Dim0, size_t... Dims>
155 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
156  static constexpr const char *name = "sycl-work-group-size";
157  static constexpr const char *value = SizeListToStr<Dim0, Dims...>::value;
158 };
159 template <size_t Dim0, size_t... Dims>
160 struct PropertyMetaInfo<work_group_size_hint_key::value_t<Dim0, Dims...>> {
161  static constexpr const char *name = "sycl-work-group-size-hint";
162  static constexpr const char *value = SizeListToStr<Dim0, Dims...>::value;
163 };
164 template <uint32_t Size>
165 struct PropertyMetaInfo<sub_group_size_key::value_t<Size>> {
166  static constexpr const char *name = "sycl-sub-group-size";
167  static constexpr uint32_t value = Size;
168 };
169 template <aspect... Aspects>
170 struct PropertyMetaInfo<device_has_key::value_t<Aspects...>> {
171  static constexpr const char *name = "sycl-device-has";
172  static constexpr const char *value =
173  SizeListToStr<static_cast<size_t>(Aspects)...>::value;
174 };
175 
176 template <typename T, typename = void>
177 struct HasKernelPropertiesGetMethod : std::false_type {};
178 
179 template <typename T>
181  std::void_t<decltype(std::declval<T>().get(
182  std::declval<properties_tag>()))>>
183  : std::true_type {
184  using properties_t =
185  decltype(std::declval<T>().get(std::declval<properties_tag>()));
186 };
187 
188 } // namespace detail
189 } // namespace ext::oneapi::experimental
190 } // namespace _V1
191 } // namespace sycl
192 
193 #ifdef __SYCL_DEVICE_ONLY__
194 #define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) \
195  [[__sycl_detail__::add_ir_attributes_function( \
196  {"sycl-device-has"}, \
197  sycl::ext::oneapi::experimental::detail::PropertyMetaInfo< \
198  std::remove_cv_t<std::remove_reference_t<decltype(PROP)>>>::name, \
199  sycl::ext::oneapi::experimental::detail::PropertyMetaInfo< \
200  std::remove_cv_t<std::remove_reference_t<decltype(PROP)>>>::value)]]
201 #else
202 #define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP)
203 #endif
constexpr sub_group_size_key::value_t< Size > sub_group_size
Definition: properties.hpp:121
constexpr work_group_size_key::value_t< Dim0, Dims... > work_group_size
Definition: properties.hpp:114
constexpr device_has_key::value_t< Aspects... > device_has
Definition: properties.hpp:124
constexpr work_group_size_hint_key::value_t< Dim0, Dims... > work_group_size_hint
Definition: properties.hpp:118
Definition: access.hpp:18