DPC++ Runtime
Runtime libraries for oneAPI DPC++
uniform.hpp
Go to the documentation of this file.
1 //==------ uniform.hpp - SYCL uniform extension --------*- C++ -*-----------==//
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 // Implemenation of the sycl_ext_oneapi_uniform extension.
9 // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_uniform.asciidoc
10 // ===--------------------------------------------------------------------=== //
11 
12 #pragma once
13 
14 // SYCL extension macro definition as required by the SYCL specification.
15 // 1 - Initial extension version. Base features are supported.
16 #define SYCL_EXT_ONEAPI_UNIFORM 1
17 
18 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_INLINE_NAMESPACE
19 
20 #include <type_traits>
21 
22 // Forward declarations of types not allowed to be wrapped in uniform:
23 namespace sycl {
25 namespace ext {
26 namespace oneapi {
27 
28 struct sub_group;
29 
30 } // namespace oneapi
31 } // namespace ext
32 
33 template <int, bool> class item;
34 template <int> class id;
35 template <int> class nd_item;
36 template <int> class h_item;
37 template <int> class group;
38 template <int> class nd_range;
40 
41 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
42 } // namespace sycl
43 
44 namespace sycl {
46 namespace ext {
47 namespace oneapi {
48 namespace experimental {
49 namespace detail {
50 
51 template <class T, template <int> class Tmpl>
52 struct is_instance_of_tmpl_int : std::false_type {};
53 template <int N, template <int> class T, template <int> class Tmpl>
54 struct is_instance_of_tmpl_int<T<N>, Tmpl>
55  : std::conditional<std::is_same_v<T<N>, Tmpl<N>>, std::true_type,
56  std::false_type> {};
57 template <class T, template <int> class Tmpl>
58 static inline constexpr bool is_instance_of_tmpl_int_v =
60 
61 template <class T, template <int, bool> class Tmpl>
62 struct is_instance_of_tmpl_int_bool : std::false_type {};
63 template <int N, bool X, template <int, bool> class T,
64  template <int, bool> class Tmpl>
65 struct is_instance_of_tmpl_int_bool<T<N, X>, Tmpl>
66  : std::conditional<std::is_same_v<T<N, X>, Tmpl<N, X>>, std::true_type,
67  std::false_type> {};
68 template <class T, template <int, bool> class Tmpl>
69 static inline constexpr bool is_instance_of_tmpl_int_bool_v =
71 } // namespace detail
72 
73 template <class T> class uniform {
74  template <class U> static constexpr bool can_be_uniform() {
75  return !detail::is_instance_of_tmpl_int_v<U, sycl::id> &&
76  !detail::is_instance_of_tmpl_int_bool_v<U, sycl::item> &&
77  !detail::is_instance_of_tmpl_int_v<U, sycl::nd_item> &&
78  !detail::is_instance_of_tmpl_int_v<U, sycl::h_item> &&
79  !detail::is_instance_of_tmpl_int_v<U, sycl::group> &&
80  !detail::is_instance_of_tmpl_int_v<U, sycl::nd_range> &&
81  !std::is_same_v<U, sycl::sub_group>;
82  }
83  static_assert(can_be_uniform<T>() && "type not allowed to be `uniform`");
84 
85 public:
86  explicit uniform(T x) noexcept : Val(x) {}
87 
88  // TODO provide a ways to reflect this conversion from uniform to T in the IR
89  // so that the compiler can take advantage of uniformness. Could be marked
90  // with some intrinsic call like `__builtin_uniform_unwrap(Val);`
91 
93  operator const T() const { return Val; }
94 
95  uniform &operator=(const uniform &) = delete;
96 
97  /* Other explicitly deleted operators improve error messages
98  if a user incorrectly attempts to modify a uniform */
99  uniform &operator+=(const T &) = delete;
100  uniform &operator-=(const T &) = delete;
101  uniform &operator*=(const T &) = delete;
102  uniform &operator/=(const T &) = delete;
103  uniform &operator%=(const T &) = delete;
104  uniform &operator&=(const T &) = delete;
105  uniform &operator|=(const T &) = delete;
106  uniform &operator^=(const T &) = delete;
107  uniform &operator<<=(const T &) = delete;
108  uniform &operator>>=(const T &) = delete;
109  uniform &operator++() = delete;
110  uniform &operator++(int) = delete;
111  uniform &operator--() = delete;
112  uniform &operator--(int) = delete;
113 
114 private:
115  T Val;
116 };
117 
118 } // namespace experimental
119 } // namespace oneapi
120 } // namespace ext
121 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
122 } // namespace sycl
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:13
sycl::_V1::ext::oneapi::experimental::uniform::uniform
uniform(T x) noexcept
Definition: uniform.hpp:86
defines_elementary.hpp
sycl::_V1::ext::oneapi::experimental::detail::is_instance_of_tmpl_int_v
static constexpr bool is_instance_of_tmpl_int_v
Definition: uniform.hpp:58
sycl::_V1::ext::oneapi::experimental::uniform
Definition: uniform.hpp:73
sycl::_V1::ext::oneapi::experimental::detail::is_instance_of_tmpl_int_bool_v
static constexpr bool is_instance_of_tmpl_int_bool_v
Definition: uniform.hpp:69
sycl::_V1::nd_range
Defines the iteration domain of both the work-groups and the overall dispatch.
Definition: uniform.hpp:38
sycl::_V1::ext::oneapi::experimental::detail::is_instance_of_tmpl_int_bool
Definition: uniform.hpp:62
sycl::_V1::ext::oneapi::experimental::detail::is_instance_of_tmpl_int
Definition: uniform.hpp:52
sycl::_V1::ext::oneapi::sub_group
Definition: sub_group.hpp:108