DPC++ Runtime
Runtime libraries for oneAPI DPC++
image_accessor_util.hpp
Go to the documentation of this file.
1 //==------------ image_accessor_util.hpp -----------------------------------==//
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 // This file includes some utilities that are used by image accessors on host
9 // device
10 //
11 
12 #pragma once
13 
14 #ifndef __SYCL_DEVICE_ONLY__
15 
16 #include <sycl/aliases.hpp> // for float4, int4, uint4
17 #include <sycl/builtins.hpp> // for clamp, fmax, min
18 #include <sycl/detail/array.hpp> // for array
19 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
20 #include <sycl/detail/generic_type_traits.hpp> // for max_v, min_v, TryToGe...
21 #include <sycl/detail/pi.h> // for PI_ERROR_INVALID_VALUE
22 #include <sycl/detail/type_list.hpp> // for is_contained, type_list
23 #include <sycl/exception.hpp> // for invalid_parameter_error
24 #include <sycl/id.hpp> // for id
25 #include <sycl/image.hpp> // for image_channel_type
26 #include <sycl/range.hpp> // for range
27 #include <sycl/sampler.hpp> // for addressing_mode, coor...
28 #include <sycl/types.hpp> // for vec, operator*, round...
29 
30 #include <cstdint> // for int32_t, uint16_t
31 #include <stddef.h> // for size_t
32 #include <type_traits> // for enable_if_t
33 
34 namespace sycl {
35 inline namespace _V1 {
36 namespace detail {
37 
38 template <typename T>
39 using IsValidCoordType = typename is_contained<
40  T, boost::mp11::mp_unique<type_list<opencl::cl_int, opencl::cl_float,
41  std::int32_t, float>>>::type;
42 
43 // The formula for unnormalization coordinates:
44 // NormalizedCoords = [UnnormalizedCoords[i] * Range[i] for i in range(0, 3)]
45 template <typename T>
46 std::enable_if_t<IsValidCoordType<T>::value, T>
47 UnnormalizeCoordinates(const T &Coords, const range<3> &Range) {
48  return Coords * Range[0];
49 }
50 
51 template <typename T>
52 std::enable_if_t<IsValidCoordType<T>::value, vec<T, 2>>
53 UnnormalizeCoordinates(const vec<T, 2> &Coords, const range<3> &Range) {
54  return {Coords.x() * Range[0], Coords.y() * Range[1]};
55 }
56 
57 template <typename T>
58 std::enable_if_t<IsValidCoordType<T>::value, vec<T, 4>>
59 UnnormalizeCoordinates(const vec<T, 4> &Coords, const range<3> &Range) {
60  return {Coords.x() * Range[0], Coords.y() * Range[1], Coords.z() * Range[2],
61  0};
62 }
63 
64 // Converts the Coordinates from any dimensions into float4.
65 // valid but unused coordinates are written as 0.5 so the Int_uvwsubhalf
66 // calculation won't pass 0.
67 // Non-valid coordinates are written as 0.
68 template <typename T>
69 std::enable_if_t<IsValidCoordType<T>::value, float4> convertToFloat4(T Coords) {
70  return {static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
71 }
72 
73 template <typename T>
74 std::enable_if_t<IsValidCoordType<T>::value, float4>
76  return {static_cast<float>(Coords.x()), static_cast<float>(Coords.y()), 0.5f,
77  0.f};
78 }
79 
80 template <typename T>
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};
85 }
86 
87 // This method compute an offset in bytes for a given Coords.
88 // Retured offset is used to find the address location of a pixel from a base
89 // ptr.
90 template <typename T>
91 std::enable_if_t<std::is_integral_v<T>, size_t>
92 getImageOffset(const T &Coords, const id<3>, const uint8_t ElementSize) {
93  return Coords * ElementSize;
94 }
95 
96 template <typename T>
97 std::enable_if_t<std::is_integral_v<T>, size_t>
98 getImageOffset(const vec<T, 2> &Coords, const id<3> ImgPitch,
99  const uint8_t ElementSize) {
100  return Coords.x() * ElementSize + Coords.y() * ImgPitch[0];
101 }
102 
103 template <typename T>
104 std::enable_if_t<std::is_integral_v<T>, size_t>
105 getImageOffset(const vec<T, 4> &Coords, const id<3> ImgPitch,
106  const uint8_t ElementSize) {
107  return Coords.x() * ElementSize + Coords.y() * ImgPitch[0] +
108  Coords.z() * ImgPitch[1];
109 }
110 
111 // Process float4 Coordinates and return the appropriate Pixel
112 // Coordinates to read from based on Addressing Mode for Nearest filter mode.
113 __SYCL_EXPORT int4 getPixelCoordNearestFiltMode(float4, const addressing_mode,
114  const range<3>);
115 
116 // Process float4 Coordinates and return the appropriate Pixel
117 // Coordinates to read from based on Addressing Mode for Linear filter mode.
118 __SYCL_EXPORT int8 getPixelCoordLinearFiltMode(float4, const addressing_mode,
119  const range<3>, float4 &);
120 
121 // Check if PixelCoord are out of range for Sampler with clamp adressing mode.
122 __SYCL_EXPORT bool isOutOfRange(const int4 PixelCoord,
123  const addressing_mode SmplAddrMode,
124  const range<3> ImgRange);
125 
126 // Get Border Color for the image_channel_order, the border color values are
127 // only used when the sampler has clamp addressing mode.
128 __SYCL_EXPORT float4 getBorderColor(const image_channel_order ImgChannelOrder);
129 
130 // Reads data from a pixel at Ptr location, based on the number of Channels in
131 // Order and returns the data.
132 // The datatype used to read from the Ptr is based on the T of the
133 // image. This datatype is computed by the calling API.
134 template <typename T>
135 vec<T, 4> readPixel(T *Ptr, const image_channel_order ChannelOrder,
136  const image_channel_type ChannelType) {
137  vec<T, 4> Pixel(0);
138 
139  switch (ChannelOrder) {
141  Pixel.w() = Ptr[0];
142  break;
145  Pixel.x() = Ptr[0];
146  Pixel.w() = 1;
147  break;
149  Pixel.x() = Ptr[0];
150  Pixel.y() = Ptr[0];
151  Pixel.z() = Ptr[0];
152  Pixel.w() = Ptr[0];
153  break;
155  Pixel.x() = Ptr[0];
156  Pixel.y() = Ptr[0];
157  Pixel.z() = Ptr[0];
158  Pixel.w() = 1.0;
159  break;
162  Pixel.x() = Ptr[0];
163  Pixel.y() = Ptr[1];
164  Pixel.w() = 1.0;
165  break;
167  Pixel.x() = Ptr[0];
168  Pixel.w() = Ptr[1];
169  break;
172  if (ChannelType == image_channel_type::unorm_short_565 ||
173  ChannelType == image_channel_type::unorm_short_555 ||
174  ChannelType == image_channel_type::unorm_int_101010) {
175  Pixel.x() = Ptr[0];
176  } else {
177  Pixel.x() = Ptr[0];
178  Pixel.y() = Ptr[1];
179  Pixel.z() = Ptr[2];
180  Pixel.w() = 1.0;
181  }
182  break;
185  Pixel.x() = Ptr[0]; // r
186  Pixel.y() = Ptr[1]; // g
187  Pixel.z() = Ptr[2]; // b
188  Pixel.w() = Ptr[3]; // a
189  break;
191  Pixel.w() = Ptr[0]; // a
192  Pixel.x() = Ptr[1]; // r
193  Pixel.y() = Ptr[2]; // g
194  Pixel.z() = Ptr[3]; // b
195  break;
197  Pixel.z() = Ptr[0]; // b
198  Pixel.y() = Ptr[1]; // g
199  Pixel.x() = Ptr[2]; // r
200  Pixel.w() = Ptr[3]; // a
201  break;
203  Pixel.w() = Ptr[0]; // a
204  Pixel.z() = Ptr[1]; // b
205  Pixel.y() = Ptr[2]; // g
206  Pixel.x() = Ptr[3]; // r
207  break;
208  }
209 
210  return Pixel;
211 }
212 
213 // Write data to a pixel at Ptr location, based on the number of Channels in
214 // ImageChannelOrder. The data passed to this API in 'Pixel' is already
215 // converted to Datatype of the Channel based on ImageChannelType by the calling
216 // API.
217 template <typename T>
218 void writePixel(const vec<T, 4> Pixel, T *Ptr,
219  const image_channel_order ChannelOrder,
220  const image_channel_type ChannelType) {
221 
222  // Data is written based on section 6.12.14.6 of openCL spec.
223  switch (ChannelOrder) {
225  Ptr[0] = Pixel.w();
226  break;
231  Ptr[0] = Pixel.x();
232  break;
235  Ptr[0] = Pixel.x();
236  Ptr[1] = Pixel.y();
237  break;
239  Ptr[0] = Pixel.x();
240  Ptr[1] = Pixel.w();
241  break;
244  if (ChannelType == image_channel_type::unorm_short_565 ||
245  ChannelType == image_channel_type::unorm_short_555 ||
246  ChannelType == image_channel_type::unorm_int_101010) {
247  Ptr[0] = Pixel.x();
248  } else {
249  Ptr[0] = Pixel.x();
250  Ptr[1] = Pixel.y();
251  Ptr[2] = Pixel.z();
252  }
253  break;
256  Ptr[0] = Pixel.x(); // r
257  Ptr[1] = Pixel.y(); // g
258  Ptr[2] = Pixel.z(); // b
259  Ptr[3] = Pixel.w(); // a
260  break;
262  Ptr[0] = Pixel.w(); // a
263  Ptr[1] = Pixel.x(); // r
264  Ptr[2] = Pixel.y(); // g
265  Ptr[3] = Pixel.z(); // b
266  break;
268  Ptr[0] = Pixel.z(); // b
269  Ptr[1] = Pixel.y(); // g
270  Ptr[2] = Pixel.x(); // r
271  Ptr[3] = Pixel.w(); // a
272  break;
274  Ptr[0] = Pixel.w(); // a
275  Ptr[1] = Pixel.z(); // b
276  Ptr[2] = Pixel.y(); // g
277  Ptr[3] = Pixel.x(); // r
278  break;
279  }
280 }
281 
282 // Converts read pixel data into return datatype based on the channel type of
283 // the image.
284 // Conversion rules used as given in the OpenCL
285 // Spec section 8.3. The conversion rules may be handled differently for each
286 // return datatype - float, int32, uint32, half. ImageChannelType is passed to
287 // the function to use appropriate conversion rules.
288 
289 template <typename ChannelType>
291  const image_channel_type ImageChannelType,
292  uint4 &RetData) {
293 
294  switch (ImageChannelType) {
298  RetData = PixelData.template convert<std::uint32_t>();
299  break;
300  default:
301  // OpenCL Spec section 6.12.14.2 does not allow reading uint4 data from an
302  // image with channel datatype other than unsigned_int8,unsigned_int16 and
303  // unsigned_int32.
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);
308  }
309 }
310 
311 template <typename ChannelType>
313  const image_channel_type ImageChannelType, int4 &RetData) {
314 
315  switch (ImageChannelType) {
319  RetData = PixelData.template convert<std::int32_t>();
320  break;
321  default:
322  // OpenCL Spec section 6.12.14.2 does not allow reading int4 data from an
323  // image with channel datatype other than signed_int8,signed_int16 and
324  // signed_int32.
325  throw sycl::invalid_parameter_error(
326  "Datatype of read data - cl_int4 is incompatible with "
327  "the "
328  "image_channel_type of the image.",
329  PI_ERROR_INVALID_VALUE);
330  }
331 }
332 
333 template <typename ChannelType>
335  const image_channel_type ImageChannelType,
336  float4 &RetData) {
337 
338  switch (ImageChannelType) {
340  // max(-1.0f, (float)c / 127.0f)
341  RetData = (PixelData.template convert<float>()) / 127.0f;
342  RetData = sycl::fmax(RetData, -1);
343  break;
345  // max(-1.0f, (float)c / 32767.0f)
346  RetData = (PixelData.template convert<float>()) / 32767.0f;
347  RetData = sycl::fmax(RetData, -1);
348  break;
350  // (float)c / 255.0f
351  RetData = (PixelData.template convert<float>()) / 255.0f;
352  break;
354  // (float)c / 65535.0f
355  RetData = (PixelData.template convert<float>()) / 65535.0f;
356  break;
358  // TODO: Missing information in OpenCL spec. check if the below code is
359  // correct after the spec is updated.
360  // Assuming: (float)c / 31.0f; c represents the 5-bit integer.
361  // (float)c / 63.0f; c represents the 6-bit integer.
362  // PixelData.x will be of type std::uint16_t.
363  ushort4 Temp(PixelData.x());
364  ushort4 MaskBits(0xF800 /*r:bits 11-15*/, 0x07E0 /*g:bits 5-10*/,
365  0x001F /*b:bits 0-4*/, 0x0000);
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;
370  break;
371  }
373  // TODO: Missing information in OpenCL spec. check if the below code is
374  // correct after the spec is updated.
375  // Assuming: (float)c / 31.0f; c represents the 5-bit integer.
376 
377  // Extracting each 5-bit channel data.
378  // PixelData.x will be of type std::uint16_t.
379  ushort4 Temp(PixelData.x());
380  ushort4 MaskBits(0x7C00 /*r:bits 10-14*/, 0x03E0 /*g:bits 5-9*/,
381  0x001F /*b:bits 0-4*/, 0x0000);
382  ushort4 ShiftBits(10, 5, 0, 0);
383  Temp = (Temp & MaskBits) >> ShiftBits;
384  RetData = (Temp.template convert<float>()) / 31.0f;
385  break;
386  }
388  // Extracting each 10-bit channel data.
389  // PixelData.x will be of type std::uint32_t.
390  uint4 Temp(PixelData.x());
391  uint4 MaskBits(0x3FF00000 /*r:bits 20-29*/, 0x000FFC00 /*g:bits 10-19*/,
392  0x000003FF /*b:bits 0-9*/, 0x00000000);
393  uint4 ShiftBits(20, 10, 0, 0);
394  Temp = (Temp & MaskBits) >> ShiftBits;
395  RetData = (Temp.template convert<float>()) / 1023.0f;
396  break;
397  }
404  // OpenCL Spec section 6.12.14.2 does not allow reading float4 data from an
405  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
406  // and signed/unsigned_int32.
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);
412  // Host has conversion from float to half with accuracy as required in
413  // section 8.3.2 OpenCL spec.
414  RetData = PixelData.template convert<float>();
415  break;
417  RetData = PixelData.template convert<float>();
418  break;
419  }
420 }
421 
422 template <typename ChannelType>
424  const image_channel_type ImageChannelType,
425  half4 &RetData) {
426  float4 RetDataFloat;
427  switch (ImageChannelType) {
429  // max(-1.0f, (half)c / 127.0f)
430  RetDataFloat = (PixelData.template convert<float>()) / 127.0f;
431  RetDataFloat = sycl::fmax(RetDataFloat, -1);
432  break;
434  // max(-1.0f, (half)c / 32767.0f)
435  RetDataFloat = (PixelData.template convert<float>()) / 32767.0f;
436  RetDataFloat = sycl::fmax(RetDataFloat, -1);
437  break;
439  // (half)c / 255.0f
440  RetDataFloat = (PixelData.template convert<float>()) / 255.0f;
441  break;
443  // (half)c / 65535.0f
444  RetDataFloat = (PixelData.template convert<float>()) / 65535.0f;
445  break;
449  // TODO: Missing information in OpenCL spec.
450  throw sycl::feature_not_supported(
451  "Currently unsupported datatype conversion from image_channel_type "
452  "to cl_half4.",
453  PI_ERROR_INVALID_OPERATION);
460  // OpenCL Spec section 6.12.14.2 does not allow reading float4 data to an
461  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
462  // and signed/unsigned_int32.
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>();
469  return;
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);
475  }
476  RetData = RetDataFloat.template convert<half>();
477 }
478 
479 // Converts data to write into appropriate datatype based on the channel of the
480 // image.
481 // The conversion rules used are as given in OpenCL Spec Section 8.3. The
482 // conversion rules are different for each return datatype - float,
483 // int32, uint32, half. ImageChannelType is passed to the function to use
484 // appropriate conversion rules.
485 template <typename ChannelType>
487 convertWriteData(const uint4 WriteData,
488  const image_channel_type ImageChannelType) {
489  switch (ImageChannelType) {
491  // convert_uchar_sat(Data)
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>();
496  }
498  // convert_ushort_sat(Data)
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>();
503  }
505  // no conversion is performed.
506  return WriteData.convert<ChannelType>();
507  default:
508  // OpenCL Spec section 6.12.14.4 does not allow writing uint4 data to an
509  // image with channel datatype other than unsigned_int8,unsigned_int16 and
510  // unsigned_int32.
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);
515  }
516 }
517 
518 template <typename ChannelType>
520 convertWriteData(const int4 WriteData,
521  const image_channel_type ImageChannelType) {
522 
523  switch (ImageChannelType) {
525  // convert_char_sat(Data)
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>();
530  }
532  // convert_short_sat(Data)
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>();
537  }
539  return WriteData.convert<ChannelType>();
540  default:
541  // OpenCL Spec section 6.12.14.4 does not allow writing int4 data to an
542  // image with channel datatype other than signed_int8,signed_int16 and
543  // signed_int32.
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);
548  }
549 }
550 
551 template <typename ChannelType>
552 vec<ChannelType, 4> processFloatDataToPixel(float4 WriteData, float MulFactor) {
553  float4 Temp = WriteData * MulFactor;
554  int4 TempInInt = Temp.convert<int, rounding_mode::rte>();
555  int4 TempInIntSaturated =
556  sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
557  return TempInIntSaturated.convert<ChannelType>();
558 }
559 
560 template <typename ChannelType>
562 convertWriteData(const float4 WriteData,
563  const image_channel_type ImageChannelType) {
564 
565  vec<ChannelType, 4> PixelData;
566 
567  switch (ImageChannelType) {
569  // convert_char_sat_rte(f * 127.0f)
570  return processFloatDataToPixel<ChannelType>(WriteData, 127.0f);
572  // convert_short_sat_rte(f * 32767.0f)
573  return processFloatDataToPixel<ChannelType>(WriteData, 32767.0f);
575  // convert_uchar_sat_rte(f * 255.0f)
576  return processFloatDataToPixel<ChannelType>(WriteData, 255.0f);
578  // convert_ushort_sat_rte(f * 65535.0f)
579  return processFloatDataToPixel<ChannelType>(WriteData, 65535.0f);
581  // TODO: Missing information in OpenCL spec.
582  throw sycl::feature_not_supported(
583  "Currently unsupported datatype conversion from image_channel_type "
584  "to cl_float4.",
585  PI_ERROR_INVALID_OPERATION);
587  // TODO: Missing information in OpenCL spec.
588  // Check if the below code is correct after the spec is updated.
589  // Assuming: min(convert_ushort_sat_rte(f * 32.0f), 0x1f)
590  // bits 9:5 and B in bits 4:0.
591  {
592  ushort4 PixelData =
593  processFloatDataToPixel<std::uint16_t>(WriteData, 32.0f);
594  PixelData = sycl::min(PixelData, static_cast<ChannelType>(0x1f));
595  // Compressing the data into the first element of PixelData.
596  // This is needed so that the data can be directly stored into the pixel
597  // location from the first element.
598  // For CL_UNORM_SHORT_555, bit 15 is undefined, R is in bits 14:10, G
599  // in bits 9:5 and B in bits 4:0
600  PixelData.x() =
601  (PixelData.x() << 10) | (PixelData.y() << 5) | PixelData.z();
602  return PixelData.convert<ChannelType>();
603  }
605  // min(convert_ushort_sat_rte(f * 1023.0f), 0x3ff)
606  // For CL_UNORM_INT_101010, bits 31:30 are undefined, R is in bits 29:20, G
607  // in bits 19:10 and B in bits 9:0
608  {
609  uint4 PixelData =
610  processFloatDataToPixel<std::uint32_t>(WriteData, 1023.0f);
611  PixelData = sycl::min(PixelData, static_cast<ChannelType>(0x3ff));
612  PixelData.x() =
613  (PixelData.x() << 20) | (PixelData.y() << 10) | PixelData.z();
614  return PixelData.convert<ChannelType>();
615  }
622  // OpenCL Spec section 6.12.14.4 does not allow writing float4 data to an
623  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
624  // and signed/unsigned_int32.
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);
630  // Host has conversion from float to half with accuracy as required in
631  // section 8.3.2 OpenCL spec.
632  return WriteData.convert<ChannelType>();
634  return WriteData.convert<ChannelType>();
635  }
636 }
637 
638 template <typename ChannelType>
640 convertWriteData(const half4 WriteData,
641  const image_channel_type ImageChannelType) {
642  float4 WriteDataFloat = WriteData.convert<float>();
643  switch (ImageChannelType) {
645  // convert_char_sat_rte(h * 127.0f)
646  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
648  // convert_short_sat_rte(h * 32767.0f)
649  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
651  // convert_uchar_sat_rte(h * 255.0f)
652  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
654  // convert_ushort_sat_rte(h * 65535.0f)
655  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 65535.0f);
659  // TODO: Missing information in OpenCL spec.
660  throw sycl::feature_not_supported(
661  "Currently unsupported datatype conversion from image_channel_type "
662  "to cl_half4.",
663  PI_ERROR_INVALID_OPERATION);
670  // OpenCL Spec section 6.12.14.4 does not allow writing float4 data to an
671  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
672  // and signed/unsigned_int32.
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);
684  }
685 }
686 
687 // imageWriteHostImpl method is called by the write API in image accessors for
688 // host device. Steps:
689 // 1. Calculates the offset from the base ptr of the image where the pixel
690 // denoted by Coord is located.(getImageOffset method.)
691 // 2. Converts the ptr to the appropriate datatype based on
692 // ImageChannelType.(reinterpret_cast)
693 // 3. The data is converted to the image pixel data based on conversion rules in
694 // the spec.(convertWriteData)
695 // 4. The converted data is then written to the pixel at Ptr, based on Number of
696 // Channels in the Image.(writePixel)
697 // Note: We assume that Coords are in the appropriate image range. OpenCL
698 // Spec says that the behaviour is undefined when the Coords are passed outside
699 // the image range. In the current implementation, the data gets written to the
700 // calculated Ptr.
701 template <typename CoordT, typename WriteDataT>
702 void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color,
703  id<3> ImgPitch, uint8_t ElementSize,
704  image_channel_type ImgChannelType,
705  image_channel_order ImgChannelOrder, void *BasePtr) {
706  // Calculate position to write
707  auto Ptr = static_cast<unsigned char *>(BasePtr) +
708  getImageOffset(Coords, ImgPitch, ElementSize);
709 
710  switch (ImgChannelType) {
712  writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
713  reinterpret_cast<std::int8_t *>(Ptr), ImgChannelOrder,
714  ImgChannelType);
715  break;
717  writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
718  reinterpret_cast<std::int16_t *>(Ptr), ImgChannelOrder,
719  ImgChannelType);
720  break;
722  writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
723  reinterpret_cast<std::uint8_t *>(Ptr), ImgChannelOrder,
724  ImgChannelType);
725  break;
727  writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
728  reinterpret_cast<std::uint16_t *>(Ptr), ImgChannelOrder,
729  ImgChannelType);
730  break;
732  writePixel(convertWriteData<short>(Color, ImgChannelType),
733  reinterpret_cast<short *>(Ptr), ImgChannelOrder, ImgChannelType);
734  break;
736  writePixel(convertWriteData<short>(Color, ImgChannelType),
737  reinterpret_cast<short *>(Ptr), ImgChannelOrder, ImgChannelType);
738  break;
740  writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
741  reinterpret_cast<std::uint32_t *>(Ptr), ImgChannelOrder,
742  ImgChannelType);
743  break;
745  writePixel(convertWriteData<std::int8_t>(Color, ImgChannelType),
746  reinterpret_cast<std::int8_t *>(Ptr), ImgChannelOrder,
747  ImgChannelType);
748  break;
750  writePixel(convertWriteData<std::int16_t>(Color, ImgChannelType),
751  reinterpret_cast<std::int16_t *>(Ptr), ImgChannelOrder,
752  ImgChannelType);
753  break;
755  writePixel(convertWriteData<std::int32_t>(Color, ImgChannelType),
756  reinterpret_cast<std::int32_t *>(Ptr), ImgChannelOrder,
757  ImgChannelType);
758  break;
760  writePixel(convertWriteData<std::uint8_t>(Color, ImgChannelType),
761  reinterpret_cast<std::uint8_t *>(Ptr), ImgChannelOrder,
762  ImgChannelType);
763  break;
765  writePixel(convertWriteData<std::uint16_t>(Color, ImgChannelType),
766  reinterpret_cast<std::uint16_t *>(Ptr), ImgChannelOrder,
767  ImgChannelType);
768  break;
770  writePixel(convertWriteData<std::uint32_t>(Color, ImgChannelType),
771  reinterpret_cast<std::uint32_t *>(Ptr), ImgChannelOrder,
772  ImgChannelType);
773  break;
775  writePixel(
776  convertWriteData<half>(Color, ImgChannelType),
777  reinterpret_cast<half *>(Ptr), ImgChannelOrder, ImgChannelType);
778  break;
780  writePixel(convertWriteData<float>(Color, ImgChannelType),
781  reinterpret_cast<float *>(Ptr), ImgChannelOrder, ImgChannelType);
782  break;
783  }
784 }
785 
786 // Method called to read a Coord by getColor function when the Coord is
787 // in-range. This method takes Unnormalized Coords - 'PixelCoord' as int4.
788 // Invalid Coord are denoted by 0. Steps:
789 // 1. Compute Offset for given Unnormalised Coordinates using ImagePitch and
790 // ElementSize.(getImageOffset)
791 // 2. Add this Offset to BasePtr to compute the location of the Image.
792 // 3. Convert this Ptr to the appropriate datatype pointer based on
793 // ImageChannelType. (reinterpret_cast)
794 // 4. Read the appropriate number of channels(computed using
795 // ImageChannelOrder) of the appropriate Channel datatype into Color
796 // variable.(readPixel)
797 // 5. Convert the Read Data into Return DataT based on conversion rules in
798 // the Spec.(convertReadData)
799 // Possible DataT are int4, uint4, float4, half4;
800 template <typename DataT>
801 DataT ReadPixelData(const int4 PixelCoord, const id<3> ImgPitch,
802  const image_channel_type ImageChannelType,
803  const image_channel_order ImageChannelOrder, void *BasePtr,
804  const uint8_t ElementSize) {
805  DataT Color(0);
806  auto Ptr = static_cast<unsigned char *>(BasePtr) +
807  getImageOffset(PixelCoord, ImgPitch,
808  ElementSize); // Utility to compute offset in
809  // image_accessor_util.hpp
810 
811  switch (ImageChannelType) {
812  // TODO: Pass either ImageChannelType or the exact channel type to the
813  // readPixel Function.
815  convertReadData<std::int8_t>(readPixel(reinterpret_cast<std::int8_t *>(Ptr),
816  ImageChannelOrder, ImageChannelType),
818  break;
820  convertReadData<std::int16_t>(
821  readPixel(reinterpret_cast<std::int16_t *>(Ptr), ImageChannelOrder,
822  ImageChannelType),
824  break;
826  convertReadData<std::uint8_t>(
827  readPixel(reinterpret_cast<std::uint8_t *>(Ptr), ImageChannelOrder,
828  ImageChannelType),
830  break;
832  convertReadData<std::uint16_t>(
833  readPixel(reinterpret_cast<std::uint16_t *>(Ptr), ImageChannelOrder,
834  ImageChannelType),
836  break;
838  convertReadData<std::uint16_t>(
839  readPixel(reinterpret_cast<std::uint16_t *>(Ptr), ImageChannelOrder,
840  ImageChannelType),
842  break;
844  convertReadData<std::uint16_t>(
845  readPixel(reinterpret_cast<std::uint16_t *>(Ptr), ImageChannelOrder,
846  ImageChannelType),
848  break;
850  convertReadData<std::uint32_t>(
851  readPixel(reinterpret_cast<std::uint32_t *>(Ptr), ImageChannelOrder,
852  ImageChannelType),
854  break;
856  convertReadData<std::int8_t>(readPixel(reinterpret_cast<std::int8_t *>(Ptr),
857  ImageChannelOrder, ImageChannelType),
859  break;
861  convertReadData<std::int16_t>(
862  readPixel(reinterpret_cast<std::int16_t *>(Ptr), ImageChannelOrder,
863  ImageChannelType),
865  break;
867  convertReadData<std::int32_t>(
868  readPixel(reinterpret_cast<std::int32_t *>(Ptr), ImageChannelOrder,
869  ImageChannelType),
871  break;
873  convertReadData<std::uint8_t>(
874  readPixel(reinterpret_cast<std::uint8_t *>(Ptr), ImageChannelOrder,
875  ImageChannelType),
877  break;
879  convertReadData<std::uint16_t>(
880  readPixel(reinterpret_cast<std::uint16_t *>(Ptr), ImageChannelOrder,
881  ImageChannelType),
883  break;
885  convertReadData<std::uint32_t>(
886  readPixel(reinterpret_cast<std::uint32_t *>(Ptr), ImageChannelOrder,
887  ImageChannelType),
889  break;
891  convertReadData<half>(readPixel(reinterpret_cast<half *>(Ptr),
892  ImageChannelOrder, ImageChannelType),
893  image_channel_type::fp16, Color);
894  break;
896  convertReadData<float>(readPixel(reinterpret_cast<float *>(Ptr),
897  ImageChannelOrder, ImageChannelType),
898  image_channel_type::fp32, Color);
899  break;
900  }
901 
902  return Color;
903 }
904 
905 // Checks if the PixelCoord is out-of-range, and returns appropriate border or
906 // color value at the PixelCoord.
907 template <typename DataT>
908 DataT getColor(const int4 PixelCoord, const addressing_mode SmplAddrMode,
909  const range<3> ImgRange, const id<3> ImgPitch,
910  const image_channel_type ImgChannelType,
911  const image_channel_order ImgChannelOrder, void *BasePtr,
912  const uint8_t ElementSize) {
913  DataT RetData;
914  if (isOutOfRange(PixelCoord, SmplAddrMode, ImgRange)) {
915  float4 BorderColor = getBorderColor(ImgChannelOrder);
916  RetData = BorderColor.convert<get_elem_type_t<DataT>>();
917  } else {
918  RetData = ReadPixelData<DataT>(PixelCoord, ImgPitch, ImgChannelType,
919  ImgChannelOrder, BasePtr, ElementSize);
920  }
921  return RetData;
922 }
923 
924 // Computes and returns color value with Linear Filter Mode.
925 // Steps:
926 // 1. Computes the 8 coordinates using all combinations of i0/i1,j0/j1,k0/k1.
927 // 2. Calls getColor() on each Coordinate.(Ci*j*k*)
928 // 3. Computes the return Color Value using a,b,c and the Color values.
929 template <typename DataT>
930 DataT ReadPixelDataLinearFiltMode(const int8 CoordValues, const float4 abc,
931  const addressing_mode SmplAddrMode,
932  const range<3> ImgRange, id<3> ImgPitch,
933  const image_channel_type ImgChannelType,
934  const image_channel_order ImgChannelOrder,
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();
939 
940  auto getColorInFloat = [&](int4 V) {
941  DataT Res =
942  getColor<DataT>(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
943  ImgChannelOrder, BasePtr, ElementSize);
944  return Res.template convert<float>();
945  };
946 
947  // Get Color Values at each Coordinate.
948  float4 Ci0j0k0 = getColorInFloat(int4{i0, j0, k0, 0});
949 
950  float4 Ci1j0k0 = getColorInFloat(int4{i1, j0, k0, 0});
951 
952  float4 Ci0j1k0 = getColorInFloat(int4{i0, j1, k0, 0});
953 
954  float4 Ci1j1k0 = getColorInFloat(int4{i1, j1, k0, 0});
955 
956  float4 Ci0j0k1 = getColorInFloat(int4{i0, j0, k1, 0});
957 
958  float4 Ci1j0k1 = getColorInFloat(int4{i1, j0, k1, 0});
959 
960  float4 Ci0j1k1 = getColorInFloat(int4{i0, j1, k1, 0});
961 
962  float4 Ci1j1k1 = getColorInFloat(int4{i1, j1, k1, 0});
963 
964  float a = abc.x();
965  float b = abc.y();
966  float c = abc.z();
967 
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;
976 
977  float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 + Ci1j0k1 +
978  Ci0j1k1 + Ci1j1k1;
979 
980  // For 2D image:k0 = 0, k1 = 0, c = 0.5
981  // RetData = (1 – a) * (1 – b) * Ci0j0 + a * (1 – b) * Ci1j0 +
982  // (1 – a) * b * Ci0j1 + a * b * Ci1j1;
983  // For 1D image: j0 = 0, j1 = 0, k0 = 0, k1 = 0, b = 0.5, c = 0.5.
984  // RetData = (1 – a) * Ci0 + a * Ci1;
985  return RetData.convert<get_elem_type_t<DataT>>();
986 }
987 
988 // imageReadSamplerHostImpl method is called by the read API in image accessors
989 // for host device.
990 // Algorithm used: The Algorithm is based on OpenCL spec section 8.2.
991 // It can be broken down into three major steps:
992 // Step 1.
993 // Check for valid sampler options and Compute u,v,w coordinates:
994 // These coordinates are used to compute the Pixel Coordinates that will be
995 // read from to compute the return values.
996 // u,v,w are normalized for AddrMode:mirror_repeat and repeat.
997 // u,v,w are unnormalized for AddrMode:clamp_to_edge, clamp, none.
998 // Convert normalized into unnormalized coords using image range.
999 // note: When dims=1, u,v,w={u,0,0}
1000 // dims=2, u,v,w={u,v,0}
1001 // dims=3, u,v,w-{u,v,w}
1002 // Step 2.
1003 // Process u,v,w, to find the exact Coordinates to read from:
1004 // if(Nearest Filtering Mode)
1005 // compute i,j,k pixel Coordinates based on AddrMode.
1006 // else(Linear Filtering Mode)
1007 // compute i0,j0,k0,i1,j1,k1,a,b,c values.
1008 // Used to load following number of pixels in Step 3.
1009 // 2x2x2 image for Dims=3
1010 // 2x2 image for Dims=2
1011 // 1 pixel for Dims=1 // I think same value should be
1012 // returned as nearest case.
1013 // Step 3.
1014 // Load Image Data, Different for Linear and Nearest Mode:
1015 // Offset = getOffset based on Coord, ImageRange,ImagePitch.
1016 // Read values in the appropriate format based on ImgChannelOrder and
1017 // ImgChannelType.
1018 // Convert to DataT as per conversion rules in section 8.3 in OpenCL Spec.
1019 //
1020 // TODO: Add additional check for half datatype read.
1021 // Based on OpenCL spec 2.0:
1022 // "The read_imageh calls that take integer coordinates must use a sampler with
1023 // filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to
1024 // CLK_NORMALIZED_COORDS_FALSE and addressing mode set to
1025 // CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise
1026 // the values returned are undefined."
1027 
1028 template <typename CoordT, typename DataT>
1030  const CoordT &Coords, coordinate_normalization_mode SmplNormMode,
1031  addressing_mode SmplAddrMode, filtering_mode SmplFiltMode,
1032  /*All image information*/ range<3> ImgRange, id<3> ImgPitch,
1033  image_channel_type ImgChannelType, image_channel_order ImgChannelOrder,
1034  void *BasePtr, uint8_t ElementSize) {
1035 
1036  CoordT Coorduvw;
1037  float4 FloatCoorduvw;
1038  DataT RetData;
1039 
1040  // Step 1:
1041  // switch-case code is used for a better view on value of Coorduvw for all
1042  // combinations of Addressing Modes and Normalization Mode.
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 "
1051  "coordinates. ",
1052  PI_ERROR_INVALID_OPERATION);
1055  case addressing_mode::none:
1056  // Continue with the unnormalized coordinates in Coorduvw.
1057  Coorduvw = Coords;
1058  break;
1059  }
1060  break; // Break for coordinate_normalization_mode::unnormalized.
1062  switch (SmplAddrMode) {
1065  // Continue with the normalized coordinates in Coorduvw.
1066  // Based on Section 8.2 Normalised coordinates are used to compute pixel
1067  // coordinates for addressing_mode::repeat and mirrored_repeat.
1068  Coorduvw = Coords;
1069  break;
1072  case addressing_mode::none:
1073  // Unnormalize these coordinates.
1074  // Based on Section 8.2 Normalised coordinats are used to compute pixel
1075  // coordinates for addressing_mode::clamp/clamp_to_edge and none.
1076  Coorduvw = UnnormalizeCoordinates(Coords, ImgRange);
1077  break;
1078  }
1079  break; // Break for coordinate_normalization_mode::normalized.
1080  }
1081 
1082  // Step 2 & Step 3:
1083 
1084  // converToFloat4 converts CoordT of any kind - std::int32_t, int2, int4,
1085  // float, float2 and float4 into Coordinates of kind float4 with no loss of
1086  // precision. For pixel_coordinates already in float4 format, the function
1087  // returns the same values. This conversion is done to enable implementation
1088  // of one common function getPixelCoordXXXMode, for any datatype of CoordT
1089  // passed.
1090  FloatCoorduvw = convertToFloat4(Coorduvw);
1091  switch (SmplFiltMode) {
1092  case filtering_mode::nearest: {
1093  // Get Pixel Coordinates in integers that will be read from in the Image.
1094  int4 PixelCoord =
1095  getPixelCoordNearestFiltMode(FloatCoorduvw, SmplAddrMode, ImgRange);
1096 
1097  // Return Border Color for out-of-range coordinates when Sampler has
1098  // addressing_mode::clamp. For all other cases and for in-range coordinates
1099  // read the color and return in DataT type.
1100  RetData =
1101  getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1102  ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1103  break;
1104  }
1105  case filtering_mode::linear: {
1106  float4 Retabc;
1107  // Get Pixel Coordinates in integers that will be read from in the Image.
1108  // Return i0,j0,k0,0,i1,j1,k1,0 to form 8 coordinates in a 3D image and
1109  // multiplication factors a,b,c
1110  int8 CoordValues = getPixelCoordLinearFiltMode(FloatCoorduvw, SmplAddrMode,
1111  ImgRange, Retabc);
1112 
1113  // Find the 8 coordinates with the values in CoordValues.
1114  // Computes the Color Value to return.
1115  RetData = ReadPixelDataLinearFiltMode<DataT>(
1116  CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1117  ImgChannelOrder, BasePtr, ElementSize);
1118 
1119  break;
1120  }
1121  }
1122 
1123  return RetData;
1124 }
1125 
1126 // SYCL 1.2.1 sampler overload.
1127 template <typename CoordT, typename DataT>
1128 DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,
1129  /*All image information*/ range<3> ImgRange,
1130  id<3> ImgPitch,
1131  image_channel_type ImgChannelType,
1132  image_channel_order ImgChannelOrder,
1133  void *BasePtr, uint8_t ElementSize) {
1134 
1135  coordinate_normalization_mode SmplNormMode =
1136  Smpl.get_coordinate_normalization_mode();
1137  addressing_mode SmplAddrMode = Smpl.get_addressing_mode();
1138  filtering_mode SmplFiltMode = Smpl.get_filtering_mode();
1139 
1140  return imageReadSamplerHostImpl<CoordT, DataT>(
1141  Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1142  ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1143 }
1144 
1145 // SYCL 2020 image_sampler overload.
1146 template <typename CoordT, typename DataT>
1147 DataT imageReadSamplerHostImpl(const CoordT &Coords, const image_sampler &Smpl,
1148  /*All image information*/ range<3> ImgRange,
1149  id<3> ImgPitch,
1150  image_channel_type ImgChannelType,
1151  image_channel_order ImgChannelOrder,
1152  void *BasePtr, uint8_t ElementSize) {
1153  coordinate_normalization_mode SmplNormMode = Smpl.coordinate;
1154  addressing_mode SmplAddrMode = Smpl.addressing;
1155  filtering_mode SmplFiltMode = Smpl.filtering;
1156 
1157  return imageReadSamplerHostImpl<CoordT, DataT>(
1158  Coords, SmplNormMode, SmplAddrMode, SmplFiltMode, ImgRange, ImgPitch,
1159  ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1160 }
1161 
1162 } // namespace detail
1163 } // namespace _V1
1164 } // namespace sycl
1165 #endif
defined(__INTEL_PREVIEW_BREAKING_CHANGES)
Definition: types.hpp:346
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
Definition: type_list.hpp:32
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
Definition: type_list.hpp:22
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)
std::int32_t cl_int
Definition: aliases.hpp:134
filtering_mode
Definition: sampler.hpp:31
coordinate_normalization_mode
Definition: sampler.hpp:36
T detail::marray_element_t< T > y T T T maxval[i] T T T a
y y maxval[j] maxval b
y y maxval[j] maxval c
std::enable_if_t< detail::is_vgenfloat_v< T >, T > fmax(T x, typename T::element_type y)
addressing_mode
Definition: sampler.hpp:23
image_channel_order
Definition: image.hpp:57
image_channel_type
Definition: image.hpp:75
Definition: access.hpp:18
addressing_mode addressing
Definition: sampler.hpp:133
coordinate_normalization_mode coordinate
Definition: sampler.hpp:134
filtering_mode filtering
Definition: sampler.hpp:135