DPC++ Runtime
Runtime libraries for oneAPI DPC++
bindless_images.hpp
Go to the documentation of this file.
1 //==----------- bindless_images.hpp --- SYCL bindless images ---------------==//
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/context.hpp> // for context
12 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
13 #include <sycl/detail/pi.h> // for pi_uint64
14 #include <sycl/device.hpp> // for device
15 #include <sycl/ext/oneapi/bindless_images_descriptor.hpp> // for image_desc...
16 #include <sycl/ext/oneapi/bindless_images_interop.hpp> // for interop_me...
17 #include <sycl/ext/oneapi/bindless_images_memory.hpp> // for image_mem_...
18 #include <sycl/ext/oneapi/bindless_images_sampler.hpp> // for bindless_i...
19 #include <sycl/image.hpp> // for image_chan...
20 #include <sycl/queue.hpp> // for queue
21 #include <sycl/range.hpp> // for range
22 
23 #include <assert.h> // for assert
24 #include <stddef.h> // for size_t
25 #include <type_traits> // for is_scalar
26 
27 #ifdef __SYCL_DEVICE_ONLY__
28 #include <sycl/detail/image_ocl_types.hpp> // for __invoke__*
29 #endif
30 
31 namespace sycl {
32 inline namespace _V1 {
33 namespace ext::oneapi::experimental {
34 
38 
40 
42  : raw_handle(raw_image_handle) {}
43 
45 };
46 
50 
52 
54 
56 };
57 
66 __SYCL_EXPORT image_mem_handle
67 alloc_image_mem(const image_descriptor &desc, const sycl::device &syclDevice,
68  const sycl::context &syclContext);
69 
77 __SYCL_EXPORT image_mem_handle alloc_image_mem(const image_descriptor &desc,
78  const sycl::queue &syclQueue);
79 
88 __SYCL_EXPORT void free_image_mem(image_mem_handle handle, image_type imageType,
89  const sycl::device &syclDevice,
90  const sycl::context &syclContext);
91 
99 __SYCL_EXPORT void free_image_mem(image_mem_handle handle, image_type imageType,
100  const sycl::queue &syclQueue);
101 
112  const image_mem_handle mipMem, unsigned int level,
113  const sycl::device &syclDevice, const sycl::context &syclContext);
114 
123 __SYCL_EXPORT image_mem_handle
124 get_mip_level_mem_handle(const image_mem_handle mipMem, unsigned int level,
125  const sycl::queue &syclQueue);
126 
139 template <typename ExternalMemHandleType>
142  const sycl::device &syclDevice, const sycl::context &syclContext);
143 
155 template <typename ExternalMemHandleType>
158  const sycl::queue &syclQueue);
159 
170 __SYCL_EXPORT
172  const image_descriptor &desc,
173  const sycl::device &syclDevice,
174  const sycl::context &syclContext);
175 
185 __SYCL_EXPORT
187  const image_descriptor &desc,
188  const sycl::queue &syclQueue);
189 
203 template <typename ExternalSemaphoreHandleType>
206  externalSemaphoreDesc,
207  const sycl::device &syclDevice, const sycl::context &syclContext);
208 
220 template <typename ExternalSemaphoreHandleType>
223  externalSemaphoreDesc,
224  const sycl::queue &syclQueue);
225 
235 __SYCL_EXPORT void
237  const sycl::device &syclDevice,
238  const sycl::context &syclContext);
239 
247 __SYCL_EXPORT void
249  const sycl::queue &syclQueue);
250 
260 __SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle,
261  const sycl::device &syclDevice,
262  const sycl::context &syclContext);
263 
271 __SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle,
272  const sycl::queue &syclQueue);
273 
283 __SYCL_EXPORT unsampled_image_handle
284 create_image(image_mem &memHandle, const image_descriptor &desc,
285  const sycl::device &syclDevice, const sycl::context &syclContext);
286 
295 __SYCL_EXPORT unsampled_image_handle create_image(image_mem &memHandle,
296  const image_descriptor &desc,
297  const sycl::queue &syclQueue);
298 
308 __SYCL_EXPORT unsampled_image_handle
309 create_image(image_mem_handle memHandle, const image_descriptor &desc,
310  const sycl::device &syclDevice, const sycl::context &syclContext);
311 
321  const image_descriptor &desc,
322  const sycl::queue &syclQueue);
323 
335 __SYCL_EXPORT sampled_image_handle
336 create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler,
337  const image_descriptor &desc, const sycl::device &syclDevice,
338  const sycl::context &syclContext);
339 
350 __SYCL_EXPORT sampled_image_handle
351 create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler,
352  const image_descriptor &desc, const sycl::queue &syclQueue);
353 
364 __SYCL_EXPORT sampled_image_handle
365 create_image(image_mem &memHandle, const bindless_image_sampler &sampler,
366  const image_descriptor &desc, const sycl::device &syclDevice,
367  const sycl::context &syclContext);
368 
378 __SYCL_EXPORT sampled_image_handle
379 create_image(image_mem &memHandle, const bindless_image_sampler &sampler,
380  const image_descriptor &desc, const sycl::queue &syclQueue);
381 
392 __SYCL_EXPORT sampled_image_handle
393 create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler,
394  const image_descriptor &desc, const sycl::device &syclDevice,
395  const sycl::context &syclContext);
396 
406 __SYCL_EXPORT sampled_image_handle
407 create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler,
408  const image_descriptor &desc, const sycl::queue &syclQueue);
409 
418 __SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle,
419  const sycl::device &syclDevice,
420  const sycl::context &syclContext);
421 
429 __SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle,
430  const sycl::queue &syclQueue);
431 
440 __SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle,
441  const sycl::device &syclDevice,
442  const sycl::context &syclContext);
443 
451 __SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle,
452  const sycl::queue &syclQueue);
453 
464 __SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
465  size_t widthInBytes, size_t height,
466  unsigned int elementSizeBytes,
467  const sycl::queue &syclQueue);
468 
480 __SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
481  size_t widthInBytes, size_t height,
482  unsigned int elementSizeBytes,
483  const sycl::device &syclDevice,
484  const sycl::context &syclContext);
485 
494 __SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
495  const image_descriptor &desc,
496  const sycl::queue &syclQueue);
497 
507 __SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
508  const image_descriptor &desc,
509  const sycl::device &syclDevice,
510  const sycl::context &syclContext);
511 
520 __SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle,
521  const sycl::device &syclDevice,
522  const sycl::context &syclContext);
523 
531 __SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle,
532  const sycl::queue &syclQueue);
533 
542 __SYCL_EXPORT sycl::image_channel_type
544  const sycl::device &syclDevice,
545  const sycl::context &syclContext);
546 
554 __SYCL_EXPORT sycl::image_channel_type
556  const sycl::queue &syclQueue);
557 
566 __SYCL_EXPORT unsigned int
568  const sycl::device &syclDevice,
569  const sycl::context &syclContext);
570 
578 __SYCL_EXPORT unsigned int
580  const sycl::queue &syclQueue);
581 
582 namespace detail {
583 
584 // is sycl::vec
585 template <typename T> struct is_vec {
586  static constexpr bool value = false;
587 };
588 template <typename T, int N> struct is_vec<sycl::vec<T, N>> {
589  static constexpr bool value = true;
590 };
591 template <typename T> inline constexpr bool is_vec_v = is_vec<T>::value;
592 
593 // Get the number of coordinates
594 template <typename CoordT> constexpr size_t coord_size() {
595  if constexpr (std::is_scalar_v<CoordT>) {
596  return 1;
597  } else {
598  return CoordT::size();
599  }
600 }
601 
602 // bit_cast Color to a type the backend is known to accept
603 template <typename DataT> constexpr auto convert_color(DataT Color) {
604  constexpr size_t dataSize = sizeof(DataT);
605  static_assert(
606  dataSize == 1 || dataSize == 2 || dataSize == 4 || dataSize == 8 ||
607  dataSize == 16,
608  "Expected input data type to be of size 1, 2, 4, 8, or 16 bytes.");
609 
610  if constexpr (dataSize == 1) {
611  return sycl::bit_cast<uint8_t>(Color);
612  } else if constexpr (dataSize == 2) {
613  return sycl::bit_cast<uint16_t>(Color);
614  } else if constexpr (dataSize == 4) {
615  return sycl::bit_cast<uint32_t>(Color);
616  } else if constexpr (dataSize == 8) {
617  return sycl::bit_cast<sycl::vec<uint32_t, 2>>(Color);
618  } else { // dataSize == 16
619  return sycl::bit_cast<sycl::vec<uint32_t, 4>>(Color);
620  }
621 }
622 
623 // assert coords or elements of coords is of an integer type
624 template <typename CoordT> constexpr void assert_unsampled_coords() {
625  if constexpr (std::is_scalar_v<CoordT>) {
626  static_assert(std::is_same_v<CoordT, int>,
627  "Expected integer coordinate data type");
628  } else {
629  static_assert(is_vec_v<CoordT>, "Expected sycl::vec coordinates");
630  static_assert(std::is_same_v<typename CoordT::element_type, int>,
631  "Expected integer coordinates data type");
632  }
633 }
634 
635 template <typename CoordT> constexpr bool are_floating_coords() {
636  if constexpr (is_vec_v<CoordT>) {
637  return std::is_same_v<typename CoordT::element_type, float>;
638  } else {
639  return std::is_same_v<CoordT, float>;
640  }
641 }
642 
643 template <typename CoordT> constexpr bool are_integer_coords() {
644  if constexpr (is_vec_v<CoordT>) {
645  return std::is_same_v<typename CoordT::element_type, int>;
646  } else {
647  return std::is_same_v<CoordT, int>;
648  }
649 }
650 
651 template <typename CoordT> constexpr void assert_coords_type() {
652  static_assert(are_floating_coords<CoordT>() || are_integer_coords<CoordT>(),
653  "Expected coordinates to be of `float` or `int` type, or "
654  "vectors of these types.");
655 }
656 
657 // assert coords or elements of coords is of a float type
658 template <typename CoordT> constexpr void assert_sample_coords() {
659  if constexpr (std::is_scalar_v<CoordT>) {
660  static_assert(std::is_same_v<CoordT, float>,
661  "Expected float coordinate data type");
662  } else {
663  static_assert(is_vec_v<CoordT>, "Expected sycl::vec coordinates");
664  static_assert(std::is_same_v<typename CoordT::element_type, float>,
665  "Expected float coordinates data type");
666  }
667 }
668 
669 // assert coords or elements of coords is of a int type
670 template <typename CoordT> constexpr void assert_fetch_coords() {
671  if constexpr (std::is_scalar_v<CoordT>) {
672  static_assert(std::is_same_v<CoordT, int>,
673  "Expected int coordinate data type");
674  } else {
675  static_assert(is_vec_v<CoordT>, "Expected sycl::vec coordinates");
676  static_assert(std::is_same_v<typename CoordT::element_type, int>,
677  "Expected int coordinates data type");
678  }
679 }
680 
681 template <typename DataT> constexpr bool is_data_size_valid() {
682  return (sizeof(DataT) == 1) || (sizeof(DataT) == 2) || (sizeof(DataT) == 4) ||
683  (sizeof(DataT) == 8) || (sizeof(DataT) == 16);
684 }
685 
686 template <typename DataT> constexpr bool is_recognized_standard_type() {
687  return is_data_size_valid<DataT>() &&
688  (is_vec_v<DataT> || std::is_scalar_v<DataT> ||
689  std::is_floating_point_v<DataT> || std::is_same_v<DataT, sycl::half>);
690 }
691 
692 #ifdef __SYCL_DEVICE_ONLY__
693 
694 // Image types used for generating SPIR-V
695 template <int NDims>
696 using OCLImageTyRead =
697  typename sycl::detail::opencl_image_type<NDims, sycl::access::mode::read,
698  sycl::access::target::image>::type;
699 
700 template <int NDims>
701 using OCLImageTyWrite =
702  typename sycl::detail::opencl_image_type<NDims, sycl::access::mode::write,
703  sycl::access::target::image>::type;
704 
705 template <int NDims>
706 using OCLImageArrayTyRead = typename sycl::detail::opencl_image_type<
707  NDims, sycl::access::mode::read, sycl::access::target::image_array>::type;
708 
709 template <int NDims>
710 using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type<
711  NDims, sycl::access::mode::write, sycl::access::target::image_array>::type;
712 
713 template <int NDims>
714 using OCLSampledImageArrayTyRead =
715  typename sycl::detail::sampled_opencl_image_type<
716  detail::OCLImageArrayTyRead<NDims>>::type;
717 
718 // Macros are required because it is not legal for a function to return
719 // a variable of type 'opencl_image_type'.
720 #if defined(__SPIR__)
721 #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) \
722  __spirv_ConvertHandleToImageINTEL<ImageType>(raw_handle)
723 
724 #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) \
725  __spirv_ConvertHandleToSampledImageINTEL< \
726  typename sycl::detail::sampled_opencl_image_type< \
727  detail::OCLImageTyRead<NDims>>::type>(raw_handle)
728 
729 #define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) \
730  __spirv_ConvertHandleToSampledImageINTEL< \
731  typename sycl::detail::sampled_opencl_image_type< \
732  detail::OCLImageArrayTyRead<NDims>>::type>(raw_handle)
733 
734 #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \
735  __invoke__ImageRead<DataT>(raw_handle, coords)
736 
737 #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \
738  __invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)
739 
740 #define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \
741  __invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)
742 
743 #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \
744  __invoke__ImageRead<DataT>(raw_handle, coordsLayer)
745 
746 #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \
747  __invoke__ImageWrite(raw_handle, coordsLayer, color)
748 
749 #define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
750  coordsLayer) \
751  __invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)
752 
753 #define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
754  coordsLayer) \
755  __invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)
756 
757 #else
758 #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle
759 
760 #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle
761 
762 #define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) raw_handle
763 
764 #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \
765  __invoke__ImageFetch<DataT>(raw_handle, coords)
766 
767 #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \
768  __invoke__SampledImageFetch<DataT>(raw_handle, coords)
769 
770 #define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \
771  __invoke__ImageRead<DataT>(raw_handle, coords)
772 
773 #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \
774  __invoke__ImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)
775 
776 #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \
777  __invoke__ImageArrayWrite(raw_handle, coords, arrayLayer, color)
778 
779 #define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
780  coordsLayer) \
781  __invoke__SampledImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)
782 
783 #define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
784  coordsLayer) \
785  __invoke__ImageArrayRead<DataT>(raw_handle, coords, arrayLayer)
786 
787 #endif
788 
789 #endif // __SYCL_DEVICE_ONLY__
790 
791 } // namespace detail
792 
814 template <typename DataT, typename HintT = DataT, typename CoordT>
815 DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
816  const CoordT &coords [[maybe_unused]]) {
817  detail::assert_fetch_coords<CoordT>();
818  constexpr size_t coordSize = detail::coord_size<CoordT>();
819  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
820  "Expected input coordinate to be have 1, 2, or 3 components "
821  "for 1D, 2D and 3D images, respectively.");
822 
823 #ifdef __SYCL_DEVICE_ONLY__
824  if constexpr (detail::is_recognized_standard_type<DataT>()) {
825  return FETCH_UNSAMPLED_IMAGE(
826  DataT,
827  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
828  detail::OCLImageTyRead<coordSize>),
829  coords);
830 
831  } else {
832  static_assert(sizeof(HintT) == sizeof(DataT),
833  "When trying to read a user-defined type, HintT must be of "
834  "the same size as the user-defined DataT.");
835  static_assert(detail::is_recognized_standard_type<HintT>(),
836  "HintT must always be a recognized standard type");
837  return sycl::bit_cast<DataT>(FETCH_UNSAMPLED_IMAGE(
838  HintT,
839  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
840  detail::OCLImageTyRead<coordSize>),
841  coords));
842  }
843 #else
844  assert(false); // Bindless images not yet implemented on host
845 #endif
846 }
847 
869 template <typename DataT, typename HintT = DataT, typename CoordT>
870 DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]],
871  const CoordT &coords [[maybe_unused]]) {
872  detail::assert_fetch_coords<CoordT>();
873  constexpr size_t coordSize = detail::coord_size<CoordT>();
874  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
875  "Expected input coordinate to be have 1, 2, or 3 components "
876  "for 1D, 2D and 3D images, respectively.");
877  static_assert(sizeof(HintT) == sizeof(DataT),
878  "When trying to read a user-defined type, HintT must be of "
879  "the same size as the user-defined DataT.");
880  static_assert(detail::is_recognized_standard_type<HintT>(),
881  "HintT must always be a recognized standard type");
882 
883 #ifdef __SYCL_DEVICE_ONLY__
884  if constexpr (detail::is_recognized_standard_type<DataT>()) {
885  return FETCH_SAMPLED_IMAGE(
886  DataT,
887  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
888  coords);
889  } else {
890  return sycl::bit_cast<DataT>(FETCH_SAMPLED_IMAGE(
891  HintT,
892  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
893  coords));
894  }
895 #else
896  assert(false); // Bindless images not yet implemented on host.
897 #endif
898 }
899 
921 template <typename DataT, typename HintT = DataT, typename CoordT>
922 DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]],
923  const CoordT &coords [[maybe_unused]]) {
924  detail::assert_sample_coords<CoordT>();
925  constexpr size_t coordSize = detail::coord_size<CoordT>();
926  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
927  "Expected input coordinate to be have 1, 2, or 3 components "
928  "for 1D, 2D and 3D images, respectively.");
929  static_assert(sizeof(HintT) == sizeof(DataT),
930  "When trying to read a user-defined type, HintT must be of "
931  "the same size as the user-defined DataT.");
932  static_assert(detail::is_recognized_standard_type<HintT>(),
933  "HintT must always be a recognized standard type");
934 
935 #ifdef __SYCL_DEVICE_ONLY__
936  if constexpr (detail::is_recognized_standard_type<DataT>()) {
937  return SAMPLE_IMAGE_READ(
938  DataT,
939  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
940  coords);
941  } else {
942  return sycl::bit_cast<DataT>(SAMPLE_IMAGE_READ(
943  HintT,
944  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
945  coords));
946  }
947 #else
948  assert(false); // Bindless images not yet implemented on host.
949 #endif
950 }
951 
967 template <typename DataT, typename HintT = DataT, typename CoordT>
968 DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
969  const CoordT &coords [[maybe_unused]],
970  const float level [[maybe_unused]]) {
971  detail::assert_sample_coords<CoordT>();
972  constexpr size_t coordSize = detail::coord_size<CoordT>();
973  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
974  "Expected input coordinate to be have 1, 2, or 3 components "
975  "for 1D, 2D and 3D images, respectively.");
976 
977 #ifdef __SYCL_DEVICE_ONLY__
978  if constexpr (detail::is_recognized_standard_type<DataT>()) {
979  return __invoke__ImageReadLod<DataT>(
980  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
981  coords, level);
982  } else {
983  static_assert(sizeof(HintT) == sizeof(DataT),
984  "When trying to read a user-defined type, HintT must be of "
985  "the same size as the user-defined DataT.");
986  static_assert(detail::is_recognized_standard_type<HintT>(),
987  "HintT must always be a recognized standard type");
988  return sycl::bit_cast<DataT>(__invoke__ImageReadLod<HintT>(
989  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
990  coords, level));
991  }
992 #else
993  assert(false); // Bindless images not yet implemented on host
994 #endif
995 }
996 
1013 template <typename DataT, typename HintT = DataT, typename CoordT>
1014 DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
1015  const CoordT &coords [[maybe_unused]],
1016  const CoordT &dX [[maybe_unused]],
1017  const CoordT &dY [[maybe_unused]]) {
1018  detail::assert_sample_coords<CoordT>();
1019  constexpr size_t coordSize = detail::coord_size<CoordT>();
1020  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
1021  "Expected input coordinates and gradients to have 1, 2, or 3 "
1022  "components for 1D, 2D, and 3D images, respectively.");
1023 
1024 #ifdef __SYCL_DEVICE_ONLY__
1025  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1026  return __invoke__ImageReadGrad<DataT>(
1027  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
1028  coords, dX, dY);
1029  } else {
1030  static_assert(sizeof(HintT) == sizeof(DataT),
1031  "When trying to read a user-defined type, HintT must be of "
1032  "the same size as the user-defined DataT.");
1033  static_assert(detail::is_recognized_standard_type<HintT>(),
1034  "HintT must always be a recognized standard type");
1035  return sycl::bit_cast<DataT>(__invoke__ImageReadGrad<HintT>(
1036  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
1037  coords, dX, dY));
1038  }
1039 #else
1040  assert(false); // Bindless images not yet implemented on host
1041 #endif
1042 }
1043 
1059 template <typename DataT, typename HintT = DataT, typename CoordT>
1061  [[maybe_unused]],
1062  const CoordT &coords [[maybe_unused]],
1063  unsigned int arrayLayer [[maybe_unused]]) {
1064  detail::assert_unsampled_coords<CoordT>();
1065  constexpr size_t coordSize = detail::coord_size<CoordT>();
1066  static_assert(coordSize == 1 || coordSize == 2,
1067  "Expected input coordinate to be have 1 or 2 components for 1D "
1068  "and 2D images respectively.");
1069 
1070 #ifdef __SYCL_DEVICE_ONLY__
1071  sycl::vec<int, coordSize + 1> coordsLayer{coords, arrayLayer};
1072  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1073  return FETCH_IMAGE_ARRAY(
1074  DataT,
1075  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1076  detail::OCLImageArrayTyRead<coordSize>),
1077  coords, arrayLayer, coordsLayer);
1078  } else {
1079  static_assert(sizeof(HintT) == sizeof(DataT),
1080  "When trying to fetch a user-defined type, HintT must be of "
1081  "the same size as the user-defined DataT.");
1082  static_assert(detail::is_recognized_standard_type<HintT>(),
1083  "HintT must always be a recognized standard type");
1084  return sycl::bit_cast<DataT>(FETCH_IMAGE_ARRAY(
1085  HintT,
1086  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1087  detail::OCLImageArrayTyRead<coordSize>),
1088  coords, arrayLayer, coordsLayer));
1089  }
1090 #else
1091  assert(false); // Bindless images not yet implemented on host.
1092 #endif
1093 }
1094 
1109 template <typename DataT, typename HintT = DataT>
1110 DataT fetch_cubemap(const unsampled_image_handle &imageHandle,
1111  const int2 &coords, unsigned int face) {
1112  return fetch_image_array<DataT, HintT>(imageHandle, coords, face);
1113 }
1114 
1129 template <typename DataT, typename HintT = DataT>
1130 DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]],
1131  const sycl::float3 &dirVec [[maybe_unused]]) {
1132  [[maybe_unused]] constexpr size_t NDims = 2;
1133 
1134 #ifdef __SYCL_DEVICE_ONLY__
1135  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1136  return __invoke__ImageReadCubemap<DataT, uint64_t>(
1137  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, NDims), dirVec);
1138  } else {
1139  static_assert(sizeof(HintT) == sizeof(DataT),
1140  "When trying to read a user-defined type, HintT must be of "
1141  "the same size as the user-defined DataT.");
1142  static_assert(detail::is_recognized_standard_type<HintT>(),
1143  "HintT must always be a recognized standard type");
1144  return sycl::bit_cast<DataT>(__invoke__ImageReadCubemap<HintT, uint64_t>(
1145  CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, NDims),
1146  dirVec));
1147  }
1148 #else
1149  assert(false); // Bindless images not yet implemented on host
1150 #endif
1151 }
1152 
1168 template <typename DataT, typename HintT = DataT, typename CoordT>
1169 DataT fetch_image_array(const sampled_image_handle &imageHandle
1170  [[maybe_unused]],
1171  const CoordT &coords [[maybe_unused]],
1172  unsigned int arrayLayer [[maybe_unused]]) {
1173  detail::assert_unsampled_coords<CoordT>();
1174  constexpr size_t coordSize = detail::coord_size<CoordT>();
1175  static_assert(coordSize == 1 || coordSize == 2,
1176  "Expected input coordinate to be have 1 or 2 components for 1D "
1177  "and 2D images respectively.");
1178 
1179 #ifdef __SYCL_DEVICE_ONLY__
1180  sycl::vec<int, coordSize + 1> coordsLayer{coords, arrayLayer};
1181  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1182  return FETCH_SAMPLED_IMAGE_ARRAY(DataT,
1183  CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
1184  imageHandle.raw_handle, coordSize),
1185  coords, arrayLayer, coordsLayer);
1186  } else {
1187  static_assert(sizeof(HintT) == sizeof(DataT),
1188  "When trying to fetch a user-defined type, HintT must be of "
1189  "the same size as the user-defined DataT.");
1190  static_assert(detail::is_recognized_standard_type<HintT>(),
1191  "HintT must always be a recognized standard type");
1192  return sycl::bit_cast<DataT>(
1193  FETCH_SAMPLED_IMAGE_ARRAY(HintT,
1194  CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
1195  imageHandle.raw_handle, coordSize),
1196  coords, arrayLayer, coordsLayer));
1197  }
1198 #else
1199  assert(false); // Bindless images not yet implemented on host.
1200 #endif
1201 }
1202 
1218 template <typename DataT, typename HintT = DataT, typename CoordT>
1220  [[maybe_unused]],
1221  const CoordT &coords [[maybe_unused]],
1222  unsigned int arrayLayer [[maybe_unused]]) {
1223  detail::assert_sample_coords<CoordT>();
1224  constexpr size_t coordSize = detail::coord_size<CoordT>();
1225  static_assert(coordSize == 1 || coordSize == 2,
1226  "Expected input coordinate to be have 1 or 2 components for 1D "
1227  "and 2D images respectively.");
1228 
1229 #ifdef __SYCL_DEVICE_ONLY__
1230  sycl::vec<float, coordSize + 1> coordsLayer{coords, arrayLayer};
1231  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1232  return READ_SAMPLED_IMAGE_ARRAY(DataT,
1233  CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
1234  imageHandle.raw_handle, coordSize),
1235  coords, arrayLayer, coordsLayer);
1236  } else {
1237  static_assert(sizeof(HintT) == sizeof(DataT),
1238  "When trying to fetch a user-defined type, HintT must be of "
1239  "the same size as the user-defined DataT.");
1240  static_assert(detail::is_recognized_standard_type<HintT>(),
1241  "HintT must always be a recognized standard type");
1242  return sycl::bit_cast<DataT>(
1243  READ_SAMPLED_IMAGE_ARRAY(HintT,
1244  CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
1245  imageHandle.raw_handle, coordSize),
1246  coords, arrayLayer, coordsLayer));
1247  }
1248 #else
1249  assert(false); // Bindless images not yet implemented on host.
1250 #endif
1251 }
1252 
1263 template <typename DataT, typename CoordT>
1264 void write_image(unsampled_image_handle imageHandle [[maybe_unused]],
1265  const CoordT &coords [[maybe_unused]],
1266  const DataT &color [[maybe_unused]]) {
1267  detail::assert_unsampled_coords<CoordT>();
1268  constexpr size_t coordSize = detail::coord_size<CoordT>();
1269  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
1270  "Expected input coordinate to be have 1, 2, or 3 components "
1271  "for 1D, 2D and 3D images, respectively.");
1272 
1273 #ifdef __SYCL_DEVICE_ONLY__
1274  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1275  __invoke__ImageWrite(
1276  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1277  detail::OCLImageTyWrite<coordSize>),
1278  coords, color);
1279  } else {
1280  // Convert DataT to a supported backend write type when user-defined type is
1281  // passed
1282  __invoke__ImageWrite(
1283  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1284  detail::OCLImageTyWrite<coordSize>),
1285  coords, detail::convert_color(color));
1286  }
1287 #else
1288  assert(false); // Bindless images not yet implemented on host
1289 #endif
1290 }
1291 
1303 template <typename DataT, typename CoordT>
1304 void write_image_array(unsampled_image_handle imageHandle [[maybe_unused]],
1305  const CoordT &coords [[maybe_unused]],
1306  unsigned int arrayLayer [[maybe_unused]],
1307  const DataT &color [[maybe_unused]]) {
1308  detail::assert_unsampled_coords<CoordT>();
1309  constexpr size_t coordSize = detail::coord_size<CoordT>();
1310  static_assert(coordSize == 1 || coordSize == 2,
1311  "Expected input coordinate to be have 1 or 2 components for 1D "
1312  "and 2D images respectively.");
1313 
1314 #ifdef __SYCL_DEVICE_ONLY__
1315  sycl::vec<int, coordSize + 1> coordsLayer{coords, arrayLayer};
1316  if constexpr (detail::is_recognized_standard_type<DataT>()) {
1317  WRITE_IMAGE_ARRAY(
1318  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1319  detail::OCLImageArrayTyWrite<coordSize>),
1320  coords, arrayLayer, coordsLayer, color);
1321  } else {
1322  // Convert DataT to a supported backend write type when user-defined type is
1323  // passed
1324  WRITE_IMAGE_ARRAY(
1325  CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
1326  detail::OCLImageArrayTyWrite<coordSize>),
1327  coords, arrayLayer, coordsLayer, detail::convert_color(color));
1328  }
1329 #else
1330  assert(false); // Bindless images not yet implemented on host.
1331 #endif
1332 }
1333 
1344 template <typename DataT>
1345 void write_cubemap(unsampled_image_handle imageHandle, const sycl::int2 &coords,
1346  int face, const DataT &color) {
1347  return write_image_array(imageHandle, coords, face, color);
1348 }
1349 
1350 } // namespace ext::oneapi::experimental
1351 
1353  const void *Src, ext::oneapi::experimental::image_mem_handle Dest,
1355  const detail::code_location &CodeLoc) {
1356  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1357  return submit(
1358  [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); },
1359  CodeLoc);
1360 }
1361 
1363  const void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
1366  sycl::range<3> CopyExtent, const detail::code_location &CodeLoc) {
1367  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1368  return submit(
1369  [&](handler &CGH) {
1370  CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
1371  DestImgDesc, CopyExtent);
1372  },
1373  CodeLoc);
1374 }
1375 
1377  const void *Src, ext::oneapi::experimental::image_mem_handle Dest,
1379  event DepEvent, const detail::code_location &CodeLoc) {
1380  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1381  return submit(
1382  [&](handler &CGH) {
1383  CGH.depends_on(DepEvent);
1384  CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
1385  },
1386  CodeLoc);
1387 }
1388 
1390  const void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
1393  sycl::range<3> CopyExtent, event DepEvent,
1394  const detail::code_location &CodeLoc) {
1395  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1396  return submit(
1397  [&](handler &CGH) {
1398  CGH.depends_on(DepEvent);
1399  CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
1400  DestImgDesc, CopyExtent);
1401  },
1402  CodeLoc);
1403 }
1404 
1406  const void *Src, ext::oneapi::experimental::image_mem_handle Dest,
1408  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1409  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1410  return submit(
1411  [&](handler &CGH) {
1412  CGH.depends_on(DepEvents);
1413  CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
1414  },
1415  CodeLoc);
1416 }
1417 
1419  const void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent,
1422  sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
1423  const detail::code_location &CodeLoc) {
1424  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1425  return submit(
1426  [&](handler &CGH) {
1427  CGH.depends_on(DepEvents);
1428  CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
1429  DestImgDesc, CopyExtent);
1430  },
1431  CodeLoc);
1432 }
1433 
1435  const ext::oneapi::experimental::image_mem_handle Src, void *Dest,
1437  const detail::code_location &CodeLoc) {
1438  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1439  return submit(
1440  [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); },
1441  CodeLoc);
1442 }
1443 
1446  sycl::range<3> SrcOffset,
1447  const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
1448  sycl::range<3> DestOffset, sycl::range<3> DestExtent,
1449  sycl::range<3> CopyExtent, const detail::code_location &CodeLoc) {
1450  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1451  return submit(
1452  [&](handler &CGH) {
1453  CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
1454  DestExtent, CopyExtent);
1455  },
1456  CodeLoc);
1457 }
1458 
1460  const ext::oneapi::experimental::image_mem_handle Src, void *Dest,
1462  event DepEvent, const detail::code_location &CodeLoc) {
1463  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1464  return submit(
1465  [&](handler &CGH) {
1466  CGH.depends_on(DepEvent);
1467  CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
1468  },
1469  CodeLoc);
1470 }
1471 
1474  sycl::range<3> SrcOffset,
1475  const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
1476  sycl::range<3> DestOffset, sycl::range<3> DestExtent,
1477  sycl::range<3> CopyExtent, event DepEvent,
1478  const detail::code_location &CodeLoc) {
1479  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1480  return submit(
1481  [&](handler &CGH) {
1482  CGH.depends_on(DepEvent);
1483  CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
1484  DestExtent, CopyExtent);
1485  },
1486  CodeLoc);
1487 }
1488 
1490  const ext::oneapi::experimental::image_mem_handle Src, void *Dest,
1492  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1493  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1494  return submit(
1495  [&](handler &CGH) {
1496  CGH.depends_on(DepEvents);
1497  CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
1498  },
1499  CodeLoc);
1500 }
1501 
1504  sycl::range<3> SrcOffset,
1505  const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest,
1506  sycl::range<3> DestOffset, sycl::range<3> DestExtent,
1507  sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
1508  const detail::code_location &CodeLoc) {
1509  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1510  return submit(
1511  [&](handler &CGH) {
1512  CGH.depends_on(DepEvents);
1513  CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
1514  DestExtent, CopyExtent);
1515  },
1516  CodeLoc);
1517 }
1518 
1520  const void *Src, void *Dest,
1521  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1522  size_t DeviceRowPitch, const detail::code_location &CodeLoc) {
1523  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1524  return submit(
1525  [&](handler &CGH) {
1526  CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
1527  },
1528  CodeLoc);
1529 }
1530 
1532  const void *Src, sycl::range<3> SrcOffset, void *Dest,
1533  sycl::range<3> DestOffset,
1534  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1535  size_t DeviceRowPitch, sycl::range<3> HostExtent, sycl::range<3> CopyExtent,
1536  const detail::code_location &CodeLoc) {
1537  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1538  return submit(
1539  [&](handler &CGH) {
1540  CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
1541  DeviceRowPitch, HostExtent, CopyExtent);
1542  },
1543  CodeLoc);
1544 }
1545 
1547  const void *Src, void *Dest,
1548  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1549  size_t DeviceRowPitch, event DepEvent,
1550  const detail::code_location &CodeLoc) {
1551  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1552  return submit(
1553  [&](handler &CGH) {
1554  CGH.depends_on(DepEvent);
1555  CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
1556  },
1557  CodeLoc);
1558 }
1559 
1564  event DepEvent, const detail::code_location &CodeLoc) {
1565  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1566  return submit(
1567  [&](handler &CGH) {
1568  CGH.depends_on(DepEvent);
1569  CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
1570  },
1571  CodeLoc);
1572 }
1573 
1578  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1579  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1580  return submit(
1581  [&](handler &CGH) {
1582  CGH.depends_on(DepEvents);
1583  CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
1584  },
1585  CodeLoc);
1586 }
1587 
1592  const detail::code_location &CodeLoc) {
1593  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1594  return submit(
1595  [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, ImageDesc); },
1596  CodeLoc);
1597 }
1598 
1600  const void *Src, sycl::range<3> SrcOffset, void *Dest,
1601  sycl::range<3> DestOffset,
1602  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1603  size_t DeviceRowPitch, sycl::range<3> HostExtent, sycl::range<3> CopyExtent,
1604  event DepEvent, const detail::code_location &CodeLoc) {
1605  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1606  return submit(
1607  [&](handler &CGH) {
1608  CGH.depends_on(DepEvent);
1609  CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
1610  DeviceRowPitch, HostExtent, CopyExtent);
1611  },
1612  CodeLoc);
1613 }
1614 
1616  const void *Src, void *Dest,
1617  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1618  size_t DeviceRowPitch, const std::vector<event> &DepEvents,
1619  const detail::code_location &CodeLoc) {
1620  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1621  return submit(
1622  [&](handler &CGH) {
1623  CGH.depends_on(DepEvents);
1624  CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
1625  },
1626  CodeLoc);
1627 }
1628 
1630  const void *Src, sycl::range<3> SrcOffset, void *Dest,
1631  sycl::range<3> DestOffset,
1632  const ext::oneapi::experimental::image_descriptor &DeviceImgDesc,
1633  size_t DeviceRowPitch, sycl::range<3> HostExtent, sycl::range<3> CopyExtent,
1634  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1635  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1636  return submit(
1637  [&](handler &CGH) {
1638  CGH.depends_on(DepEvents);
1639  CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
1640  DeviceRowPitch, HostExtent, CopyExtent);
1641  },
1642  CodeLoc);
1643 }
1644 
1647  event DepEvent, const detail::code_location &CodeLoc) {
1648  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1649  return submit(
1650  [&](handler &CGH) {
1651  CGH.depends_on(DepEvent);
1652  CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
1653  },
1654  CodeLoc);
1655 }
1656 
1659  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1660  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1661  return submit(
1662  [&](handler &CGH) {
1663  CGH.depends_on(DepEvents);
1664  CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
1665  },
1666  CodeLoc);
1667 }
1668 
1671  uint64_t WaitValue, const detail::code_location &CodeLoc) {
1672  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1673  return submit(
1674  [&](handler &CGH) {
1675  CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
1676  },
1677  CodeLoc);
1678 }
1679 
1682  uint64_t WaitValue, event DepEvent, const detail::code_location &CodeLoc) {
1683  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1684  return submit(
1685  [&](handler &CGH) {
1686  CGH.depends_on(DepEvent);
1687  CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
1688  },
1689  CodeLoc);
1690 }
1691 
1694  uint64_t WaitValue, const std::vector<event> &DepEvents,
1695  const detail::code_location &CodeLoc) {
1696  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1697  return submit(
1698  [&](handler &CGH) {
1699  CGH.depends_on(DepEvents);
1700  CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
1701  },
1702  CodeLoc);
1703 }
1704 
1707  const detail::code_location &CodeLoc) {
1708  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1709  return submit(
1710  [&](handler &CGH) {
1711  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
1712  },
1713  CodeLoc);
1714 }
1715 
1718  event DepEvent, const detail::code_location &CodeLoc) {
1719  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1720  return submit(
1721  [&](handler &CGH) {
1722  CGH.depends_on(DepEvent);
1723  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
1724  },
1725  CodeLoc);
1726 }
1727 
1730  const std::vector<event> &DepEvents, const detail::code_location &CodeLoc) {
1731  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1732  return submit(
1733  [&](handler &CGH) {
1734  CGH.depends_on(DepEvents);
1735  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
1736  },
1737  CodeLoc);
1738 }
1739 
1742  uint64_t SignalValue, const detail::code_location &CodeLoc) {
1743  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1744  return submit(
1745  [&](handler &CGH) {
1746  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
1747  },
1748  CodeLoc);
1749 }
1750 
1753  uint64_t SignalValue, event DepEvent,
1754  const detail::code_location &CodeLoc) {
1755  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1756  return submit(
1757  [&](handler &CGH) {
1758  CGH.depends_on(DepEvent);
1759  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
1760  },
1761  CodeLoc);
1762 }
1763 
1766  uint64_t SignalValue, const std::vector<event> &DepEvents,
1767  const detail::code_location &CodeLoc) {
1768  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
1769  return submit(
1770  [&](handler &CGH) {
1771  CGH.depends_on(DepEvents);
1772  CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
1773  },
1774  CodeLoc);
1775 }
1776 
1777 } // namespace _V1
1778 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:50
Data type that manages the code_location information in TLS.
Definition: common.hpp:129
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:64
An event object can be used to synchronize memory transfers, enqueues of kernels and signaling barrie...
Definition: event.hpp:44
Command group handler class.
Definition: handler.hpp:468
void depends_on(event Event)
Registers event dependencies on this command group.
Definition: handler.cpp:1491
void ext_oneapi_signal_external_semaphore(ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle)
Instruct the queue to signal the external semaphore once all previous commands submitted to the queue...
Definition: handler.cpp:1423
void ext_oneapi_wait_external_semaphore(ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle)
Submit a non-blocking device-side wait on an external.
Definition: handler.cpp:1381
void ext_oneapi_copy(const void *Src, ext::oneapi::experimental::image_mem_handle Dest, const ext::oneapi::experimental::image_descriptor &DestImgDesc)
Copies data from one memory region to another, where Src is a USM pointer and Dest is an opaque image...
Definition: handler.cpp:1004
Encapsulates a single SYCL queue which schedules kernels on a SYCL device.
Definition: queue.hpp:110
event ext_oneapi_wait_external_semaphore(sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, const detail::code_location &CodeLoc=detail::code_location::current())
Instruct the queue with a non-blocking wait on an external semaphore.
Definition: queue.hpp:1848
event ext_oneapi_signal_external_semaphore(sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, const detail::code_location &CodeLoc=detail::code_location::current())
Instruct the queue to signal the external semaphore once all previous commands have completed executi...
event ext_oneapi_copy(const void *Src, ext::oneapi::experimental::image_mem_handle Dest, const ext::oneapi::experimental::image_descriptor &DestImgDesc, const detail::code_location &CodeLoc=detail::code_location::current())
Copies data from one memory region to another, where Src is a USM pointer and Dest is an opaque image...
std::enable_if_t< std::is_invocable_r_v< void, T, handler & >, event > submit(T CGF, const detail::code_location &CodeLoc=detail::code_location::current())
Submits a command group function object to the queue, in order to be scheduled for execution on the d...
Definition: queue.hpp:340
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: range.hpp:26
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:131
void release_external_semaphore(interop_semaphore_handle semaphoreHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Release the external semaphore.
image_mem_handle map_external_image_memory(interop_mem_handle memHandle, const image_descriptor &desc, const sycl::device &syclDevice, const sycl::context &syclContext)
Maps an interop memory handle to an image memory handle (which may have a device optimized memory lay...
sycl::range< 3 > get_image_range(const image_mem_handle memHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Get the range that describes the image's dimensions.
unsigned int get_image_num_channels(const image_mem_handle memHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Get the number of channels that describes the image memory.
interop_mem_handle import_external_memory(external_mem_descriptor< ExternalMemHandleType > externalMem, const sycl::device &syclDevice, const sycl::context &syclContext)
Import external memory taking an external memory handle (the type of which is dependent on the OS & e...
DataT sample_image(const sampled_image_handle &imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]])
Sample data from a sampled image using its handle.
DataT sample_cubemap(const sampled_image_handle &imageHandle[[maybe_unused]], const sycl::float3 &dirVec[[maybe_unused]])
Sample a cubemap image using its handle.
unsampled_image_handle create_image(image_mem &memHandle, const image_descriptor &desc, const sycl::device &syclDevice, const sycl::context &syclContext)
Create an image and return the device image handle.
DataT fetch_image_array(const unsampled_image_handle &imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]], unsigned int arrayLayer[[maybe_unused]])
Fetch data from an unsampled image array using its handle.
void free_image_mem(image_mem_handle handle, image_type imageType, const sycl::device &syclDevice, const sycl::context &syclContext)
Free image memory.
image_mem_handle get_mip_level_mem_handle(const image_mem_handle mipMem, unsigned int level, const sycl::device &syclDevice, const sycl::context &syclContext)
Retrieve the memory handle to an individual mipmap image.
void write_cubemap(unsampled_image_handle imageHandle, const sycl::int2 &coords, int face, const DataT &color)
Write to an unsampled cubemap using its handle.
void * pitched_alloc_device(size_t *resultPitch, size_t widthInBytes, size_t height, unsigned int elementSizeBytes, const sycl::queue &syclQueue)
Allocate pitched USM image memory.
void write_image_array(unsampled_image_handle imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]], unsigned int arrayLayer[[maybe_unused]], const DataT &color[[maybe_unused]])
Write to an unsampled image array using its handle.
DataT fetch_image(const unsampled_image_handle &imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]])
Fetch data from an unsampled image using its handle.
DataT sample_mipmap(const sampled_image_handle &imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]], const float level[[maybe_unused]])
Sample a mipmap image using its handle with LOD filtering.
image_mem_handle alloc_image_mem(const image_descriptor &desc, const sycl::device &syclDevice, const sycl::context &syclContext)
Allocate image memory based on image_descriptor.
void release_external_memory(interop_mem_handle interopHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Release external memory.
DataT sample_image_array(const sampled_image_handle &imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]], unsigned int arrayLayer[[maybe_unused]])
Sample data from a sampled image array using its handle.
interop_semaphore_handle import_external_semaphore(external_semaphore_descriptor< ExternalSemaphoreHandleType > externalSemaphoreDesc, const sycl::device &syclDevice, const sycl::context &syclContext)
Import external semaphore taking an external semaphore handle (the type of which is dependent on the ...
void destroy_image_handle(unsampled_image_handle &imageHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Destroy an unsampled image handle.
DataT fetch_cubemap(const unsampled_image_handle &imageHandle, const int2 &coords, unsigned int face)
Fetch data from an unsampled cubemap image using its handle.
void write_image(unsampled_image_handle imageHandle[[maybe_unused]], const CoordT &coords[[maybe_unused]], const DataT &color[[maybe_unused]])
Write to an unsampled image using its handle.
sycl::image_channel_type get_image_channel_type(const image_mem_handle memHandle, const sycl::device &syclDevice, const sycl::context &syclContext)
Get the channel type that describes the image memory.
image_channel_type
Definition: image.hpp:74
Definition: access.hpp:18
uint64_t pi_uint64
Definition: pi.h:253
A struct to describe the properties of an image.
unsampled_image_handle(raw_image_handle_type raw_image_handle)