DPC++ Runtime
Runtime libraries for oneAPI DPC++
simd_mask_impl.hpp
Go to the documentation of this file.
1 //==------------ - simd_mask_impl.hpp - DPC++ Explicit SIMD API ----------==//
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 // Implementation detail of Explicit SIMD mask class.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
15 
17 namespace __ESIMD_DNS {
18 
21 
54 template <typename T, int N>
56  : public simd_obj_impl<
57  T, N, simd_mask_impl<T, N>,
58  std::enable_if_t<std::is_same_v<simd_mask_elem_type, T>>> {
62 
63 public:
65  using raw_element_type = T;
67  using element_type = T;
69  using raw_vector_type = typename base_type::raw_vector_type;
70  static_assert(std::is_same_v<raw_vector_type, simd_mask_storage_t<N>> &&
71  "mask impl type mismatch");
72 
74  simd_mask_impl() = default;
75 
77  simd_mask_impl(const simd_mask_impl &other) : base_type(other) {}
78 
81  template <class T1, class = std::enable_if_t<std::is_integral_v<T1>>>
82  simd_mask_impl(T1 Val) : base_type((T)Val) {}
83 
85  // TODO this should be made inaccessible from user code.
86  simd_mask_impl(const raw_vector_type &Val) : base_type(Val) {}
87 
89  template <int N1, class = std::enable_if_t<N1 == N>>
90  simd_mask_impl(const raw_element_type (&&Arr)[N1]) {
91  base_type::template init_from_array<false>(std::move(Arr));
92  }
93 
95  simd_mask_impl(const simd<T, N> &Val) : base_type(Val.data()) {}
96 
97 private:
99  static inline constexpr bool mask_size_ok_for_mem_io() {
100  constexpr unsigned Sz = sizeof(element_type) * N;
101  return (Sz >= OperandSize::OWORD) && (Sz % OperandSize::OWORD == 0) &&
102  isPowerOf2(Sz / OperandSize::OWORD) &&
103  (Sz <= 8 * OperandSize::OWORD);
104  }
106 
107 public:
108  // TODO add accessor-based mask memory operations.
109 
111  // Implementation note: use SFINAE to avoid overload ambiguity:
112  // 1) with 'simd_mask(element_type v)' in 'simd_mask<N> m(0)'
113  // 2) with 'simd_mask(const T1(&&arr)[N])' in simd_mask<N>
114  // m((element_type*)p)'
115  template <typename T1,
116  typename = std::enable_if_t<mask_size_ok_for_mem_io() &&
117  std::is_same_v<T1, element_type>>>
118  explicit simd_mask_impl(const T1 *ptr) {
119  base_type::copy_from(ptr);
120  }
121 
124  base_type::set(val);
125  return *this;
126  }
127 
130  template <class T1 = simd_mask_impl,
131  class = std::enable_if_t<T1::length == 1>>
132  operator bool() const {
133  return base_type::data()[0] != 0;
134  }
135 };
136 
137 #undef __ESIMD_MASK_DEPRECATION_MSG
138 
140 
141 } // namespace __ESIMD_DNS
142 } // __SYCL_INLINE_NAMESPACE(cl)
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(T1 Val)
Broadcast constructor with conversion.
Definition: simd_mask_impl.hpp:82
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(const raw_element_type(&&Arr)[N1])
Construct from an array. To allow e.g. simd_mask<N> m({1,0,0,1,...}).
Definition: simd_mask_impl.hpp:90
simd_obj_impl
cl::__ESIMD_DNS::simd_mask_impl::element_type
T element_type
Element type, same as raw.
Definition: simd_mask_impl.hpp:67
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(const raw_vector_type &Val)
Implicit conversion constructor from a raw vector object.
Definition: simd_mask_impl.hpp:86
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(const simd< T, N > &Val)
Implicit conversion from simd.
Definition: simd_mask_impl.hpp:95
cl::__ESIMD_DNS::simd_mask_impl::operator=
simd_mask_impl & operator=(element_type val) noexcept
Broadcast assignment operator to support simd_mask_impl<N> n = a > b;.
Definition: simd_mask_impl.hpp:123
simd_obj_impl.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::__ESIMD_DNS::simd_mask_impl
This class is a simd_obj_impl specialization representing a simd mask, which is basically a simd_obj_...
Definition: simd_mask_impl.hpp:55
types.hpp
cl::__ESIMD_DNS::simd_mask_impl::raw_element_type
T raw_element_type
Raw element type actually used for storage.
Definition: simd_mask_impl.hpp:65
simd
Definition: simd.hpp:1027
cl::__ESIMD_DNS::simd_mask_impl::raw_vector_type
typename base_type::raw_vector_type raw_vector_type
Underlying storage type for the entire vector.
Definition: simd_mask_impl.hpp:69
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(const T1 *ptr)
Load constructor.
Definition: simd_mask_impl.hpp:118
cl::__ESIMD_DNS::simd_mask_impl::simd_mask_impl
simd_mask_impl(const simd_mask_impl &other)
Copy constructor.
Definition: simd_mask_impl.hpp:77
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12