14 #ifndef __SYCL_DEVICE_ONLY__
31 #include <type_traits>
34 inline namespace _V1 {
40 std::int32_t,
float>>>::type;
45 std::enable_if_t<IsValidCoordType<T>::value, T>
47 return Coords * Range[0];
51 std::enable_if_t<IsValidCoordType<T>::value,
vec<T, 2>>
53 return {Coords.x() * Range[0], Coords.y() * Range[1]};
57 std::enable_if_t<IsValidCoordType<T>::value,
vec<T, 4>>
59 return {Coords.x() * Range[0], Coords.y() * Range[1], Coords.z() * Range[2],
69 return {
static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
73 std::enable_if_t<IsValidCoordType<T>::value, float4>
75 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()), 0.5f,
80 std::enable_if_t<IsValidCoordType<T>::value, float4>
82 return {
static_cast<float>(Coords.x()),
static_cast<float>(Coords.y()),
83 static_cast<float>(Coords.z()), 0.f};
90 std::enable_if_t<std::is_integral_v<T>,
size_t>
92 return Coords * ElementSize;
96 std::enable_if_t<std::is_integral_v<T>,
size_t>
98 const uint8_t ElementSize) {
99 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0];
102 template <
typename T>
103 std::enable_if_t<std::is_integral_v<T>,
size_t>
105 const uint8_t ElementSize) {
106 return Coords.x() * ElementSize + Coords.y() * ImgPitch[0] +
107 Coords.z() * ImgPitch[1];
133 template <
typename T>
138 switch (ChannelOrder) {
216 template <
typename T>
222 switch (ChannelOrder) {
288 template <
typename ChannelType>
293 switch (ImageChannelType) {
297 RetData = PixelData.template convert<std::uint32_t>();
304 "Datatype of read data - cl_uint4 is incompatible "
305 "with the image_channel_type of the image.");
309 template <
typename ChannelType>
313 switch (ImageChannelType) {
317 RetData = PixelData.template convert<std::int32_t>();
324 "Datatype of read data - cl_int4 is incompatible "
325 "with the image_channel_type of the image.");
329 template <
typename ChannelType>
334 switch (ImageChannelType) {
337 RetData = (PixelData.template convert<float>()) / 127.0f;
338 RetData = sycl::fmax(RetData, -1);
342 RetData = (PixelData.template convert<float>()) / 32767.0f;
343 RetData = sycl::fmax(RetData, -1);
347 RetData = (PixelData.template convert<float>()) / 255.0f;
351 RetData = (PixelData.template convert<float>()) / 65535.0f;
359 ushort4 Temp(PixelData.x());
360 ushort4 MaskBits(0xF800 , 0x07E0 ,
362 ushort4 ShiftBits(11, 5, 0, 0);
363 float4 DivisorToNormalise(31.0f, 63.0f, 31.0f, 1);
364 Temp = (Temp & MaskBits) >> ShiftBits;
365 RetData = (Temp.template convert<float>()) / DivisorToNormalise;
375 ushort4 Temp(PixelData.x());
376 ushort4 MaskBits(0x7C00 , 0x03E0 ,
378 ushort4 ShiftBits(10, 5, 0, 0);
379 Temp = (Temp & MaskBits) >> ShiftBits;
380 RetData = (Temp.template convert<float>()) / 31.0f;
386 uint4 Temp(PixelData.x());
387 uint4 MaskBits(0x3FF00000 , 0x000FFC00 ,
388 0x000003FF , 0x00000000);
389 uint4 ShiftBits(20, 10, 0, 0);
390 Temp = (Temp & MaskBits) >> ShiftBits;
391 RetData = (Temp.template convert<float>()) / 1023.0f;
404 "Datatype of read data - cl_float4 is incompatible "
405 "with the image_channel_type of the image.");
409 RetData = PixelData.template convert<float>();
412 RetData = PixelData.template convert<float>();
417 template <
typename ChannelType>
422 switch (ImageChannelType) {
425 RetDataFloat = (PixelData.template convert<float>()) / 127.0f;
426 RetDataFloat = sycl::fmax(RetDataFloat, -1);
430 RetDataFloat = (PixelData.template convert<float>()) / 32767.0f;
431 RetDataFloat = sycl::fmax(RetDataFloat, -1);
435 RetDataFloat = (PixelData.template convert<float>()) / 255.0f;
439 RetDataFloat = (PixelData.template convert<float>()) / 65535.0f;
446 "Currently unsupported datatype conversion from "
447 "image_channel_type to cl_half4.");
458 "Datatype to read- cl_half4 is incompatible with the "
459 "image_channel_type of the image.");
461 RetData = PixelData.template convert<half>();
465 "Datatype to read - cl_half4 is incompatible with "
466 "the image_channel_type of the image.");
468 RetData = RetDataFloat.template convert<half>();
477 template <
typename ChannelType>
481 switch (ImageChannelType) {
484 std::uint32_t MinVal = min_v<std::uint8_t>();
485 std::uint32_t MaxVal = max_v<std::uint8_t>();
486 uint4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
487 return PixelData.convert<ChannelType>();
491 std::uint32_t MinVal = min_v<std::uint16_t>();
492 std::uint32_t MaxVal = max_v<std::uint16_t>();
493 uint4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
494 return PixelData.convert<ChannelType>();
498 return WriteData.convert<ChannelType>();
505 "Datatype of data to write - cl_uint4 is incompatible with the "
506 "image_channel_type of the image.");
510 template <
typename ChannelType>
515 switch (ImageChannelType) {
518 std::int32_t MinVal = min_v<std::int8_t>();
519 std::int32_t MaxVal = max_v<std::int8_t>();
520 int4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
521 return PixelData.convert<ChannelType>();
525 std::int32_t MinVal = min_v<std::int16_t>();
526 std::int32_t MaxVal = max_v<std::int16_t>();
527 int4 PixelData =
sycl::clamp(WriteData, MinVal, MaxVal);
528 return PixelData.convert<ChannelType>();
531 return WriteData.convert<ChannelType>();
537 "Datatype of data to write - cl_int4 is incompatible "
538 "with the image_channel_type of the image.");
542 template <
typename ChannelType>
544 float4 Temp = WriteData * MulFactor;
546 int4 TempInIntSaturated =
547 sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
548 return TempInIntSaturated.convert<ChannelType>();
551 template <
typename ChannelType>
558 switch (ImageChannelType) {
561 return processFloatDataToPixel<ChannelType>(WriteData, 127.0f);
564 return processFloatDataToPixel<ChannelType>(WriteData, 32767.0f);
567 return processFloatDataToPixel<ChannelType>(WriteData, 255.0f);
570 return processFloatDataToPixel<ChannelType>(WriteData, 65535.0f);
574 "Currently unsupported datatype conversion from "
575 "image_channel_type to cl_float4.");
583 processFloatDataToPixel<std::uint16_t>(WriteData, 32.0f);
584 PixelData =
sycl::min(PixelData,
static_cast<ChannelType
>(0x1f));
591 (PixelData.x() << 10) | (PixelData.y() << 5) | PixelData.z();
592 return PixelData.convert<ChannelType>();
600 processFloatDataToPixel<std::uint32_t>(WriteData, 1023.0f);
601 PixelData =
sycl::min(PixelData,
static_cast<ChannelType
>(0x3ff));
603 (PixelData.x() << 20) | (PixelData.y() << 10) | PixelData.z();
604 return PixelData.convert<ChannelType>();
617 "Datatype of data to write - cl_float4 is incompatible with the "
618 "image_channel_type of the image.");
622 return WriteData.convert<ChannelType>();
624 return WriteData.convert<ChannelType>();
628 template <
typename ChannelType>
632 float4 WriteDataFloat = WriteData.convert<
float>();
633 switch (ImageChannelType) {
636 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
639 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
642 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
645 return processFloatDataToPixel<ChannelType>(WriteDataFloat, 65535.0f);
651 "Currently unsupported datatype conversion from "
652 "image_channel_type to cl_half4.");
664 "Datatype of data to write - cl_float4 is incompatible with the "
665 "image_channel_type of the image.");
667 return WriteData.convert<ChannelType>();
671 "Datatype of data to write - cl_float4 is incompatible with the "
672 "image_channel_type of the image.");
690 template <
typename CoordT,
typename WriteDataT>
692 id<3> ImgPitch, uint8_t ElementSize,
696 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
699 switch (ImgChannelType) {
701 writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
702 reinterpret_cast<std::int8_t *
>(Ptr), ImgChannelOrder,
706 writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
707 reinterpret_cast<std::int16_t *
>(Ptr), ImgChannelOrder,
711 writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
712 reinterpret_cast<std::uint8_t *
>(Ptr), ImgChannelOrder,
716 writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
717 reinterpret_cast<std::uint16_t *
>(Ptr), ImgChannelOrder,
721 writePixel(convertWriteData<short>(Color, ImgChannelType),
722 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
725 writePixel(convertWriteData<short>(Color, ImgChannelType),
726 reinterpret_cast<short *
>(Ptr), ImgChannelOrder, ImgChannelType);
729 writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
730 reinterpret_cast<std::uint32_t *
>(Ptr), ImgChannelOrder,
734 writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
735 reinterpret_cast<std::int8_t *
>(Ptr), ImgChannelOrder,
739 writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
740 reinterpret_cast<std::int16_t *
>(Ptr), ImgChannelOrder,
744 writePixel(convertWriteData<std::int32_t>(Color, ImgChannelType),
745 reinterpret_cast<std::int32_t *
>(Ptr), ImgChannelOrder,
749 writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
750 reinterpret_cast<std::uint8_t *
>(Ptr), ImgChannelOrder,
754 writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
755 reinterpret_cast<std::uint16_t *
>(Ptr), ImgChannelOrder,
759 writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
760 reinterpret_cast<std::uint32_t *
>(Ptr), ImgChannelOrder,
764 writePixel(convertWriteData<half>(Color, ImgChannelType),
765 reinterpret_cast<half *
>(Ptr), ImgChannelOrder, ImgChannelType);
768 writePixel(convertWriteData<float>(Color, ImgChannelType),
769 reinterpret_cast<float *
>(Ptr), ImgChannelOrder, ImgChannelType);
788 template <
typename DataT>
792 const uint8_t ElementSize) {
794 auto Ptr =
static_cast<unsigned char *
>(BasePtr) +
799 switch (ImageChannelType) {
803 convertReadData<std::int8_t>(
readPixel(
reinterpret_cast<std::int8_t *
>(Ptr),
804 ImageChannelOrder, ImageChannelType),
808 convertReadData<std::int16_t>(
809 readPixel(
reinterpret_cast<std::int16_t *
>(Ptr), ImageChannelOrder,
814 convertReadData<std::uint8_t>(
815 readPixel(
reinterpret_cast<std::uint8_t *
>(Ptr), ImageChannelOrder,
820 convertReadData<std::uint16_t>(
821 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
826 convertReadData<std::uint16_t>(
827 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
832 convertReadData<std::uint16_t>(
833 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
838 convertReadData<std::uint32_t>(
839 readPixel(
reinterpret_cast<std::uint32_t *
>(Ptr), ImageChannelOrder,
844 convertReadData<std::int8_t>(
readPixel(
reinterpret_cast<std::int8_t *
>(Ptr),
845 ImageChannelOrder, ImageChannelType),
849 convertReadData<std::int16_t>(
850 readPixel(
reinterpret_cast<std::int16_t *
>(Ptr), ImageChannelOrder,
855 convertReadData<std::int32_t>(
856 readPixel(
reinterpret_cast<std::int32_t *
>(Ptr), ImageChannelOrder,
861 convertReadData<std::uint8_t>(
862 readPixel(
reinterpret_cast<std::uint8_t *
>(Ptr), ImageChannelOrder,
867 convertReadData<std::uint16_t>(
868 readPixel(
reinterpret_cast<std::uint16_t *
>(Ptr), ImageChannelOrder,
873 convertReadData<std::uint32_t>(
874 readPixel(
reinterpret_cast<std::uint32_t *
>(Ptr), ImageChannelOrder,
879 convertReadData<half>(
readPixel(
reinterpret_cast<half *
>(Ptr),
880 ImageChannelOrder, ImageChannelType),
884 convertReadData<float>(
readPixel(
reinterpret_cast<float *
>(Ptr),
885 ImageChannelOrder, ImageChannelType),
895 template <
typename DataT>
900 const uint8_t ElementSize) {
906 RetData = ReadPixelData<DataT>(PixelCoord, ImgPitch, ImgChannelType,
907 ImgChannelOrder, BasePtr, ElementSize);
917 template <
typename DataT>
923 void *BasePtr,
const uint8_t ElementSize) {
924 std::int32_t i0 = CoordValues.s0(), j0 = CoordValues.s1(),
925 k0 = CoordValues.s2(), i1 = CoordValues.s4(),
926 j1 = CoordValues.s5(), k1 = CoordValues.s6();
928 auto getColorInFloat = [&](int4 V) {
930 getColor<DataT>(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
931 ImgChannelOrder, BasePtr, ElementSize);
932 return Res.template convert<float>();
936 float4 Ci0j0k0 = getColorInFloat(int4{i0, j0, k0, 0});
938 float4 Ci1j0k0 = getColorInFloat(int4{i1, j0, k0, 0});
940 float4 Ci0j1k0 = getColorInFloat(int4{i0, j1, k0, 0});
942 float4 Ci1j1k0 = getColorInFloat(int4{i1, j1, k0, 0});
944 float4 Ci0j0k1 = getColorInFloat(int4{i0, j0, k1, 0});
946 float4 Ci1j0k1 = getColorInFloat(int4{i1, j0, k1, 0});
948 float4 Ci0j1k1 = getColorInFloat(int4{i0, j1, k1, 0});
950 float4 Ci1j1k1 = getColorInFloat(int4{i1, j1, k1, 0});
956 Ci0j0k0 = (1 -
a) * (1 -
b) * (1 - c) * Ci0j0k0;
957 Ci1j0k0 =
a * (1 -
b) * (1 - c) * Ci1j0k0;
958 Ci0j1k0 = (1 -
a) *
b * (1 - c) * Ci0j1k0;
959 Ci1j1k0 =
a *
b * (1 - c) * Ci1j1k0;
960 Ci0j0k1 = (1 -
a) * (1 -
b) * c * Ci0j0k1;
961 Ci1j0k1 =
a * (1 -
b) * c * Ci1j0k1;
962 Ci0j1k1 = (1 -
a) *
b * c * Ci0j1k1;
963 Ci1j1k1 =
a *
b * c * Ci1j1k1;
965 float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 + Ci1j0k1 +
1016 template <
typename CoordT,
typename DataT>
1022 void *BasePtr, uint8_t ElementSize) {
1025 float4 FloatCoorduvw;
1031 switch (SmplNormMode) {
1033 switch (SmplAddrMode) {
1037 "Sampler used with unsupported configuration of "
1038 "mirrored_repeat/repeat filtering mode with "
1039 "unnormalized coordinates. ");
1049 switch (SmplAddrMode) {
1078 switch (SmplFiltMode) {
1088 getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1089 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1102 RetData = ReadPixelDataLinearFiltMode<DataT>(
1103 CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1104 ImgChannelOrder, BasePtr, ElementSize);
1114 template <
typename CoordT,
typename DataT>
1120 void *BasePtr, uint8_t ElementSize) {
1123 Smpl.get_coordinate_normalization_mode();
1127 return imageReadSamplerHostImpl<CoordT, DataT>(
1128 Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1129 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1133 template <
typename CoordT,
typename DataT>
1139 void *BasePtr, uint8_t ElementSize) {
1144 return imageReadSamplerHostImpl<CoordT, DataT>(
1145 Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1146 ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
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
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
addressing_mode addressing
coordinate_normalization_mode coordinate