14 #ifndef __SYCL_DEVICE_ONLY__
32 #include <type_traits>
35 inline namespace _V1 {
41 std::int32_t,
float>>>::type;
46 std::enable_if_t<IsValidCoordType<T>::value, T>
48 return Coords * Range[0];
52 std::enable_if_t<IsValidCoordType<T>::value,
vec<T, 2>>
54 return {Coords.x() * Range[0], Coords.y() * Range[1]};
58 std::enable_if_t<IsValidCoordType<T>::value,
vec<T, 4>>
60 return {Coords.x() * Range[0], Coords.y() * Range[1], Coords.z() * Range[2],
70 return {
static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
74 std::enable_if_t<IsValidCoordType<T>::value, float4>
76 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()), 0.5f,
81 std::enable_if_t<IsValidCoordType<T>::value, float4>
83 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()),
84 static_cast<float>(Coords.z()), 0.f};
91 std::enable_if_t<std::is_integral_v<T>,
size_t>
93 return Coords * ElementSize;
97 std::enable_if_t<std::is_integral_v<T>,
size_t>
99 const uint8_t ElementSize) {
100 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0];
103 template <
typename T>
104 std::enable_if_t<std::is_integral_v<T>,
size_t>
106 const uint8_t ElementSize) {
107 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0] +
108 Coords.z() * ImgPitch[1];
134 template <
typename T>
139 switch (ChannelOrder) {
217 template <
typename T>
223 switch (ChannelOrder) {
289 template <
typename ChannelType>
294 switch (ImageChannelType) {
298 RetData = PixelData.template convert<std::uint32_t>();
304 throw sycl::invalid_parameter_error(
305 "Datatype of read data - cl_uint4 is incompatible with the "
306 "image_channel_type of the image.",
307 PI_ERROR_INVALID_VALUE);
311 template <
typename ChannelType>
315 switch (ImageChannelType) {
319 RetData = PixelData.template convert<std::int32_t>();
325 throw sycl::invalid_parameter_error(
326 "Datatype of read data - cl_int4 is incompatible with "
328 "image_channel_type of the image.",
329 PI_ERROR_INVALID_VALUE);
333 template <
typename ChannelType>
338 switch (ImageChannelType) {
341 RetData = (PixelData.template convert<float>()) / 127.0f;
342 RetData = sycl::fmax(RetData, -1);
346 RetData = (PixelData.template convert<float>()) / 32767.0f;
347 RetData = sycl::fmax(RetData, -1);
351 RetData = (PixelData.template convert<float>()) / 255.0f;
355 RetData = (PixelData.template convert<float>()) / 65535.0f;
363 ushort4 Temp(PixelData.x());
364 ushort4 MaskBits(0xF800 , 0x07E0 ,
366 ushort4 ShiftBits(11, 5, 0, 0);
367 float4 DivisorToNormalise(31.0f, 63.0f, 31.0f, 1);
368 Temp = (Temp & MaskBits) >> ShiftBits;
369 RetData = (Temp.template convert<float>()) / DivisorToNormalise;
379 ushort4 Temp(PixelData.x());
380 ushort4 MaskBits(0x7C00 , 0x03E0 ,
382 ushort4 ShiftBits(10, 5, 0, 0);
383 Temp = (Temp & MaskBits) >> ShiftBits;
384 RetData = (Temp.template convert<float>()) / 31.0f;
390 uint4 Temp(PixelData.x());
391 uint4 MaskBits(0x3FF00000 , 0x000FFC00 ,
392 0x000003FF , 0x00000000);
393 uint4 ShiftBits(20, 10, 0, 0);
394 Temp = (Temp & MaskBits) >> ShiftBits;
395 RetData = (Temp.template convert<float>()) / 1023.0f;
407 throw sycl::invalid_parameter_error(
408 "Datatype of read data - cl_float4 is incompatible with the "
409 "image_channel_type of the image.",
410 PI_ERROR_INVALID_VALUE);
414 RetData = PixelData.template convert<float>();
417 RetData = PixelData.template convert<float>();
422 template <
typename ChannelType>
427 switch (ImageChannelType) {
430 RetDataFloat = (PixelData.template convert<float>()) / 127.0f;
431 RetDataFloat = sycl::fmax(RetDataFloat, -1);
435 RetDataFloat = (PixelData.template convert<float>()) / 32767.0f;
436 RetDataFloat = sycl::fmax(RetDataFloat, -1);
440 RetDataFloat = (PixelData.template convert<float>()) / 255.0f;
444 RetDataFloat = (PixelData.template convert<float>()) / 65535.0f;
450 throw sycl::feature_not_supported(
451 "Currently unsupported datatype conversion from image_channel_type "
453 PI_ERROR_INVALID_OPERATION);
463 throw sycl::invalid_parameter_error(
464 "Datatype to read- cl_half4 is incompatible with the "
465 "image_channel_type of the image.",
466 PI_ERROR_INVALID_VALUE);
468 RetData = PixelData.template convert<half>();
471 throw sycl::invalid_parameter_error(
472 "Datatype to read - cl_half4 is incompatible with the "
473 "image_channel_type of the image.",
474 PI_ERROR_INVALID_VALUE);
476 RetData = RetDataFloat.template convert<half>();
485 template <
typename ChannelType>
489 switch (ImageChannelType) {
492 std::uint32_t MinVal = min_v<std::uint8_t>();
493 std::uint32_t MaxVal = max_v<std::uint8_t>();
494 uint4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
495 return PixelData.convert<ChannelType>();
499 std::uint32_t MinVal = min_v<std::uint16_t>();
500 std::uint32_t MaxVal = max_v<std::uint16_t>();
501 uint4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
502 return PixelData.convert<ChannelType>();
506 return WriteData.convert<ChannelType>();
511 throw sycl::invalid_parameter_error(
512 "Datatype of data to write - cl_uint4 is incompatible with the "
513 "image_channel_type of the image.",
514 PI_ERROR_INVALID_VALUE);
518 template <
typename ChannelType>
523 switch (ImageChannelType) {
526 std::int32_t MinVal = min_v<std::int8_t>();
527 std::int32_t MaxVal = max_v<std::int8_t>();
528 int4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
529 return PixelData.convert<ChannelType>();
533 std::int32_t MinVal = min_v<std::int16_t>();
534 std::int32_t MaxVal = max_v<std::int16_t>();
535 int4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
536 return PixelData.convert<ChannelType>();
539 return WriteData.convert<ChannelType>();
544 throw sycl::invalid_parameter_error(
545 "Datatype of data to write - cl_int4 is incompatible with the "
546 "image_channel_type of the image.",
547 PI_ERROR_INVALID_VALUE);
551 template <
typename ChannelType>
553 float4 Temp = WriteData * MulFactor;
555 int4 TempInIntSaturated =
556 sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
557 return TempInIntSaturated.convert<ChannelType>();
560 template <
typename ChannelType>
567 switch (ImageChannelType) {
570 return processFloatDataToPixel<ChannelType>(WriteData, 127.0f);
573 return processFloatDataToPixel<ChannelType>(WriteData, 32767.0f);
576 return processFloatDataToPixel<ChannelType>(WriteData, 255.0f);
579 return processFloatDataToPixel<ChannelType>(WriteData, 65535.0f);
582 throw sycl::feature_not_supported(
583 "Currently unsupported datatype conversion from image_channel_type "
585 PI_ERROR_INVALID_OPERATION);
593 processFloatDataToPixel<std::uint16_t>(WriteData, 32.0f);
594 PixelData =
sycl::min(PixelData,
static_cast<ChannelType
>(0x1f));
601 (PixelData.x() << 10) | (PixelData.y() << 5) | PixelData.z();
602 return PixelData.convert<ChannelType>();
610 processFloatDataToPixel<std::uint32_t>(WriteData, 1023.0f);
611 PixelData =
sycl::min(PixelData,
static_cast<ChannelType
>(0x3ff));
613 (PixelData.x() << 20) | (PixelData.y() << 10) | PixelData.z();
614 return PixelData.convert<ChannelType>();
625 throw sycl::invalid_parameter_error(
626 "Datatype of data to write - cl_float4 is incompatible with the "
627 "image_channel_type of the image.",
628 PI_ERROR_INVALID_VALUE);
632 return WriteData.convert<ChannelType>();
634 return WriteData.convert<ChannelType>();
638 template <
typename ChannelType>
642 float4 WriteDataFloat = WriteData.convert<
float>();
643 switch (ImageChannelType) {
646 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
649 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
652 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
655 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 65535.0f);
660 throw sycl::feature_not_supported(
661 "Currently unsupported datatype conversion from image_channel_type "
663 PI_ERROR_INVALID_OPERATION);
673 throw sycl::invalid_parameter_error(
674 "Datatype of data to write - cl_float4 is incompatible with the "
675 "image_channel_type of the image.",
676 PI_ERROR_INVALID_VALUE);
678 return WriteData.convert<ChannelType>();
680 throw sycl::invalid_parameter_error(
681 "Datatype of data to write - cl_float4 is incompatible with the "
682 "image_channel_type of the image.",
683 PI_ERROR_INVALID_VALUE);
701 template <
typename CoordT,
typename WriteDataT>
703 id<3> ImgPitch, uint8_t ElementSize,
707 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
710 switch (ImgChannelType) {
712 writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
713 reinterpret_cast<std::int8_t *
>(Ptr), ImgChannelOrder,
717 writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
718 reinterpret_cast<std::int16_t *
>(Ptr), ImgChannelOrder,
722 writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
723 reinterpret_cast<std::uint8_t *
>(Ptr), ImgChannelOrder,
727 writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
728 reinterpret_cast<std::uint16_t *
>(Ptr), ImgChannelOrder,
732 writePixel(convertWriteData<short>(Color, ImgChannelType),
733 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
736 writePixel(convertWriteData<short>(Color, ImgChannelType),
737 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
740 writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
741 reinterpret_cast<std::uint32_t *
>(Ptr), ImgChannelOrder,
745 writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
746 reinterpret_cast<std::int8_t *
>(Ptr), ImgChannelOrder,
750 writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
751 reinterpret_cast<std::int16_t *
>(Ptr), ImgChannelOrder,
755 writePixel(convertWriteData<std::int32_t>(Color, ImgChannelType),
756 reinterpret_cast<std::int32_t *
>(Ptr), ImgChannelOrder,
760 writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
761 reinterpret_cast<std::uint8_t *
>(Ptr), ImgChannelOrder,
765 writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
766 reinterpret_cast<std::uint16_t *
>(Ptr), ImgChannelOrder,
770 writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
771 reinterpret_cast<std::uint32_t *
>(Ptr), ImgChannelOrder,
776 convertWriteData<half>(Color, ImgChannelType),
777 reinterpret_cast<half *
>(Ptr), ImgChannelOrder, ImgChannelType);
780 writePixel(convertWriteData<float>(Color, ImgChannelType),
781 reinterpret_cast<float *
>(Ptr), ImgChannelOrder, ImgChannelType);
800 template <
typename DataT>
804 const uint8_t ElementSize) {
806 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
811 switch (ImageChannelType) {
815 convertReadData<std::int8_t>(
readPixel(
reinterpret_cast<std::int8_t *
>(Ptr),
816 ImageChannelOrder, ImageChannelType),
820 convertReadData<std::int16_t>(
821 readPixel(
reinterpret_cast<std::int16_t *
>(Ptr), ImageChannelOrder,
826 convertReadData<std::uint8_t>(
827 readPixel(
reinterpret_cast<std::uint8_t *
>(Ptr), ImageChannelOrder,
832 convertReadData<std::uint16_t>(
833 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
838 convertReadData<std::uint16_t>(
839 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
844 convertReadData<std::uint16_t>(
845 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
850 convertReadData<std::uint32_t>(
851 readPixel(
reinterpret_cast<std::uint32_t *
>(Ptr), ImageChannelOrder,
856 convertReadData<std::int8_t>(
readPixel(
reinterpret_cast<std::int8_t *
>(Ptr),
857 ImageChannelOrder, ImageChannelType),
861 convertReadData<std::int16_t>(
862 readPixel(
reinterpret_cast<std::int16_t *
>(Ptr), ImageChannelOrder,
867 convertReadData<std::int32_t>(
868 readPixel(
reinterpret_cast<std::int32_t *
>(Ptr), ImageChannelOrder,
873 convertReadData<std::uint8_t>(
874 readPixel(
reinterpret_cast<std::uint8_t *
>(Ptr), ImageChannelOrder,
879 convertReadData<std::uint16_t>(
880 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
885 convertReadData<std::uint32_t>(
886 readPixel(
reinterpret_cast<std::uint32_t *
>(Ptr), ImageChannelOrder,
891 convertReadData<half>(
readPixel(
reinterpret_cast<half *
>(Ptr),
892 ImageChannelOrder, ImageChannelType),
896 convertReadData<float>(
readPixel(
reinterpret_cast<float *
>(Ptr),
897 ImageChannelOrder, ImageChannelType),
907 template <
typename DataT>
912 const uint8_t ElementSize) {
918 RetData = ReadPixelData<DataT>(PixelCoord, ImgPitch, ImgChannelType,
919 ImgChannelOrder, BasePtr, ElementSize);
929 template <
typename DataT>
935 void *BasePtr,
const uint8_t ElementSize) {
936 std::int32_t i0 = CoordValues.s0(), j0 = CoordValues.s1(),
937 k0 = CoordValues.s2(), i1 = CoordValues.s4(),
938 j1 = CoordValues.s5(), k1 = CoordValues.s6();
940 auto getColorInFloat = [&](int4 V) {
942 getColor<DataT>(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
943 ImgChannelOrder, BasePtr, ElementSize);
944 return Res.template convert<float>();
948 float4 Ci0j0k0 = getColorInFloat(int4{i0, j0, k0, 0});
950 float4 Ci1j0k0 = getColorInFloat(int4{i1, j0, k0, 0});
952 float4 Ci0j1k0 = getColorInFloat(int4{i0, j1, k0, 0});
954 float4 Ci1j1k0 = getColorInFloat(int4{i1, j1, k0, 0});
956 float4 Ci0j0k1 = getColorInFloat(int4{i0, j0, k1, 0});
958 float4 Ci1j0k1 = getColorInFloat(int4{i1, j0, k1, 0});
960 float4 Ci0j1k1 = getColorInFloat(int4{i0, j1, k1, 0});
962 float4 Ci1j1k1 = getColorInFloat(int4{i1, j1, k1, 0});
968 Ci0j0k0 = (1 -
a) * (1 -
b) * (1 - c) * Ci0j0k0;
969 Ci1j0k0 =
a * (1 -
b) * (1 - c) * Ci1j0k0;
970 Ci0j1k0 = (1 -
a) *
b * (1 - c) * Ci0j1k0;
971 Ci1j1k0 =
a *
b * (1 - c) * Ci1j1k0;
972 Ci0j0k1 = (1 -
a) * (1 -
b) * c * Ci0j0k1;
973 Ci1j0k1 =
a * (1 -
b) * c * Ci1j0k1;
974 Ci0j1k1 = (1 -
a) *
b * c * Ci0j1k1;
975 Ci1j1k1 =
a *
b * c * Ci1j1k1;
977 float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 + Ci1j0k1 +
1028 template <
typename CoordT,
typename DataT>
1034 void *BasePtr, uint8_t ElementSize) {
1037 float4 FloatCoorduvw;
1043 switch (SmplNormMode) {
1045 switch (SmplAddrMode) {
1048 throw sycl::feature_not_supported(
1049 "Sampler used with unsupported configuration of "
1050 "mirrored_repeat/repeat filtering mode with unnormalized "
1052 PI_ERROR_INVALID_OPERATION);
1062 switch (SmplAddrMode) {
1091 switch (SmplFiltMode) {
1101 getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1102 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1115 RetData = ReadPixelDataLinearFiltMode<DataT>(
1116 CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1117 ImgChannelOrder, BasePtr, ElementSize);
1127 template <
typename CoordT,
typename DataT>
1133 void *BasePtr, uint8_t ElementSize) {
1136 Smpl.get_coordinate_normalization_mode();
1140 return imageReadSamplerHostImpl<CoordT, DataT>(
1141 Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1142 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1146 template <
typename CoordT,
typename DataT>
1152 void *BasePtr, uint8_t ElementSize) {
1157 return imageReadSamplerHostImpl<CoordT, DataT>(
1158 Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1159 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
float4 getBorderColor(const image_channel_order ImgChannelOrder)
int4 getPixelCoordNearestFiltMode(float4, const addressing_mode, const range< 3 >)
DataT imageReadSamplerHostImpl(const CoordT &Coords, coordinate_normalization_mode SmplNormMode, addressing_mode SmplAddrMode, filtering_mode SmplFiltMode, range< 3 > ImgRange, id< 3 > ImgPitch, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr, uint8_t ElementSize)
bool isOutOfRange(const int4 PixelCoord, const addressing_mode SmplAddrMode, const range< 3 > ImgRange)
DataT getColor(const int4 PixelCoord, const addressing_mode SmplAddrMode, const range< 3 > ImgRange, const id< 3 > ImgPitch, const image_channel_type ImgChannelType, const image_channel_order ImgChannelOrder, void *BasePtr, const uint8_t ElementSize)
void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color, id< 3 > ImgPitch, uint8_t ElementSize, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr)
void writePixel(const vec< T, 4 > Pixel, T *Ptr, const image_channel_order ChannelOrder, const image_channel_type ChannelType)
boost::mp11::mp_set_contains< TypeList, std::remove_cv_t< T > > is_contained
typename get_elem_type< T >::type get_elem_type_t
int8 getPixelCoordLinearFiltMode(float4, const addressing_mode, const range< 3 >, float4 &)
vec< T, 4 > readPixel(T *Ptr, const image_channel_order ChannelOrder, const image_channel_type ChannelType)
boost::mp11::mp_list< T... > type_list
std::enable_if_t< std::is_integral_v< T >, size_t > getImageOffset(const T &Coords, const id< 3 >, const uint8_t ElementSize)
vec< ChannelType, 4 > processFloatDataToPixel(float4 WriteData, float MulFactor)
DataT ReadPixelData(const int4 PixelCoord, const id< 3 > ImgPitch, const image_channel_type ImageChannelType, const image_channel_order ImageChannelOrder, void *BasePtr, const uint8_t ElementSize)
typename is_contained< T, boost::mp11::mp_unique< type_list< opencl::cl_int, opencl::cl_float, std::int32_t, float > >>::type IsValidCoordType
void convertReadData(const vec< ChannelType, 4 > PixelData, const image_channel_type ImageChannelType, uint4 &RetData)
DataT ReadPixelDataLinearFiltMode(const int8 CoordValues, const float4 abc, const addressing_mode SmplAddrMode, const range< 3 > ImgRange, id< 3 > ImgPitch, const image_channel_type ImgChannelType, const image_channel_order ImgChannelOrder, void *BasePtr, const uint8_t ElementSize)
std::enable_if_t< IsValidCoordType< T >::value, float4 > convertToFloat4(T Coords)
std::enable_if_t< IsValidCoordType< T >::value, T > UnnormalizeCoordinates(const T &Coords, const range< 3 > &Range)
vec< ChannelType, 4 > convertWriteData(const uint4 WriteData, const image_channel_type ImageChannelType)
coordinate_normalization_mode
addressing_mode addressing
coordinate_normalization_mode coordinate