DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_bundle_impl.hpp
Go to the documentation of this file.
1 //==------- kernel_bundle_impl.hpp - SYCL kernel_bundle_impl ---------------==//
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 <detail/kernel_impl.hpp>
14 #include <sycl/backend_types.hpp>
15 #include <sycl/context.hpp>
16 #include <sycl/detail/common.hpp>
17 #include <sycl/detail/pi.h>
18 #include <sycl/device.hpp>
19 #include <sycl/kernel_bundle.hpp>
20 
21 #include <algorithm>
22 #include <cassert>
23 #include <cstring>
24 #include <memory>
25 #include <vector>
26 
27 namespace sycl {
29 namespace detail {
30 
31 static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,
32  const context &Context) {
33  return std::all_of(
34  Devices.begin(), Devices.end(), [&Context](const device &Dev) {
35  return getSyclObjImpl(Context)->isDeviceValid(getSyclObjImpl(Dev));
36  });
37 }
38 
39 static bool checkAllDevicesHaveAspect(const std::vector<device> &Devices,
40  aspect Aspect) {
41  return std::all_of(Devices.begin(), Devices.end(),
42  [&Aspect](const device &Dev) { return Dev.has(Aspect); });
43 }
44 
45 // The class is an impl counterpart of the sycl::kernel_bundle.
46 // It provides an access and utilities to manage set of sycl::device_images
47 // objects.
49 
50  using SpecConstMapT = std::map<std::string, std::vector<unsigned char>>;
51 
52  void common_ctor_checks(bundle_state State) {
53  const bool AllDevicesInTheContext =
54  checkAllDevicesAreInContext(MDevices, MContext);
55  if (MDevices.empty() || !AllDevicesInTheContext)
56  throw sycl::exception(
57  make_error_code(errc::invalid),
58  "Not all devices are associated with the context or "
59  "vector of devices is empty");
60 
61  if (bundle_state::input == State &&
62  !checkAllDevicesHaveAspect(MDevices, aspect::online_compiler))
63  throw sycl::exception(make_error_code(errc::invalid),
64  "Not all devices have aspect::online_compiler");
65 
66  if (bundle_state::object == State &&
67  !checkAllDevicesHaveAspect(MDevices, aspect::online_linker))
68  throw sycl::exception(make_error_code(errc::invalid),
69  "Not all devices have aspect::online_linker");
70  }
71 
72 public:
73  kernel_bundle_impl(context Ctx, std::vector<device> Devs, bundle_state State)
74  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
75 
76  common_ctor_checks(State);
77 
78  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
79  MContext, MDevices, State);
80  }
81 
82  // Interop constructor used by make_kernel
83  kernel_bundle_impl(context Ctx, std::vector<device> Devs)
84  : MContext(Ctx), MDevices(Devs), MState(bundle_state::executable) {
85  if (!checkAllDevicesAreInContext(Devs, Ctx))
86  throw sycl::exception(
87  make_error_code(errc::invalid),
88  "Not all devices are associated with the context or "
89  "vector of devices is empty");
90  MIsInterop = true;
91  }
92 
93  // Interop constructor
94  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
95  device_image_plain &DevImage)
96  : kernel_bundle_impl(Ctx, Devs) {
97  MDeviceImages.push_back(DevImage);
98  }
99 
100  // Matches sycl::build and sycl::compile
101  // Have one constructor because sycl::build and sycl::compile have the same
102  // signature
104  std::vector<device> Devs, const property_list &PropList,
105  bundle_state TargetState)
106  : MContext(InputBundle.get_context()), MDevices(std::move(Devs)),
107  MState(TargetState) {
108 
109  MSpecConstValues = getSyclObjImpl(InputBundle)->get_spec_const_map_ref();
110 
111  const std::vector<device> &InputBundleDevices =
112  getSyclObjImpl(InputBundle)->get_devices();
113  const bool AllDevsAssociatedWithInputBundle =
114  std::all_of(MDevices.begin(), MDevices.end(),
115  [&InputBundleDevices](const device &Dev) {
116  return InputBundleDevices.end() !=
117  std::find(InputBundleDevices.begin(),
118  InputBundleDevices.end(), Dev);
119  });
120  if (MDevices.empty() || !AllDevsAssociatedWithInputBundle)
121  throw sycl::exception(
122  make_error_code(errc::invalid),
123  "Not all devices are in the set of associated "
124  "devices for input bundle or vector of devices is empty");
125 
126  for (const device_image_plain &DeviceImage : InputBundle) {
127  // Skip images which are not compatible with devices provided
128  if (std::none_of(
129  MDevices.begin(), MDevices.end(),
130  [&DeviceImage](const device &Dev) {
131  return getSyclObjImpl(DeviceImage)->compatible_with_device(Dev);
132  }))
133  continue;
134 
135  switch (TargetState) {
136  case bundle_state::object:
137  MDeviceImages.push_back(detail::ProgramManager::getInstance().compile(
138  DeviceImage, MDevices, PropList));
139  break;
140  case bundle_state::executable:
141  MDeviceImages.push_back(detail::ProgramManager::getInstance().build(
142  DeviceImage, MDevices, PropList));
143  break;
144  case bundle_state::input:
145  throw sycl::runtime_error(
146  "Internal error. The target state should not be input",
147  PI_ERROR_INVALID_OPERATION);
148  break;
149  }
150  }
151  }
152 
153  // Matches sycl::link
155  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
156  std::vector<device> Devs, const property_list &PropList)
157  : MDevices(std::move(Devs)), MState(bundle_state::executable) {
158 
159  if (MDevices.empty())
160  throw sycl::exception(make_error_code(errc::invalid),
161  "Vector of devices is empty");
162 
163  if (ObjectBundles.empty())
164  return;
165 
166  MContext = ObjectBundles[0].get_context();
167  for (size_t I = 1; I < ObjectBundles.size(); ++I) {
168  if (ObjectBundles[I].get_context() != MContext)
169  throw sycl::exception(
170  make_error_code(errc::invalid),
171  "Not all input bundles have the same associated context");
172  }
173 
174  // Check if any of the devices in devs are not in the set of associated
175  // devices for any of the bundles in ObjectBundles
176  const bool AllDevsAssociatedWithInputBundles = std::all_of(
177  MDevices.begin(), MDevices.end(), [&ObjectBundles](const device &Dev) {
178  // Number of devices is expected to be small
179  return std::all_of(
180  ObjectBundles.begin(), ObjectBundles.end(),
181  [&Dev](const kernel_bundle<bundle_state::object> &KernelBundle) {
182  const std::vector<device> &BundleDevices =
183  getSyclObjImpl(KernelBundle)->get_devices();
184  return BundleDevices.end() != std::find(BundleDevices.begin(),
185  BundleDevices.end(),
186  Dev);
187  });
188  });
189  if (!AllDevsAssociatedWithInputBundles)
190  throw sycl::exception(make_error_code(errc::invalid),
191  "Not all devices are in the set of associated "
192  "devices for input bundles");
193 
194  // TODO: Unify with c'tor for sycl::comile and sycl::build by calling
195  // sycl::join on vector of kernel_bundles
196 
197  // The loop below just links each device image separately, not linking any
198  // two device images together. This is correct so long as each device image
199  // has no unresolved symbols. That's the case when device images are created
200  // from generic SYCL APIs. There's no way in generic SYCL to create a kernel
201  // which references an undefined symbol. If we decide in the future to allow
202  // a backend interop API to create a "sycl::kernel_bundle" that references
203  // undefined symbols, then the logic in this loop will need to be changed.
204  for (const kernel_bundle<bundle_state::object> &ObjectBundle :
205  ObjectBundles) {
206  for (const device_image_plain &DeviceImage : ObjectBundle) {
207 
208  // Skip images which are not compatible with devices provided
209  if (std::none_of(MDevices.begin(), MDevices.end(),
210  [&DeviceImage](const device &Dev) {
211  return getSyclObjImpl(DeviceImage)
212  ->compatible_with_device(Dev);
213  }))
214  continue;
215 
216  const std::vector<device_image_plain> VectorOfOneImage{DeviceImage};
217  std::vector<device_image_plain> LinkedResults =
218  detail::ProgramManager::getInstance().link(VectorOfOneImage,
219  MDevices, PropList);
220  MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
221  LinkedResults.end());
222  }
223  }
224 
225  for (const kernel_bundle<bundle_state::object> &Bundle : ObjectBundles) {
226  const KernelBundleImplPtr BundlePtr = getSyclObjImpl(Bundle);
227  for (const std::pair<const std::string, std::vector<unsigned char>>
228  &SpecConst : BundlePtr->MSpecConstValues) {
229  MSpecConstValues[SpecConst.first] = SpecConst.second;
230  }
231  }
232  }
233 
234  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
235  const std::vector<kernel_id> &KernelIDs,
236  bundle_state State)
237  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
238 
239  // TODO: Add a check that all kernel ids are compatible with at least one
240  // device in Devs
241  common_ctor_checks(State);
242 
243  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
244  MContext, MDevices, KernelIDs, State);
245  }
246 
247  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
248  const DevImgSelectorImpl &Selector, bundle_state State)
249  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
250 
251  common_ctor_checks(State);
252 
253  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
254  MContext, MDevices, Selector, State);
255  }
256 
257  // C'tor matches sycl::join API
258  kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
259  bundle_state State)
260  : MState(State) {
261  if (Bundles.empty())
262  return;
263 
264  MContext = Bundles[0]->MContext;
265  MDevices = Bundles[0]->MDevices;
266  for (size_t I = 1; I < Bundles.size(); ++I) {
267  if (Bundles[I]->MContext != MContext)
268  throw sycl::exception(
269  make_error_code(errc::invalid),
270  "Not all input bundles have the same associated context.");
271  if (Bundles[I]->MDevices != MDevices)
272  throw sycl::exception(
273  make_error_code(errc::invalid),
274  "Not all input bundles have the same set of associated devices.");
275  }
276 
277  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
278 
279  MDeviceImages.insert(MDeviceImages.end(), Bundle->MDeviceImages.begin(),
280  Bundle->MDeviceImages.end());
281  }
282 
283  std::sort(MDeviceImages.begin(), MDeviceImages.end(),
285 
286  if (get_bundle_state() == bundle_state::input) {
287  // Copy spec constants values from the device images to be removed.
288  auto MergeSpecConstants = [this](const device_image_plain &Img) {
289  const detail::DeviceImageImplPtr &ImgImpl = getSyclObjImpl(Img);
290  const std::map<std::string,
291  std::vector<device_image_impl::SpecConstDescT>>
292  &SpecConsts = ImgImpl->get_spec_const_data_ref();
293  const std::vector<unsigned char> &Blob =
294  ImgImpl->get_spec_const_blob_ref();
295  for (const std::pair<const std::string,
296  std::vector<device_image_impl::SpecConstDescT>>
297  &SpecConst : SpecConsts) {
298  if (SpecConst.second.front().IsSet)
299  set_specialization_constant_raw_value(
300  SpecConst.first.c_str(),
301  Blob.data() + SpecConst.second.front().BlobOffset,
302  SpecConst.second.back().CompositeOffset +
303  SpecConst.second.back().Size);
304  }
305  };
306  std::for_each(MDeviceImages.begin(), MDeviceImages.end(),
307  MergeSpecConstants);
308  }
309 
310  const auto DevImgIt =
311  std::unique(MDeviceImages.begin(), MDeviceImages.end());
312 
313  // Remove duplicate device images.
314  MDeviceImages.erase(DevImgIt, MDeviceImages.end());
315 
316  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
317  for (const std::pair<const std::string, std::vector<unsigned char>>
318  &SpecConst : Bundle->MSpecConstValues) {
319  set_specialization_constant_raw_value(SpecConst.first.c_str(),
320  SpecConst.second.data(),
321  SpecConst.second.size());
322  }
323  }
324  }
325 
326  bool empty() const noexcept { return MDeviceImages.empty(); }
327 
328  backend get_backend() const noexcept {
329  return MContext.get_platform().get_backend();
330  }
331 
332  context get_context() const noexcept { return MContext; }
333 
334  const std::vector<device> &get_devices() const noexcept { return MDevices; }
335 
336  std::vector<kernel_id> get_kernel_ids() const {
337  // Collect kernel ids from all device images, then remove duplicates
338 
339  std::vector<kernel_id> Result;
340  for (const device_image_plain &DeviceImage : MDeviceImages) {
341  const std::vector<kernel_id> &KernelIDs =
342  getSyclObjImpl(DeviceImage)->get_kernel_ids();
343 
344  Result.insert(Result.end(), KernelIDs.begin(), KernelIDs.end());
345  }
346  std::sort(Result.begin(), Result.end(), LessByNameComp{});
347 
348  auto NewIt = std::unique(Result.begin(), Result.end(), EqualByNameComp{});
349  Result.erase(NewIt, Result.end());
350 
351  return Result;
352  }
353 
354  kernel
355  get_kernel(const kernel_id &KernelID,
356  const std::shared_ptr<detail::kernel_bundle_impl> &Self) const {
357 
358  auto It = std::find_if(MDeviceImages.begin(), MDeviceImages.end(),
359  [&KernelID](const device_image_plain &DeviceImage) {
360  return DeviceImage.has_kernel(KernelID);
361  });
362 
363  if (MDeviceImages.end() == It)
364  throw sycl::exception(make_error_code(errc::invalid),
365  "The kernel bundle does not contain the kernel "
366  "identified by kernelId.");
367 
368  const std::shared_ptr<detail::device_image_impl> &DeviceImageImpl =
370 
371  RT::PiKernel Kernel = nullptr;
372  std::tie(Kernel, std::ignore) =
373  detail::ProgramManager::getInstance().getOrCreateKernel(
374  MContext, KernelID.get_name(), /*PropList=*/{},
375  DeviceImageImpl->get_program_ref());
376 
377  std::shared_ptr<kernel_impl> KernelImpl = std::make_shared<kernel_impl>(
378  Kernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self);
379 
380  return detail::createSyclObjFromImpl<kernel>(KernelImpl);
381  }
382 
383  bool has_kernel(const kernel_id &KernelID) const noexcept {
384  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
385  [&KernelID](const device_image_plain &DeviceImage) {
386  return DeviceImage.has_kernel(KernelID);
387  });
388  }
389 
390  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
391  return std::any_of(
392  MDeviceImages.begin(), MDeviceImages.end(),
393  [&KernelID, &Dev](const device_image_plain &DeviceImage) {
394  return DeviceImage.has_kernel(KernelID, Dev);
395  });
396  }
397 
398  bool contains_specialization_constants() const noexcept {
399  return std::any_of(
400  MDeviceImages.begin(), MDeviceImages.end(),
401  [](const device_image_plain &DeviceImage) {
402  return getSyclObjImpl(DeviceImage)->has_specialization_constants();
403  });
404  }
405 
406  bool native_specialization_constant() const noexcept {
407  return std::all_of(MDeviceImages.begin(), MDeviceImages.end(),
408  [](const device_image_plain &DeviceImage) {
409  return getSyclObjImpl(DeviceImage)
410  ->all_specialization_constant_native();
411  });
412  }
413 
414  bool has_specialization_constant(const char *SpecName) const noexcept {
415  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
416  [SpecName](const device_image_plain &DeviceImage) {
417  return getSyclObjImpl(DeviceImage)
418  ->has_specialization_constant(SpecName);
419  });
420  }
421 
422  void set_specialization_constant_raw_value(const char *SpecName,
423  const void *Value,
424  size_t Size) noexcept {
425  if (has_specialization_constant(SpecName))
426  for (const device_image_plain &DeviceImage : MDeviceImages)
427  getSyclObjImpl(DeviceImage)
428  ->set_specialization_constant_raw_value(SpecName, Value);
429  else {
430  const auto *DataPtr = static_cast<const unsigned char *>(Value);
431  std::vector<unsigned char> &Val = MSpecConstValues[std::string{SpecName}];
432  Val.resize(Size);
433  Val.insert(Val.begin(), DataPtr, DataPtr + Size);
434  }
435  }
436 
437  void get_specialization_constant_raw_value(const char *SpecName,
438  void *ValueRet) const noexcept {
439  for (const device_image_plain &DeviceImage : MDeviceImages)
440  if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) {
441  getSyclObjImpl(DeviceImage)
442  ->get_specialization_constant_raw_value(SpecName, ValueRet);
443  return;
444  }
445 
446  // Specialization constant wasn't found in any of the device images,
447  // try to fetch value from kernel_bundle.
448  if (MSpecConstValues.count(std::string{SpecName}) != 0) {
449  const std::vector<unsigned char> &Val =
450  MSpecConstValues.at(std::string{SpecName});
451  auto *Dest = static_cast<unsigned char *>(ValueRet);
452  std::uninitialized_copy(Val.begin(), Val.end(), Dest);
453  return;
454  }
455 
456  assert(false &&
457  "get_specialization_constant_raw_value called for missing constant");
458  }
459 
460  bool is_specialization_constant_set(const char *SpecName) const noexcept {
461  bool SetInDevImg =
462  std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
463  [SpecName](const device_image_plain &DeviceImage) {
464  return getSyclObjImpl(DeviceImage)
465  ->is_specialization_constant_set(SpecName);
466  });
467  return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0;
468  }
469 
470  const device_image_plain *begin() const { return MDeviceImages.data(); }
471 
472  const device_image_plain *end() const {
473  return MDeviceImages.data() + MDeviceImages.size();
474  }
475 
476  size_t size() const noexcept { return MDeviceImages.size(); }
477 
478  bundle_state get_bundle_state() const { return MState; }
479 
480  const SpecConstMapT &get_spec_const_map_ref() const noexcept {
481  return MSpecConstValues;
482  }
483 
484  bool isInterop() const { return MIsInterop; }
485 
486  bool add_kernel(const kernel_id &KernelID, const device &Dev) {
487  // Skip if kernel is already there
488  if (has_kernel(KernelID, Dev))
489  return true;
490 
491  // First try and get images in current bundle state
492  const bundle_state BundleState = get_bundle_state();
493  std::vector<device_image_plain> NewDevImgs =
494  detail::ProgramManager::getInstance().getSYCLDeviceImages(
495  MContext, {Dev}, {KernelID}, BundleState);
496 
497  // No images found so we report as not inserted
498  if (NewDevImgs.empty())
499  return false;
500 
501  // Propagate already set specialization constants to the new images
502  for (device_image_plain &DevImg : NewDevImgs)
503  for (auto SpecConst : MSpecConstValues)
504  getSyclObjImpl(DevImg)->set_specialization_constant_raw_value(
505  SpecConst.first.c_str(), SpecConst.second.data());
506 
507  // Add the images to the collection
508  MDeviceImages.insert(MDeviceImages.end(), NewDevImgs.begin(),
509  NewDevImgs.end());
510  return true;
511  }
512 
513 private:
514  context MContext;
515  std::vector<device> MDevices;
516  std::vector<device_image_plain> MDeviceImages;
517  // This map stores values for specialization constants, that are missing
518  // from any device image.
519  SpecConstMapT MSpecConstValues;
520  bool MIsInterop = false;
521  bundle_state MState;
522 };
523 
524 } // namespace detail
525 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
526 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:39
bool has_kernel(const kernel_id &KernelID) const noexcept
const SpecConstMapT & get_spec_const_map_ref() const noexcept
kernel_bundle_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, std::vector< device > Devs, const property_list &PropList)
kernel_bundle_impl(context Ctx, std::vector< device > Devs)
bool add_kernel(const kernel_id &KernelID, const device &Dev)
kernel_bundle_impl(context Ctx, std::vector< device > Devs, device_image_plain &DevImage)
const device_image_plain * end() const
const std::vector< device > & get_devices() const noexcept
void set_specialization_constant_raw_value(const char *SpecName, const void *Value, size_t Size) noexcept
bool contains_specialization_constants() const noexcept
std::vector< kernel_id > get_kernel_ids() const
bool is_specialization_constant_set(const char *SpecName) const noexcept
bool native_specialization_constant() const noexcept
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const std::vector< kernel_id > &KernelIDs, bundle_state State)
kernel_bundle_impl(const kernel_bundle< bundle_state::input > &InputBundle, std::vector< device > Devs, const property_list &PropList, bundle_state TargetState)
const device_image_plain * begin() const
kernel_bundle_impl(context Ctx, std::vector< device > Devs, bundle_state State)
void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet) const noexcept
bool has_specialization_constant(const char *SpecName) const noexcept
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
kernel get_kernel(const kernel_id &KernelID, const std::shared_ptr< detail::kernel_bundle_impl > &Self) const
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const DevImgSelectorImpl &Selector, bundle_state State)
kernel_bundle_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:47
The kernel_bundle class represents collection of device images in a particular state.
Objects of the class identify kernel is some kernel_bundle related APIs.
const char * get_name() const noexcept
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:69
Objects of the property_list class are containers for the SYCL properties.
#define __SYCL_INLINE_VER_NAMESPACE(X)
::pi_kernel PiKernel
Definition: pi.hpp:116
std::shared_ptr< device_image_impl > DeviceImageImplPtr
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:240
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
auto tie(Ts &...Args)
Definition: tuple.hpp:40
Function for_each(Group g, Ptr first, Ptr last, Function f)
kernel_bundle< bundle_state::executable > build(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:91
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept