DPC++ Runtime
Runtime libraries for oneAPI DPC++
boolean.hpp
Go to the documentation of this file.
1 //==----------- boolean.hpp - SYCL boolean type ----------------------------==//
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 
12 #include <CL/sycl/types.hpp>
13 
14 #include <initializer_list>
15 #include <type_traits>
16 
18 namespace sycl {
19 namespace detail {
20 
21 template <int Num> struct Assigner {
22  template <typename R, typename T> static void assign(R &r, const T x) {
24  r.template swizzle<Num>() = x.value[Num];
25  }
26 
27  template <typename R, typename T, typename ET>
28  static void init(R &r, const T x) {
29  Assigner<Num - 1>::template init<R, T, ET>(r, x);
30  ET v = x.template swizzle<Num>();
31  r.value[Num] = msbIsSet(v) * (-1);
32  }
33 };
34 
35 template <> struct Assigner<0> {
36  template <typename R, typename T> static void assign(R &r, const T x) {
37  r.template swizzle<0>() = x.value[0];
38  }
39  template <typename R, typename T, typename ET>
40  static void init(R &r, const T x) {
41  ET v = x.template swizzle<0>();
42  r.value[0] = msbIsSet(v) * (-1);
43  }
44 };
45 
46 template <int N> struct Boolean {
47  static_assert(((N == 2) || (N == 3) || (N == 4) || (N == 8) || (N == 16)),
48  "Invalid size");
49 
50  using element_type = int8_t;
51 
52 #ifdef __SYCL_DEVICE_ONLY__
53  using DataType = element_type __attribute__((ext_vector_type(N)));
54  using vector_t = DataType;
55 #else
56  using DataType = element_type[N];
57 #endif
58 
59  Boolean() : value{0} {}
60 
61  Boolean(std::initializer_list<element_type> l) {
62  for (size_t I = 0; I < N; ++I) {
63  value[I] = *(l.begin() + I) ? -1 : 0;
64  }
65  }
66 
67  Boolean(const Boolean &rhs) {
68  for (size_t I = 0; I < N; ++I) {
69  value[I] = rhs.value[I];
70  }
71  }
72 
73  template <typename T> Boolean(const T rhs) {
74  static_assert(is_vgeninteger<T>::value, "Invalid constructor");
75  Assigner<N - 1>::template init<Boolean<N>, T, typename T::element_type>(
76  *this, rhs);
77  }
78 
79 #ifdef __SYCL_DEVICE_ONLY__
80  // TODO change this to the vectors assignment when the assignment will be
81  // fixed on Intel GPU NEO OpenCL runtime
82  Boolean(const vector_t rhs) {
83  for (size_t I = 0; I < N; ++I) {
84  value[I] = rhs[I];
85  }
86  }
87 
88  operator vector_t() const { return value; }
89 #endif
90 
91  template <typename T> operator T() const {
92  static_assert(is_vgeninteger<T>::value, "Invalid conversion");
93  T r;
94  Assigner<N - 1>::assign(r, *this);
95  return r;
96  }
97 
98 private:
99  template <int Num> friend struct Assigner;
101 };
102 
103 template <> struct Boolean<1> {
104  Boolean() = default;
105 
106  // Build from a signed interger type
107  template <typename T> Boolean(T val) : value(val) {
108  static_assert(is_sgeninteger<T>::value, "Invalid constructor");
109  }
110 
111  // Cast to a signed interger type
112  template <typename T> operator T() const {
113  static_assert(is_sgeninteger<T>::value, "Invalid conversion");
114  return value;
115  }
116 
117 #ifdef __SYCL_DEVICE_ONLY__
118  // Build from a boolean type
119  Boolean(bool f) : value(f) {}
120  // Cast to a boolean type
121  operator bool() const { return value; }
122 #endif
123 
124 private:
125  alignas(1) bool value = false;
126 };
127 
128 } // namespace detail
129 } // namespace sycl
130 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::Boolean::DataType
element_type[N] DataType
Definition: boolean.hpp:56
cl::sycl::detail::Boolean::element_type
int8_t element_type
Definition: boolean.hpp:50
cl::sycl::detail::msbIsSet
constexpr bool msbIsSet(const T x)
Definition: generic_type_traits.hpp:500
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
cl::sycl::detail::Assigner::init
static void init(R &r, const T x)
Definition: boolean.hpp:28
sycl
Definition: invoke_simd.hpp:68
cl::sycl::detail::Assigner
Definition: boolean.hpp:21
cl::sycl::detail::Boolean::Boolean
Boolean(const Boolean &rhs)
Definition: boolean.hpp:67
cl::sycl::detail::Boolean< 1 >::Boolean
Boolean(T val)
Definition: boolean.hpp:107
generic_type_traits.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::image_channel_order::r
@ r
cl::sycl::detail::Boolean::Boolean
Boolean(const T rhs)
Definition: boolean.hpp:73
cl::sycl::detail::Boolean::Boolean
Boolean(std::initializer_list< element_type > l)
Definition: boolean.hpp:61
cl::sycl::detail::Assigner< 0 >::init
static void init(R &r, const T x)
Definition: boolean.hpp:40
sycl::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:293
cl::sycl::detail::Assigner::assign
static void assign(R &r, const T x)
Definition: boolean.hpp:22
cl::sycl::detail::vector_alignment
Definition: type_traits.hpp:86
cl::sycl::detail::Boolean::Boolean
Boolean()
Definition: boolean.hpp:59
types.hpp
cl::sycl::detail::Assigner< 0 >::assign
static void assign(R &r, const T x)
Definition: boolean.hpp:36
cl::sycl::detail::Boolean
Definition: boolean.hpp:46
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12