14 #ifndef __SYCL_DEVICE_ONLY__
38 return Coords * Range[0];
44 return {Coords.x() * Range[0], Coords.y() * Range[1]};
50 return {Coords.x() * Range[0], Coords.y() * Range[1], Coords.z() * Range[2],
61 return {
static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
67 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()), 0.5f,
74 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()),
75 static_cast<float>(Coords.z()), 0.f};
84 return Coords * ElementSize;
90 const uint8_t ElementSize) {
91 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0];
97 const uint8_t ElementSize) {
98 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0] +
99 Coords.z() * ImgPitch[1];
115 __SYCL_EXPORT
bool isOutOfRange(
const cl_int4 PixelCoord,
121 __SYCL_EXPORT cl_float4
128 template <
typename T>
133 switch (ChannelOrder) {
134 case image_channel_order::a:
137 case image_channel_order::r:
138 case image_channel_order::rx:
142 case image_channel_order::intensity:
148 case image_channel_order::luminance:
154 case image_channel_order::rg:
155 case image_channel_order::rgx:
160 case image_channel_order::ra:
164 case image_channel_order::rgb:
165 case image_channel_order::rgbx:
166 if (ChannelType == image_channel_type::unorm_short_565 ||
167 ChannelType == image_channel_type::unorm_short_555 ||
168 ChannelType == image_channel_type::unorm_int_101010) {
177 case image_channel_order::rgba:
178 case image_channel_order::ext_oneapi_srgba:
184 case image_channel_order::argb:
190 case image_channel_order::bgra:
196 case image_channel_order::abgr:
211 template <
typename T>
217 switch (ChannelOrder) {
218 case image_channel_order::a:
221 case image_channel_order::r:
222 case image_channel_order::rx:
223 case image_channel_order::intensity:
224 case image_channel_order::luminance:
227 case image_channel_order::rg:
228 case image_channel_order::rgx:
232 case image_channel_order::ra:
236 case image_channel_order::rgb:
237 case image_channel_order::rgbx:
238 if (ChannelType == image_channel_type::unorm_short_565 ||
239 ChannelType == image_channel_type::unorm_short_555 ||
240 ChannelType == image_channel_type::unorm_int_101010) {
248 case image_channel_order::rgba:
249 case image_channel_order::ext_oneapi_srgba:
255 case image_channel_order::argb:
261 case image_channel_order::bgra:
267 case image_channel_order::abgr:
283 template <
typename ChannelType>
288 switch (ImageChannelType) {
289 case image_channel_type::unsigned_int8:
290 case image_channel_type::unsigned_int16:
291 case image_channel_type::unsigned_int32:
292 RetData = PixelData.template convert<cl_uint>();
298 throw cl::sycl::invalid_parameter_error(
299 "Datatype of read data - cl_uint4 is incompatible with the "
300 "image_channel_type of the image.",
305 template <
typename ChannelType>
310 switch (ImageChannelType) {
311 case image_channel_type::signed_int8:
312 case image_channel_type::signed_int16:
313 case image_channel_type::signed_int32:
314 RetData = PixelData.template convert<cl_int>();
320 throw cl::sycl::invalid_parameter_error(
321 "Datatype of read data - cl_int4 is incompatible with the "
322 "image_channel_type of the image.",
327 template <
typename ChannelType>
332 switch (ImageChannelType) {
333 case image_channel_type::snorm_int8:
335 RetData = (PixelData.template convert<cl_float>()) / 127.0f;
338 case image_channel_type::snorm_int16:
340 RetData = (PixelData.template convert<cl_float>()) / 32767.0f;
343 case image_channel_type::unorm_int8:
345 RetData = (PixelData.template convert<cl_float>()) / 255.0f;
347 case image_channel_type::unorm_int16:
349 RetData = (PixelData.template convert<cl_float>()) / 65535.0f;
351 case image_channel_type::unorm_short_565: {
362 Temp = (Temp & MaskBits) >> ShiftBits;
363 RetData = (Temp.template convert<cl_float>()) / DivisorToNormalise;
366 case image_channel_type::unorm_short_555: {
377 Temp = (Temp & MaskBits) >> ShiftBits;
378 RetData = (Temp.template convert<cl_float>()) / 31.0f;
381 case image_channel_type::unorm_int_101010: {
387 0x000003FF , 0x00000000);
389 Temp = (Temp & MaskBits) >> ShiftBits;
390 RetData = (Temp.template convert<cl_float>()) / 1023.0f;
393 case image_channel_type::signed_int8:
394 case image_channel_type::signed_int16:
395 case image_channel_type::signed_int32:
396 case image_channel_type::unsigned_int8:
397 case image_channel_type::unsigned_int16:
398 case image_channel_type::unsigned_int32:
402 throw cl::sycl::invalid_parameter_error(
403 "Datatype of read data - cl_float4 is incompatible with the "
404 "image_channel_type of the image.",
406 case image_channel_type::fp16:
409 RetData = PixelData.template convert<cl_float>();
411 case image_channel_type::fp32:
412 RetData = PixelData.template convert<cl_float>();
417 template <
typename ChannelType>
422 switch (ImageChannelType) {
423 case image_channel_type::snorm_int8:
425 RetDataFloat = (PixelData.template convert<cl_float>()) / 127.0f;
428 case image_channel_type::snorm_int16:
430 RetDataFloat = (PixelData.template convert<cl_float>()) / 32767.0f;
433 case image_channel_type::unorm_int8:
435 RetDataFloat = (PixelData.template convert<cl_float>()) / 255.0f;
437 case image_channel_type::unorm_int16:
439 RetDataFloat = (PixelData.template convert<cl_float>()) / 65535.0f;
441 case image_channel_type::unorm_short_565:
442 case image_channel_type::unorm_short_555:
443 case image_channel_type::unorm_int_101010:
445 throw cl::sycl::feature_not_supported(
446 "Currently unsupported datatype conversion from image_channel_type "
449 case image_channel_type::signed_int8:
450 case image_channel_type::signed_int16:
451 case image_channel_type::signed_int32:
452 case image_channel_type::unsigned_int8:
453 case image_channel_type::unsigned_int16:
454 case image_channel_type::unsigned_int32:
458 throw cl::sycl::invalid_parameter_error(
459 "Datatype to read- cl_half4 is incompatible with the "
460 "image_channel_type of the image.",
462 case image_channel_type::fp16:
463 RetData = PixelData.template convert<cl_half>();
465 case image_channel_type::fp32:
466 throw cl::sycl::invalid_parameter_error(
467 "Datatype to read - cl_half4 is incompatible with the "
468 "image_channel_type of the image.",
471 RetData = RetDataFloat.template convert<cl_half>();
480 template <
typename ChannelType>
484 switch (ImageChannelType) {
485 case image_channel_type::unsigned_int8: {
487 cl_uint MinVal = min_v<cl_uchar>();
488 cl_uint MaxVal = max_v<cl_uchar>();
490 return PixelData.convert<ChannelType>();
492 case image_channel_type::unsigned_int16: {
494 cl_uint MinVal = min_v<cl_ushort>();
495 cl_uint MaxVal = max_v<cl_ushort>();
497 return PixelData.convert<ChannelType>();
499 case image_channel_type::unsigned_int32:
501 return WriteData.convert<ChannelType>();
506 throw cl::sycl::invalid_parameter_error(
507 "Datatype of data to write - cl_uint4 is incompatible with the "
508 "image_channel_type of the image.",
513 template <
typename ChannelType>
518 switch (ImageChannelType) {
519 case image_channel_type::signed_int8: {
521 cl_int MinVal = min_v<cl_char>();
522 cl_int MaxVal = max_v<cl_char>();
524 return PixelData.convert<ChannelType>();
526 case image_channel_type::signed_int16: {
528 cl_int MinVal = min_v<cl_short>();
529 cl_int MaxVal = max_v<cl_short>();
531 return PixelData.convert<ChannelType>();
533 case image_channel_type::signed_int32:
534 return WriteData.convert<ChannelType>();
539 throw cl::sycl::invalid_parameter_error(
540 "Datatype of data to write - cl_int4 is incompatible with the "
541 "image_channel_type of the image.",
546 template <
typename ChannelType>
550 vec<cl_int, 4> TempInInt = Temp.convert<int, rounding_mode::rte>();
552 cl::sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
553 return TempInIntSaturated.convert<ChannelType>();
556 template <
typename ChannelType>
563 switch (ImageChannelType) {
564 case image_channel_type::snorm_int8:
566 return processFloatDataToPixel<ChannelType>(WriteData, 127.0f);
567 case image_channel_type::snorm_int16:
569 return processFloatDataToPixel<ChannelType>(WriteData, 32767.0f);
570 case image_channel_type::unorm_int8:
572 return processFloatDataToPixel<ChannelType>(WriteData, 255.0f);
573 case image_channel_type::unorm_int16:
575 return processFloatDataToPixel<ChannelType>(WriteData, 65535.0f);
576 case image_channel_type::unorm_short_565:
578 throw cl::sycl::feature_not_supported(
579 "Currently unsupported datatype conversion from image_channel_type "
582 case image_channel_type::unorm_short_555:
589 processFloatDataToPixel<cl_ushort>(WriteData, 32.0f);
590 PixelData = cl::sycl::min(PixelData,
static_cast<ChannelType
>(0x1f));
597 (PixelData.x() << 10) | (PixelData.y() << 5) | PixelData.z();
598 return PixelData.convert<ChannelType>();
600 case image_channel_type::unorm_int_101010:
606 processFloatDataToPixel<cl_uint>(WriteData, 1023.0f);
607 PixelData = cl::sycl::min(PixelData,
static_cast<ChannelType
>(0x3ff));
609 (PixelData.x() << 20) | (PixelData.y() << 10) | PixelData.z();
610 return PixelData.convert<ChannelType>();
612 case image_channel_type::signed_int8:
613 case image_channel_type::signed_int16:
614 case image_channel_type::signed_int32:
615 case image_channel_type::unsigned_int8:
616 case image_channel_type::unsigned_int16:
617 case image_channel_type::unsigned_int32:
621 throw cl::sycl::invalid_parameter_error(
622 "Datatype of data to write - cl_float4 is incompatible with the "
623 "image_channel_type of the image.",
625 case image_channel_type::fp16:
628 return WriteData.convert<ChannelType>();
629 case image_channel_type::fp32:
630 return WriteData.convert<ChannelType>();
634 template <
typename ChannelType>
639 switch (ImageChannelType) {
640 case image_channel_type::snorm_int8:
642 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
643 case image_channel_type::snorm_int16:
645 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
646 case image_channel_type::unorm_int8:
648 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
649 case image_channel_type::unorm_int16:
651 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 65535.0f);
652 case image_channel_type::unorm_short_565:
653 case image_channel_type::unorm_short_555:
654 case image_channel_type::unorm_int_101010:
656 throw cl::sycl::feature_not_supported(
657 "Currently unsupported datatype conversion from image_channel_type "
660 case image_channel_type::signed_int8:
661 case image_channel_type::signed_int16:
662 case image_channel_type::signed_int32:
663 case image_channel_type::unsigned_int8:
664 case image_channel_type::unsigned_int16:
665 case image_channel_type::unsigned_int32:
669 throw cl::sycl::invalid_parameter_error(
670 "Datatype of data to write - cl_float4 is incompatible with the "
671 "image_channel_type of the image.",
673 case image_channel_type::fp16:
674 return WriteData.convert<ChannelType>();
675 case image_channel_type::fp32:
676 throw cl::sycl::invalid_parameter_error(
677 "Datatype of data to write - cl_float4 is incompatible with the "
678 "image_channel_type of the image.",
697 template <
typename CoordT,
typename WriteDataT>
699 id<3> ImgPitch, uint8_t ElementSize,
703 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
706 switch (ImgChannelType) {
707 case image_channel_type::snorm_int8:
708 writePixel(convertWriteData<cl_char>(Color, ImgChannelType),
709 reinterpret_cast<cl_char *
>(Ptr), ImgChannelOrder,
712 case image_channel_type::snorm_int16:
713 writePixel(convertWriteData<cl_short>(Color, ImgChannelType),
714 reinterpret_cast<cl_short *
>(Ptr), ImgChannelOrder,
717 case image_channel_type::unorm_int8:
718 writePixel(convertWriteData<cl_uchar>(Color, ImgChannelType),
719 reinterpret_cast<cl_uchar *
>(Ptr), ImgChannelOrder,
722 case image_channel_type::unorm_int16:
723 writePixel(convertWriteData<cl_ushort>(Color, ImgChannelType),
724 reinterpret_cast<cl_ushort *
>(Ptr), ImgChannelOrder,
727 case image_channel_type::unorm_short_565:
728 writePixel(convertWriteData<short>(Color, ImgChannelType),
729 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
731 case image_channel_type::unorm_short_555:
732 writePixel(convertWriteData<short>(Color, ImgChannelType),
733 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
735 case image_channel_type::unorm_int_101010:
736 writePixel(convertWriteData<cl_uint>(Color, ImgChannelType),
737 reinterpret_cast<cl_uint *
>(Ptr), ImgChannelOrder,
740 case image_channel_type::signed_int8:
741 writePixel(convertWriteData<cl_char>(Color, ImgChannelType),
742 reinterpret_cast<cl_char *
>(Ptr), ImgChannelOrder,
745 case image_channel_type::signed_int16:
746 writePixel(convertWriteData<cl_short>(Color, ImgChannelType),
747 reinterpret_cast<cl_short *
>(Ptr), ImgChannelOrder,
750 case image_channel_type::signed_int32:
751 writePixel(convertWriteData<cl_int>(Color, ImgChannelType),
752 reinterpret_cast<cl_int *
>(Ptr), ImgChannelOrder,
755 case image_channel_type::unsigned_int8:
756 writePixel(convertWriteData<cl_uchar>(Color, ImgChannelType),
757 reinterpret_cast<cl_uchar *
>(Ptr), ImgChannelOrder,
760 case image_channel_type::unsigned_int16:
761 writePixel(convertWriteData<cl_ushort>(Color, ImgChannelType),
762 reinterpret_cast<cl_ushort *
>(Ptr), ImgChannelOrder,
765 case image_channel_type::unsigned_int32:
766 writePixel(convertWriteData<cl_uint>(Color, ImgChannelType),
767 reinterpret_cast<cl_uint *
>(Ptr), ImgChannelOrder,
770 case image_channel_type::fp16:
774 convertWriteData<cl_half>(Color, ImgChannelType),
775 reinterpret_cast<cl_half *
>(Ptr), ImgChannelOrder, ImgChannelType);
777 case image_channel_type::fp32:
778 writePixel(convertWriteData<cl_float>(Color, ImgChannelType),
779 reinterpret_cast<cl_float *
>(Ptr), ImgChannelOrder,
799 template <
typename DataT>
803 void *BasePtr,
const uint8_t ElementSize) {
805 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
810 switch (ImageChannelType) {
813 case image_channel_type::snorm_int8:
815 ImageChannelOrder, ImageChannelType),
816 image_channel_type::snorm_int8, Color);
818 case image_channel_type::snorm_int16:
820 ImageChannelOrder, ImageChannelType),
821 image_channel_type::snorm_int16, Color);
823 case image_channel_type::unorm_int8:
825 ImageChannelOrder, ImageChannelType),
826 image_channel_type::unorm_int8, Color);
828 case image_channel_type::unorm_int16:
830 ImageChannelOrder, ImageChannelType),
831 image_channel_type::unorm_int16, Color);
833 case image_channel_type::unorm_short_565:
835 ImageChannelOrder, ImageChannelType),
836 image_channel_type::unorm_short_565, Color);
838 case image_channel_type::unorm_short_555:
840 ImageChannelOrder, ImageChannelType),
841 image_channel_type::unorm_short_555, Color);
843 case image_channel_type::unorm_int_101010:
845 ImageChannelOrder, ImageChannelType),
846 image_channel_type::unorm_int_101010, Color);
848 case image_channel_type::signed_int8:
850 ImageChannelOrder, ImageChannelType),
851 image_channel_type::signed_int8, Color);
853 case image_channel_type::signed_int16:
855 ImageChannelOrder, ImageChannelType),
856 image_channel_type::signed_int16, Color);
858 case image_channel_type::signed_int32:
860 ImageChannelOrder, ImageChannelType),
861 image_channel_type::signed_int32, Color);
863 case image_channel_type::unsigned_int8:
865 ImageChannelOrder, ImageChannelType),
866 image_channel_type::unsigned_int8, Color);
868 case image_channel_type::unsigned_int16:
870 ImageChannelOrder, ImageChannelType),
871 image_channel_type::unsigned_int16, Color);
873 case image_channel_type::unsigned_int32:
875 ImageChannelOrder, ImageChannelType),
876 image_channel_type::unsigned_int32, Color);
878 case image_channel_type::fp16:
880 ImageChannelOrder, ImageChannelType),
881 image_channel_type::fp16, Color);
883 case image_channel_type::fp32:
885 ImageChannelOrder, ImageChannelType),
886 image_channel_type::fp32, Color);
895 template <
typename DataT>
900 const uint8_t ElementSize) {
906 RetData = ReadPixelData<DataT>(PixelCoord, ImgPitch, ImgChannelType,
907 ImgChannelOrder, BasePtr, ElementSize);
917 template <
typename DataT>
924 void *BasePtr,
const uint8_t ElementSize) {
925 cl_int i0 = CoordValues.s0(), j0 = CoordValues.s1(), k0 = CoordValues.s2(),
926 i1 = CoordValues.s4(), j1 = CoordValues.s5(), k1 = CoordValues.s6();
928 auto getColorInFloat =
930 DataT Res = getColor<DataT>(V, SmplAddrMode,
931 ImgRange, ImgPitch, ImgChannelType,
932 ImgChannelOrder, BasePtr, ElementSize);
933 return Res.template convert<cl_float>();
937 cl_float4 Ci0j0k0 = getColorInFloat(cl_int4{i0, j0, k0, 0});
939 cl_float4 Ci1j0k0 = getColorInFloat(cl_int4{i1, j0, k0, 0});
941 cl_float4 Ci0j1k0 = getColorInFloat(cl_int4{i0, j1, k0, 0});
943 cl_float4 Ci1j1k0 = getColorInFloat(cl_int4{i1, j1, k0, 0});
945 cl_float4 Ci0j0k1 = getColorInFloat(cl_int4{i0, j0, k1, 0});
947 cl_float4 Ci1j0k1 = getColorInFloat(cl_int4{i1, j0, k1, 0});
949 cl_float4 Ci0j1k1 = getColorInFloat(cl_int4{i0, j1, k1, 0});
951 cl_float4 Ci1j1k1 = getColorInFloat(cl_int4{i1, j1, k1, 0});
957 Ci0j0k0 = (1 -
a) * (1 - b) * (1 - c) * Ci0j0k0;
958 Ci1j0k0 =
a * (1 - b) * (1 - c) * Ci1j0k0;
959 Ci0j1k0 = (1 -
a) * b * (1 - c) * Ci0j1k0;
960 Ci1j1k0 =
a * b * (1 - c) * Ci1j1k0;
961 Ci0j0k1 = (1 -
a) * (1 - b) * c * Ci0j0k1;
962 Ci1j0k1 =
a * (1 - b) * c * Ci1j0k1;
963 Ci0j1k1 = (1 -
a) * b * c * Ci0j1k1;
964 Ci1j1k1 =
a * b * c * Ci1j1k1;
966 cl_float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 +
967 Ci1j0k1 + Ci0j1k1 + Ci1j1k1;
1017 template <
typename CoordT,
typename DataT>
1023 void *BasePtr, uint8_t ElementSize) {
1031 cl_float4 FloatCoorduvw;
1037 switch (SmplNormMode) {
1038 case coordinate_normalization_mode::unnormalized:
1039 switch (SmplAddrMode) {
1040 case addressing_mode::mirrored_repeat:
1041 case addressing_mode::repeat:
1042 throw cl::sycl::feature_not_supported(
1043 "Sampler used with unsupported configuration of "
1044 "mirrored_repeat/repeat filtering mode with unnormalized "
1047 case addressing_mode::clamp_to_edge:
1049 case addressing_mode::none:
1055 case coordinate_normalization_mode::normalized:
1056 switch (SmplAddrMode) {
1057 case addressing_mode::mirrored_repeat:
1058 case addressing_mode::repeat:
1064 case addressing_mode::clamp_to_edge:
1066 case addressing_mode::none:
1085 switch (SmplFiltMode) {
1086 case filtering_mode::nearest: {
1088 cl_int4 PixelCoord =
1095 getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1096 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1099 case filtering_mode::linear: {
1105 FloatCoorduvw, SmplAddrMode, ImgRange, Retabc);
1109 RetData = ReadPixelDataLinearFiltMode<DataT>(
1110 CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1111 ImgChannelOrder, BasePtr, ElementSize);