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 #include <sycl/builtins.hpp>
16 #include <sycl/detail/export.hpp>
19 #include <sycl/image.hpp>
20 #include <sycl/sampler.hpp>
21 #include <sycl/types.hpp>
22 
23 #include <cmath>
24 
25 namespace sycl {
27 namespace detail {
28 
29 template <typename T>
32 
33 // The formula for unnormalization coordinates:
34 // NormalizedCoords = [UnnormalizedCoords[i] * Range[i] for i in range(0, 3)]
35 template <typename T>
36 std::enable_if_t<IsValidCoordType<T>::value, T>
37 UnnormalizeCoordinates(const T &Coords, const range<3> &Range) {
38  return Coords * Range[0];
39 }
40 
41 template <typename T>
42 std::enable_if_t<IsValidCoordType<T>::value, vec<T, 2>>
43 UnnormalizeCoordinates(const vec<T, 2> &Coords, const range<3> &Range) {
44  return {Coords.x() * Range[0], Coords.y() * Range[1]};
45 }
46 
47 template <typename T>
48 std::enable_if_t<IsValidCoordType<T>::value, vec<T, 4>>
49 UnnormalizeCoordinates(const vec<T, 4> &Coords, const range<3> &Range) {
50  return {Coords.x() * Range[0], Coords.y() * Range[1], Coords.z() * Range[2],
51  0};
52 }
53 
54 // Converts the Coordinates from any dimensions into float4.
55 // valid but unused coordinates are written as 0.5 so the Int_uvwsubhalf
56 // calculation won't pass 0.
57 // Non-valid coordinates are written as 0.
58 template <typename T>
59 std::enable_if_t<IsValidCoordType<T>::value, cl_float4>
60 convertToFloat4(T Coords) {
61  return {static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
62 }
63 
64 template <typename T>
65 std::enable_if_t<IsValidCoordType<T>::value, cl_float4>
67  return {static_cast<float>(Coords.x()), static_cast<float>(Coords.y()), 0.5f,
68  0.f};
69 }
70 
71 template <typename T>
72 std::enable_if_t<IsValidCoordType<T>::value, cl_float4>
74  return {static_cast<float>(Coords.x()), static_cast<float>(Coords.y()),
75  static_cast<float>(Coords.z()), 0.f};
76 }
77 
78 // This method compute an offset in bytes for a given Coords.
79 // Retured offset is used to find the address location of a pixel from a base
80 // ptr.
81 template <typename T>
82 std::enable_if_t<std::is_integral<T>::value, size_t>
83 getImageOffset(const T &Coords, const id<3>, const uint8_t ElementSize) {
84  return Coords * ElementSize;
85 }
86 
87 template <typename T>
88 std::enable_if_t<std::is_integral<T>::value, size_t>
89 getImageOffset(const vec<T, 2> &Coords, const id<3> ImgPitch,
90  const uint8_t ElementSize) {
91  return Coords.x() * ElementSize + Coords.y() * ImgPitch[0];
92 }
93 
94 template <typename T>
95 std::enable_if_t<std::is_integral<T>::value, size_t>
96 getImageOffset(const vec<T, 4> &Coords, const id<3> ImgPitch,
97  const uint8_t ElementSize) {
98  return Coords.x() * ElementSize + Coords.y() * ImgPitch[0] +
99  Coords.z() * ImgPitch[1];
100 }
101 
102 // Process cl_float4 Coordinates and return the appropriate Pixel Coordinates to
103 // read from based on Addressing Mode for Nearest filter mode.
104 __SYCL_EXPORT cl_int4 getPixelCoordNearestFiltMode(cl_float4,
105  const addressing_mode,
106  const range<3>);
107 
108 // Process cl_float4 Coordinates and return the appropriate Pixel Coordinates to
109 // read from based on Addressing Mode for Linear filter mode.
110 __SYCL_EXPORT cl_int8 getPixelCoordLinearFiltMode(cl_float4,
111  const addressing_mode,
112  const range<3>, cl_float4 &);
113 
114 // Check if PixelCoord are out of range for Sampler with clamp adressing mode.
115 __SYCL_EXPORT bool isOutOfRange(const cl_int4 PixelCoord,
116  const addressing_mode SmplAddrMode,
117  const range<3> ImgRange);
118 
119 // Get Border Color for the image_channel_order, the border color values are
120 // only used when the sampler has clamp addressing mode.
121 __SYCL_EXPORT cl_float4
122 getBorderColor(const image_channel_order ImgChannelOrder);
123 
124 // Reads data from a pixel at Ptr location, based on the number of Channels in
125 // Order and returns the data.
126 // The datatype used to read from the Ptr is based on the T of the
127 // image. This datatype is computed by the calling API.
128 template <typename T>
129 vec<T, 4> readPixel(T *Ptr, const image_channel_order ChannelOrder,
130  const image_channel_type ChannelType) {
131  vec<T, 4> Pixel(0);
132 
133  switch (ChannelOrder) {
134  case image_channel_order::a:
135  Pixel.w() = Ptr[0];
136  break;
137  case image_channel_order::r:
138  case image_channel_order::rx:
139  Pixel.x() = Ptr[0];
140  Pixel.w() = 1;
141  break;
142  case image_channel_order::intensity:
143  Pixel.x() = Ptr[0];
144  Pixel.y() = Ptr[0];
145  Pixel.z() = Ptr[0];
146  Pixel.w() = Ptr[0];
147  break;
148  case image_channel_order::luminance:
149  Pixel.x() = Ptr[0];
150  Pixel.y() = Ptr[0];
151  Pixel.z() = Ptr[0];
152  Pixel.w() = 1.0;
153  break;
154  case image_channel_order::rg:
155  case image_channel_order::rgx:
156  Pixel.x() = Ptr[0];
157  Pixel.y() = Ptr[1];
158  Pixel.w() = 1.0;
159  break;
160  case image_channel_order::ra:
161  Pixel.x() = Ptr[0];
162  Pixel.w() = Ptr[1];
163  break;
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) {
169  Pixel.x() = Ptr[0];
170  } else {
171  Pixel.x() = Ptr[0];
172  Pixel.y() = Ptr[1];
173  Pixel.z() = Ptr[2];
174  Pixel.w() = 1.0;
175  }
176  break;
177  case image_channel_order::rgba:
178  case image_channel_order::ext_oneapi_srgba:
179  Pixel.x() = Ptr[0]; // r
180  Pixel.y() = Ptr[1]; // g
181  Pixel.z() = Ptr[2]; // b
182  Pixel.w() = Ptr[3]; // a
183  break;
184  case image_channel_order::argb:
185  Pixel.w() = Ptr[0]; // a
186  Pixel.x() = Ptr[1]; // r
187  Pixel.y() = Ptr[2]; // g
188  Pixel.z() = Ptr[3]; // b
189  break;
190  case image_channel_order::bgra:
191  Pixel.z() = Ptr[0]; // b
192  Pixel.y() = Ptr[1]; // g
193  Pixel.x() = Ptr[2]; // r
194  Pixel.w() = Ptr[3]; // a
195  break;
196  case image_channel_order::abgr:
197  Pixel.w() = Ptr[0]; // a
198  Pixel.z() = Ptr[1]; // b
199  Pixel.y() = Ptr[2]; // g
200  Pixel.x() = Ptr[3]; // r
201  break;
202  }
203 
204  return Pixel;
205 }
206 
207 // Write data to a pixel at Ptr location, based on the number of Channels in
208 // ImageChannelOrder. The data passed to this API in 'Pixel' is already
209 // converted to Datatype of the Channel based on ImageChannelType by the calling
210 // API.
211 template <typename T>
212 void writePixel(const vec<T, 4> Pixel, T *Ptr,
213  const image_channel_order ChannelOrder,
214  const image_channel_type ChannelType) {
215 
216  // Data is written based on section 6.12.14.6 of openCL spec.
217  switch (ChannelOrder) {
218  case image_channel_order::a:
219  Ptr[0] = Pixel.w();
220  break;
221  case image_channel_order::r:
222  case image_channel_order::rx:
223  case image_channel_order::intensity:
224  case image_channel_order::luminance:
225  Ptr[0] = Pixel.x();
226  break;
227  case image_channel_order::rg:
228  case image_channel_order::rgx:
229  Ptr[0] = Pixel.x();
230  Ptr[1] = Pixel.y();
231  break;
232  case image_channel_order::ra:
233  Ptr[0] = Pixel.x();
234  Ptr[1] = Pixel.w();
235  break;
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) {
241  Ptr[0] = Pixel.x();
242  } else {
243  Ptr[0] = Pixel.x();
244  Ptr[1] = Pixel.y();
245  Ptr[2] = Pixel.z();
246  }
247  break;
248  case image_channel_order::rgba:
249  case image_channel_order::ext_oneapi_srgba:
250  Ptr[0] = Pixel.x(); // r
251  Ptr[1] = Pixel.y(); // g
252  Ptr[2] = Pixel.z(); // b
253  Ptr[3] = Pixel.w(); // a
254  break;
255  case image_channel_order::argb:
256  Ptr[0] = Pixel.w(); // a
257  Ptr[1] = Pixel.x(); // r
258  Ptr[2] = Pixel.y(); // g
259  Ptr[3] = Pixel.z(); // b
260  break;
261  case image_channel_order::bgra:
262  Ptr[0] = Pixel.z(); // b
263  Ptr[1] = Pixel.y(); // g
264  Ptr[2] = Pixel.x(); // r
265  Ptr[3] = Pixel.w(); // a
266  break;
267  case image_channel_order::abgr:
268  Ptr[0] = Pixel.w(); // a
269  Ptr[1] = Pixel.z(); // b
270  Ptr[2] = Pixel.y(); // g
271  Ptr[3] = Pixel.x(); // r
272  break;
273  }
274 }
275 
276 // Converts read pixel data into return datatype based on the channel type of
277 // the image.
278 // Conversion rules used as given in the OpenCL
279 // Spec section 8.3. The conversion rules may be handled differently for each
280 // return datatype - float, int32, uint32, half. ImageChannelType is passed to
281 // the function to use appropriate conversion rules.
282 
283 template <typename ChannelType>
285  const image_channel_type ImageChannelType,
286  vec<cl_uint, 4> &RetData) {
287 
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>();
293  break;
294  default:
295  // OpenCL Spec section 6.12.14.2 does not allow reading uint4 data from an
296  // image with channel datatype other than unsigned_int8,unsigned_int16 and
297  // unsigned_int32.
298  throw sycl::invalid_parameter_error(
299  "Datatype of read data - cl_uint4 is incompatible with the "
300  "image_channel_type of the image.",
301  PI_ERROR_INVALID_VALUE);
302  }
303 }
304 
305 template <typename ChannelType>
307  const image_channel_type ImageChannelType,
308  vec<cl_int, 4> &RetData) {
309 
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>();
315  break;
316  default:
317  // OpenCL Spec section 6.12.14.2 does not allow reading int4 data from an
318  // image with channel datatype other than signed_int8,signed_int16 and
319  // signed_int32.
320  throw sycl::invalid_parameter_error(
321  "Datatype of read data - cl_int4 is incompatible with the "
322  "image_channel_type of the image.",
323  PI_ERROR_INVALID_VALUE);
324  }
325 }
326 
327 template <typename ChannelType>
329  const image_channel_type ImageChannelType,
330  vec<cl_float, 4> &RetData) {
331 
332  switch (ImageChannelType) {
333  case image_channel_type::snorm_int8:
334  // max(-1.0f, (float)c / 127.0f)
335  RetData = (PixelData.template convert<cl_float>()) / 127.0f;
336  RetData = sycl::fmax(RetData, -1);
337  break;
338  case image_channel_type::snorm_int16:
339  // max(-1.0f, (float)c / 32767.0f)
340  RetData = (PixelData.template convert<cl_float>()) / 32767.0f;
341  RetData = sycl::fmax(RetData, -1);
342  break;
343  case image_channel_type::unorm_int8:
344  // (float)c / 255.0f
345  RetData = (PixelData.template convert<cl_float>()) / 255.0f;
346  break;
347  case image_channel_type::unorm_int16:
348  // (float)c / 65535.0f
349  RetData = (PixelData.template convert<cl_float>()) / 65535.0f;
350  break;
351  case image_channel_type::unorm_short_565: {
352  // TODO: Missing information in OpenCL spec. check if the below code is
353  // correct after the spec is updated.
354  // Assuming: (float)c / 31.0f; c represents the 5-bit integer.
355  // (float)c / 63.0f; c represents the 6-bit integer.
356  // PixelData.x will be of type cl_ushort.
357  vec<cl_ushort, 4> Temp(PixelData.x());
358  vec<cl_ushort, 4> MaskBits(0xF800 /*r:bits 11-15*/, 0x07E0 /*g:bits 5-10*/,
359  0x001F /*b:bits 0-4*/, 0x0000);
360  vec<cl_ushort, 4> ShiftBits(11, 5, 0, 0);
361  vec<cl_float, 4> DivisorToNormalise(31.0f, 63.0f, 31.0f, 1);
362  Temp = (Temp & MaskBits) >> ShiftBits;
363  RetData = (Temp.template convert<cl_float>()) / DivisorToNormalise;
364  break;
365  }
366  case image_channel_type::unorm_short_555: {
367  // TODO: Missing information in OpenCL spec. check if the below code is
368  // correct after the spec is updated.
369  // Assuming: (float)c / 31.0f; c represents the 5-bit integer.
370 
371  // Extracting each 5-bit channel data.
372  // PixelData.x will be of type cl_ushort.
373  vec<cl_ushort, 4> Temp(PixelData.x());
374  vec<cl_ushort, 4> MaskBits(0x7C00 /*r:bits 10-14*/, 0x03E0 /*g:bits 5-9*/,
375  0x001F /*b:bits 0-4*/, 0x0000);
376  vec<cl_ushort, 4> ShiftBits(10, 5, 0, 0);
377  Temp = (Temp & MaskBits) >> ShiftBits;
378  RetData = (Temp.template convert<cl_float>()) / 31.0f;
379  break;
380  }
381  case image_channel_type::unorm_int_101010: {
382  // Extracting each 10-bit channel data.
383  // PixelData.x will be of type cl_uint.
384  vec<cl_uint, 4> Temp(PixelData.x());
385  vec<cl_uint, 4> MaskBits(0x3FF00000 /*r:bits 20-29*/,
386  0x000FFC00 /*g:bits 10-19*/,
387  0x000003FF /*b:bits 0-9*/, 0x00000000);
388  vec<cl_uint, 4> ShiftBits(20, 10, 0, 0);
389  Temp = (Temp & MaskBits) >> ShiftBits;
390  RetData = (Temp.template convert<cl_float>()) / 1023.0f;
391  break;
392  }
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:
399  // OpenCL Spec section 6.12.14.2 does not allow reading float4 data from an
400  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
401  // and signed/unsigned_int32.
402  throw sycl::invalid_parameter_error(
403  "Datatype of read data - cl_float4 is incompatible with the "
404  "image_channel_type of the image.",
405  PI_ERROR_INVALID_VALUE);
406  case image_channel_type::fp16:
407  // Host has conversion from float to half with accuracy as required in
408  // section 8.3.2 OpenCL spec.
409  RetData = PixelData.template convert<cl_float>();
410  break;
411  case image_channel_type::fp32:
412  RetData = PixelData.template convert<cl_float>();
413  break;
414  }
415 }
416 
417 template <typename ChannelType>
419  const image_channel_type ImageChannelType,
420  vec<cl_half, 4> &RetData) {
421  vec<cl_float, 4> RetDataFloat;
422  switch (ImageChannelType) {
423  case image_channel_type::snorm_int8:
424  // max(-1.0f, (half)c / 127.0f)
425  RetDataFloat = (PixelData.template convert<cl_float>()) / 127.0f;
426  RetDataFloat = sycl::fmax(RetDataFloat, -1);
427  break;
428  case image_channel_type::snorm_int16:
429  // max(-1.0f, (half)c / 32767.0f)
430  RetDataFloat = (PixelData.template convert<cl_float>()) / 32767.0f;
431  RetDataFloat = sycl::fmax(RetDataFloat, -1);
432  break;
433  case image_channel_type::unorm_int8:
434  // (half)c / 255.0f
435  RetDataFloat = (PixelData.template convert<cl_float>()) / 255.0f;
436  break;
437  case image_channel_type::unorm_int16:
438  // (half)c / 65535.0f
439  RetDataFloat = (PixelData.template convert<cl_float>()) / 65535.0f;
440  break;
441  case image_channel_type::unorm_short_565:
442  case image_channel_type::unorm_short_555:
443  case image_channel_type::unorm_int_101010:
444  // TODO: Missing information in OpenCL spec.
445  throw sycl::feature_not_supported(
446  "Currently unsupported datatype conversion from image_channel_type "
447  "to cl_half4.",
448  PI_ERROR_INVALID_OPERATION);
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:
455  // OpenCL Spec section 6.12.14.2 does not allow reading float4 data to an
456  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
457  // and signed/unsigned_int32.
458  throw sycl::invalid_parameter_error(
459  "Datatype to read- cl_half4 is incompatible with the "
460  "image_channel_type of the image.",
461  PI_ERROR_INVALID_VALUE);
462  case image_channel_type::fp16:
463  RetData = PixelData.template convert<cl_half>();
464  return;
465  case image_channel_type::fp32:
466  throw sycl::invalid_parameter_error(
467  "Datatype to read - cl_half4 is incompatible with the "
468  "image_channel_type of the image.",
469  PI_ERROR_INVALID_VALUE);
470  }
471  RetData = RetDataFloat.template convert<cl_half>();
472 }
473 
474 // Converts data to write into appropriate datatype based on the channel of the
475 // image.
476 // The conversion rules used are as given in OpenCL Spec Section 8.3. The
477 // conversion rules are different for each return datatype - float,
478 // int32, uint32, half. ImageChannelType is passed to the function to use
479 // appropriate conversion rules.
480 template <typename ChannelType>
483  const image_channel_type ImageChannelType) {
484  switch (ImageChannelType) {
485  case image_channel_type::unsigned_int8: {
486  // convert_uchar_sat(Data)
487  cl_uint MinVal = min_v<cl_uchar>();
488  cl_uint MaxVal = max_v<cl_uchar>();
489  vec<cl_uint, 4> PixelData = sycl::clamp(WriteData, MinVal, MaxVal);
490  return PixelData.convert<ChannelType>();
491  }
492  case image_channel_type::unsigned_int16: {
493  // convert_ushort_sat(Data)
494  cl_uint MinVal = min_v<cl_ushort>();
495  cl_uint MaxVal = max_v<cl_ushort>();
496  vec<cl_uint, 4> PixelData = sycl::clamp(WriteData, MinVal, MaxVal);
497  return PixelData.convert<ChannelType>();
498  }
499  case image_channel_type::unsigned_int32:
500  // no conversion is performed.
501  return WriteData.convert<ChannelType>();
502  default:
503  // OpenCL Spec section 6.12.14.4 does not allow writing uint4 data to an
504  // image with channel datatype other than unsigned_int8,unsigned_int16 and
505  // unsigned_int32.
506  throw sycl::invalid_parameter_error(
507  "Datatype of data to write - cl_uint4 is incompatible with the "
508  "image_channel_type of the image.",
509  PI_ERROR_INVALID_VALUE);
510  }
511 }
512 
513 template <typename ChannelType>
516  const image_channel_type ImageChannelType) {
517 
518  switch (ImageChannelType) {
519  case image_channel_type::signed_int8: {
520  // convert_char_sat(Data)
521  cl_int MinVal = min_v<cl_char>();
522  cl_int MaxVal = max_v<cl_char>();
523  vec<cl_int, 4> PixelData = sycl::clamp(WriteData, MinVal, MaxVal);
524  return PixelData.convert<ChannelType>();
525  }
526  case image_channel_type::signed_int16: {
527  // convert_short_sat(Data)
528  cl_int MinVal = min_v<cl_short>();
529  cl_int MaxVal = max_v<cl_short>();
530  vec<cl_int, 4> PixelData = sycl::clamp(WriteData, MinVal, MaxVal);
531  return PixelData.convert<ChannelType>();
532  }
533  case image_channel_type::signed_int32:
534  return WriteData.convert<ChannelType>();
535  default:
536  // OpenCL Spec section 6.12.14.4 does not allow writing int4 data to an
537  // image with channel datatype other than signed_int8,signed_int16 and
538  // signed_int32.
539  throw sycl::invalid_parameter_error(
540  "Datatype of data to write - cl_int4 is incompatible with the "
541  "image_channel_type of the image.",
542  PI_ERROR_INVALID_VALUE);
543  }
544 }
545 
546 template <typename ChannelType>
548  float MulFactor) {
549  vec<cl_float, 4> Temp = WriteData * MulFactor;
550  vec<cl_int, 4> TempInInt = Temp.convert<int, rounding_mode::rte>();
551  vec<cl_int, 4> TempInIntSaturated =
552  sycl::clamp(TempInInt, min_v<ChannelType>(), max_v<ChannelType>());
553  return TempInIntSaturated.convert<ChannelType>();
554 }
555 
556 template <typename ChannelType>
559  const image_channel_type ImageChannelType) {
560 
561  vec<ChannelType, 4> PixelData;
562 
563  switch (ImageChannelType) {
564  case image_channel_type::snorm_int8:
565  // convert_char_sat_rte(f * 127.0f)
566  return processFloatDataToPixel<ChannelType>(WriteData, 127.0f);
567  case image_channel_type::snorm_int16:
568  // convert_short_sat_rte(f * 32767.0f)
569  return processFloatDataToPixel<ChannelType>(WriteData, 32767.0f);
570  case image_channel_type::unorm_int8:
571  // convert_uchar_sat_rte(f * 255.0f)
572  return processFloatDataToPixel<ChannelType>(WriteData, 255.0f);
573  case image_channel_type::unorm_int16:
574  // convert_ushort_sat_rte(f * 65535.0f)
575  return processFloatDataToPixel<ChannelType>(WriteData, 65535.0f);
576  case image_channel_type::unorm_short_565:
577  // TODO: Missing information in OpenCL spec.
578  throw sycl::feature_not_supported(
579  "Currently unsupported datatype conversion from image_channel_type "
580  "to cl_float4.",
581  PI_ERROR_INVALID_OPERATION);
582  case image_channel_type::unorm_short_555:
583  // TODO: Missing information in OpenCL spec.
584  // Check if the below code is correct after the spec is updated.
585  // Assuming: min(convert_ushort_sat_rte(f * 32.0f), 0x1f)
586  // bits 9:5 and B in bits 4:0.
587  {
588  vec<cl_ushort, 4> PixelData =
589  processFloatDataToPixel<cl_ushort>(WriteData, 32.0f);
590  PixelData = sycl::min(PixelData, static_cast<ChannelType>(0x1f));
591  // Compressing the data into the first element of PixelData.
592  // This is needed so that the data can be directly stored into the pixel
593  // location from the first element.
594  // For CL_UNORM_SHORT_555, bit 15 is undefined, R is in bits 14:10, G
595  // in bits 9:5 and B in bits 4:0
596  PixelData.x() =
597  (PixelData.x() << 10) | (PixelData.y() << 5) | PixelData.z();
598  return PixelData.convert<ChannelType>();
599  }
600  case image_channel_type::unorm_int_101010:
601  // min(convert_ushort_sat_rte(f * 1023.0f), 0x3ff)
602  // For CL_UNORM_INT_101010, bits 31:30 are undefined, R is in bits 29:20, G
603  // in bits 19:10 and B in bits 9:0
604  {
605  vec<cl_uint, 4> PixelData =
606  processFloatDataToPixel<cl_uint>(WriteData, 1023.0f);
607  PixelData = sycl::min(PixelData, static_cast<ChannelType>(0x3ff));
608  PixelData.x() =
609  (PixelData.x() << 20) | (PixelData.y() << 10) | PixelData.z();
610  return PixelData.convert<ChannelType>();
611  }
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:
618  // OpenCL Spec section 6.12.14.4 does not allow writing float4 data to an
619  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
620  // and signed/unsigned_int32.
621  throw sycl::invalid_parameter_error(
622  "Datatype of data to write - cl_float4 is incompatible with the "
623  "image_channel_type of the image.",
624  PI_ERROR_INVALID_VALUE);
625  case image_channel_type::fp16:
626  // Host has conversion from float to half with accuracy as required in
627  // section 8.3.2 OpenCL spec.
628  return WriteData.convert<ChannelType>();
629  case image_channel_type::fp32:
630  return WriteData.convert<ChannelType>();
631  }
632 }
633 
634 template <typename ChannelType>
637  const image_channel_type ImageChannelType) {
638  vec<cl_float, 4> WriteDataFloat = WriteData.convert<cl_float>();
639  switch (ImageChannelType) {
640  case image_channel_type::snorm_int8:
641  // convert_char_sat_rte(h * 127.0f)
642  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 127.0f);
643  case image_channel_type::snorm_int16:
644  // convert_short_sat_rte(h * 32767.0f)
645  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 32767.0f);
646  case image_channel_type::unorm_int8:
647  // convert_uchar_sat_rte(h * 255.0f)
648  return processFloatDataToPixel<ChannelType>(WriteDataFloat, 255.0f);
649  case image_channel_type::unorm_int16:
650  // convert_ushort_sat_rte(h * 65535.0f)
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:
655  // TODO: Missing information in OpenCL spec.
656  throw sycl::feature_not_supported(
657  "Currently unsupported datatype conversion from image_channel_type "
658  "to cl_half4.",
659  PI_ERROR_INVALID_OPERATION);
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:
666  // OpenCL Spec section 6.12.14.4 does not allow writing float4 data to an
667  // image with channel datatype - signed/unsigned_int8,signed/unsigned_int16
668  // and signed/unsigned_int32.
669  throw sycl::invalid_parameter_error(
670  "Datatype of data to write - cl_float4 is incompatible with the "
671  "image_channel_type of the image.",
672  PI_ERROR_INVALID_VALUE);
673  case image_channel_type::fp16:
674  return WriteData.convert<ChannelType>();
675  case image_channel_type::fp32:
676  throw sycl::invalid_parameter_error(
677  "Datatype of data to write - cl_float4 is incompatible with the "
678  "image_channel_type of the image.",
679  PI_ERROR_INVALID_VALUE);
680  }
681 }
682 
683 // imageWriteHostImpl method is called by the write API in image accessors for
684 // host device. Steps:
685 // 1. Calculates the offset from the base ptr of the image where the pixel
686 // denoted by Coord is located.(getImageOffset method.)
687 // 2. Converts the ptr to the appropriate datatype based on
688 // ImageChannelType.(reinterpret_cast)
689 // 3. The data is converted to the image pixel data based on conversion rules in
690 // the spec.(convertWriteData)
691 // 4. The converted data is then written to the pixel at Ptr, based on Number of
692 // Channels in the Image.(writePixel)
693 // Note: We assume that Coords are in the appropriate image range. OpenCL
694 // Spec says that the behaviour is undefined when the Coords are passed outside
695 // the image range. In the current implementation, the data gets written to the
696 // calculated Ptr.
697 template <typename CoordT, typename WriteDataT>
698 void imageWriteHostImpl(const CoordT &Coords, const WriteDataT &Color,
699  id<3> ImgPitch, uint8_t ElementSize,
700  image_channel_type ImgChannelType,
701  image_channel_order ImgChannelOrder, void *BasePtr) {
702  // Calculate position to write
703  auto Ptr = static_cast<unsigned char *>(BasePtr) +
704  getImageOffset(Coords, ImgPitch, ElementSize);
705 
706  switch (ImgChannelType) {
707  case image_channel_type::snorm_int8:
708  writePixel(convertWriteData<cl_char>(Color, ImgChannelType),
709  reinterpret_cast<cl_char *>(Ptr), ImgChannelOrder,
710  ImgChannelType);
711  break;
712  case image_channel_type::snorm_int16:
713  writePixel(convertWriteData<cl_short>(Color, ImgChannelType),
714  reinterpret_cast<cl_short *>(Ptr), ImgChannelOrder,
715  ImgChannelType);
716  break;
717  case image_channel_type::unorm_int8:
718  writePixel(convertWriteData<cl_uchar>(Color, ImgChannelType),
719  reinterpret_cast<cl_uchar *>(Ptr), ImgChannelOrder,
720  ImgChannelType);
721  break;
722  case image_channel_type::unorm_int16:
723  writePixel(convertWriteData<cl_ushort>(Color, ImgChannelType),
724  reinterpret_cast<cl_ushort *>(Ptr), ImgChannelOrder,
725  ImgChannelType);
726  break;
727  case image_channel_type::unorm_short_565:
728  writePixel(convertWriteData<short>(Color, ImgChannelType),
729  reinterpret_cast<short *>(Ptr), ImgChannelOrder, ImgChannelType);
730  break;
731  case image_channel_type::unorm_short_555:
732  writePixel(convertWriteData<short>(Color, ImgChannelType),
733  reinterpret_cast<short *>(Ptr), ImgChannelOrder, ImgChannelType);
734  break;
735  case image_channel_type::unorm_int_101010:
736  writePixel(convertWriteData<cl_uint>(Color, ImgChannelType),
737  reinterpret_cast<cl_uint *>(Ptr), ImgChannelOrder,
738  ImgChannelType);
739  break;
740  case image_channel_type::signed_int8:
741  writePixel(convertWriteData<cl_char>(Color, ImgChannelType),
742  reinterpret_cast<cl_char *>(Ptr), ImgChannelOrder,
743  ImgChannelType);
744  break;
745  case image_channel_type::signed_int16:
746  writePixel(convertWriteData<cl_short>(Color, ImgChannelType),
747  reinterpret_cast<cl_short *>(Ptr), ImgChannelOrder,
748  ImgChannelType);
749  break;
750  case image_channel_type::signed_int32:
751  writePixel(convertWriteData<cl_int>(Color, ImgChannelType),
752  reinterpret_cast<cl_int *>(Ptr), ImgChannelOrder,
753  ImgChannelType);
754  break;
755  case image_channel_type::unsigned_int8:
756  writePixel(convertWriteData<cl_uchar>(Color, ImgChannelType),
757  reinterpret_cast<cl_uchar *>(Ptr), ImgChannelOrder,
758  ImgChannelType);
759  break;
760  case image_channel_type::unsigned_int16:
761  writePixel(convertWriteData<cl_ushort>(Color, ImgChannelType),
762  reinterpret_cast<cl_ushort *>(Ptr), ImgChannelOrder,
763  ImgChannelType);
764  break;
765  case image_channel_type::unsigned_int32:
766  writePixel(convertWriteData<cl_uint>(Color, ImgChannelType),
767  reinterpret_cast<cl_uint *>(Ptr), ImgChannelOrder,
768  ImgChannelType);
769  break;
770  case image_channel_type::fp16:
771  writePixel(
772  // convertWriteDataToHalf<typename
773  // TryToGetElementType<WriteDataT>::type>(
774  convertWriteData<cl_half>(Color, ImgChannelType),
775  reinterpret_cast<cl_half *>(Ptr), ImgChannelOrder, ImgChannelType);
776  break;
777  case image_channel_type::fp32:
778  writePixel(convertWriteData<cl_float>(Color, ImgChannelType),
779  reinterpret_cast<cl_float *>(Ptr), ImgChannelOrder,
780  ImgChannelType);
781  break;
782  }
783 }
784 
785 // Method called to read a Coord by getColor function when the Coord is
786 // in-range. This method takes Unnormalized Coords - 'PixelCoord' as cl_int4.
787 // Invalid Coord are denoted by 0. Steps:
788 // 1. Compute Offset for given Unnormalised Coordinates using ImagePitch and
789 // ElementSize.(getImageOffset)
790 // 2. Add this Offset to BasePtr to compute the location of the Image.
791 // 3. Convert this Ptr to the appropriate datatype pointer based on
792 // ImageChannelType. (reinterpret_cast)
793 // 4. Read the appropriate number of channels(computed using
794 // ImageChannelOrder) of the appropriate Channel datatype into Color
795 // variable.(readPixel)
796 // 5. Convert the Read Data into Return DataT based on conversion rules in
797 // the Spec.(convertReadData)
798 // Possible DataT are cl_int4, cl_uint4, cl_float4, cl_half4;
799 template <typename DataT>
800 DataT ReadPixelData(const cl_int4 PixelCoord, const id<3> ImgPitch,
801  const image_channel_type ImageChannelType,
802  const image_channel_order ImageChannelOrder, void *BasePtr,
803  const uint8_t ElementSize) {
804  DataT Color(0);
805  auto Ptr = static_cast<unsigned char *>(BasePtr) +
806  getImageOffset(PixelCoord, ImgPitch,
807  ElementSize); // Utility to compute offset in
808  // image_accessor_util.hpp
809 
810  switch (ImageChannelType) {
811  // TODO: Pass either ImageChannelType or the exact channel type to the
812  // readPixel Function.
813  case image_channel_type::snorm_int8:
814  convertReadData<cl_char>(readPixel(reinterpret_cast<cl_char *>(Ptr),
815  ImageChannelOrder, ImageChannelType),
816  image_channel_type::snorm_int8, Color);
817  break;
818  case image_channel_type::snorm_int16:
819  convertReadData<cl_short>(readPixel(reinterpret_cast<cl_short *>(Ptr),
820  ImageChannelOrder, ImageChannelType),
821  image_channel_type::snorm_int16, Color);
822  break;
823  case image_channel_type::unorm_int8:
824  convertReadData<cl_uchar>(readPixel(reinterpret_cast<cl_uchar *>(Ptr),
825  ImageChannelOrder, ImageChannelType),
826  image_channel_type::unorm_int8, Color);
827  break;
828  case image_channel_type::unorm_int16:
829  convertReadData<cl_ushort>(readPixel(reinterpret_cast<cl_ushort *>(Ptr),
830  ImageChannelOrder, ImageChannelType),
831  image_channel_type::unorm_int16, Color);
832  break;
833  case image_channel_type::unorm_short_565:
834  convertReadData<cl_ushort>(readPixel(reinterpret_cast<cl_ushort *>(Ptr),
835  ImageChannelOrder, ImageChannelType),
836  image_channel_type::unorm_short_565, Color);
837  break;
838  case image_channel_type::unorm_short_555:
839  convertReadData<cl_ushort>(readPixel(reinterpret_cast<cl_ushort *>(Ptr),
840  ImageChannelOrder, ImageChannelType),
841  image_channel_type::unorm_short_555, Color);
842  break;
843  case image_channel_type::unorm_int_101010:
844  convertReadData<cl_uint>(readPixel(reinterpret_cast<cl_uint *>(Ptr),
845  ImageChannelOrder, ImageChannelType),
846  image_channel_type::unorm_int_101010, Color);
847  break;
848  case image_channel_type::signed_int8:
849  convertReadData<cl_char>(readPixel(reinterpret_cast<cl_char *>(Ptr),
850  ImageChannelOrder, ImageChannelType),
851  image_channel_type::signed_int8, Color);
852  break;
853  case image_channel_type::signed_int16:
854  convertReadData<cl_short>(readPixel(reinterpret_cast<cl_short *>(Ptr),
855  ImageChannelOrder, ImageChannelType),
856  image_channel_type::signed_int16, Color);
857  break;
858  case image_channel_type::signed_int32:
859  convertReadData<cl_int>(readPixel(reinterpret_cast<cl_int *>(Ptr),
860  ImageChannelOrder, ImageChannelType),
861  image_channel_type::signed_int32, Color);
862  break;
863  case image_channel_type::unsigned_int8:
864  convertReadData<cl_uchar>(readPixel(reinterpret_cast<cl_uchar *>(Ptr),
865  ImageChannelOrder, ImageChannelType),
866  image_channel_type::unsigned_int8, Color);
867  break;
868  case image_channel_type::unsigned_int16:
869  convertReadData<cl_ushort>(readPixel(reinterpret_cast<cl_ushort *>(Ptr),
870  ImageChannelOrder, ImageChannelType),
871  image_channel_type::unsigned_int16, Color);
872  break;
873  case image_channel_type::unsigned_int32:
874  convertReadData<cl_uint>(readPixel(reinterpret_cast<cl_uint *>(Ptr),
875  ImageChannelOrder, ImageChannelType),
876  image_channel_type::unsigned_int32, Color);
877  break;
878  case image_channel_type::fp16:
879  convertReadData<cl_half>(readPixel(reinterpret_cast<cl_half *>(Ptr),
880  ImageChannelOrder, ImageChannelType),
881  image_channel_type::fp16, Color);
882  break;
883  case image_channel_type::fp32:
884  convertReadData<cl_float>(readPixel(reinterpret_cast<cl_float *>(Ptr),
885  ImageChannelOrder, ImageChannelType),
886  image_channel_type::fp32, Color);
887  break;
888  }
889 
890  return Color;
891 }
892 
893 // Checks if the PixelCoord is out-of-range, and returns appropriate border or
894 // color value at the PixelCoord.
895 template <typename DataT>
896 DataT getColor(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode,
897  const range<3> ImgRange, const id<3> ImgPitch,
898  const image_channel_type ImgChannelType,
899  const image_channel_order ImgChannelOrder, void *BasePtr,
900  const uint8_t ElementSize) {
901  DataT RetData;
902  if (isOutOfRange(PixelCoord, SmplAddrMode, ImgRange)) {
903  cl_float4 BorderColor = getBorderColor(ImgChannelOrder);
904  RetData = BorderColor.convert<typename TryToGetElementType<DataT>::type>();
905  } else {
906  RetData = ReadPixelData<DataT>(PixelCoord, ImgPitch, ImgChannelType,
907  ImgChannelOrder, BasePtr, ElementSize);
908  }
909  return RetData;
910 }
911 
912 // Computes and returns color value with Linear Filter Mode.
913 // Steps:
914 // 1. Computes the 8 coordinates using all combinations of i0/i1,j0/j1,k0/k1.
915 // 2. Calls getColor() on each Coordinate.(Ci*j*k*)
916 // 3. Computes the return Color Value using a,b,c and the Color values.
917 template <typename DataT>
918 DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues,
919  const cl_float4 abc,
920  const addressing_mode SmplAddrMode,
921  const range<3> ImgRange, id<3> ImgPitch,
922  const image_channel_type ImgChannelType,
923  const image_channel_order ImgChannelOrder,
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();
927 
928  auto getColorInFloat = [&](cl_int4 V) {
929  DataT Res =
930  getColor<DataT>(V, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
931  ImgChannelOrder, BasePtr, ElementSize);
932  return Res.template convert<cl_float>();
933  };
934 
935  // Get Color Values at each Coordinate.
936  cl_float4 Ci0j0k0 = getColorInFloat(cl_int4{i0, j0, k0, 0});
937 
938  cl_float4 Ci1j0k0 = getColorInFloat(cl_int4{i1, j0, k0, 0});
939 
940  cl_float4 Ci0j1k0 = getColorInFloat(cl_int4{i0, j1, k0, 0});
941 
942  cl_float4 Ci1j1k0 = getColorInFloat(cl_int4{i1, j1, k0, 0});
943 
944  cl_float4 Ci0j0k1 = getColorInFloat(cl_int4{i0, j0, k1, 0});
945 
946  cl_float4 Ci1j0k1 = getColorInFloat(cl_int4{i1, j0, k1, 0});
947 
948  cl_float4 Ci0j1k1 = getColorInFloat(cl_int4{i0, j1, k1, 0});
949 
950  cl_float4 Ci1j1k1 = getColorInFloat(cl_int4{i1, j1, k1, 0});
951 
952  cl_float a = abc.x();
953  cl_float b = abc.y();
954  cl_float c = abc.z();
955 
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;
964 
965  cl_float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 +
966  Ci1j0k1 + Ci0j1k1 + Ci1j1k1;
967 
968  // For 2D image:k0 = 0, k1 = 0, c = 0.5
969  // RetData = (1 – a) * (1 – b) * Ci0j0 + a * (1 – b) * Ci1j0 +
970  // (1 – a) * b * Ci0j1 + a * b * Ci1j1;
971  // For 1D image: j0 = 0, j1 = 0, k0 = 0, k1 = 0, b = 0.5, c = 0.5.
972  // RetData = (1 – a) * Ci0 + a * Ci1;
973  return RetData.convert<typename TryToGetElementType<DataT>::type>();
974 }
975 
976 // imageReadSamplerHostImpl method is called by the read API in image accessors
977 // for host device.
978 // Algorithm used: The Algorithm is based on OpenCL spec section 8.2.
979 // It can be broken down into three major steps:
980 // Step 1.
981 // Check for valid sampler options and Compute u,v,w coordinates:
982 // These coordinates are used to compute the Pixel Coordinates that will be
983 // read from to compute the return values.
984 // u,v,w are normalized for AddrMode:mirror_repeat and repeat.
985 // u,v,w are unnormalized for AddrMode:clamp_to_edge, clamp, none.
986 // Convert normalized into unnormalized coords using image range.
987 // note: When dims=1, u,v,w={u,0,0}
988 // dims=2, u,v,w={u,v,0}
989 // dims=3, u,v,w-{u,v,w}
990 // Step 2.
991 // Process u,v,w, to find the exact Coordinates to read from:
992 // if(Nearest Filtering Mode)
993 // compute i,j,k pixel Coordinates based on AddrMode.
994 // else(Linear Filtering Mode)
995 // compute i0,j0,k0,i1,j1,k1,a,b,c values.
996 // Used to load following number of pixels in Step 3.
997 // 2x2x2 image for Dims=3
998 // 2x2 image for Dims=2
999 // 1 pixel for Dims=1 // I think same value should be
1000 // returned as nearest case.
1001 // Step 3.
1002 // Load Image Data, Different for Linear and Nearest Mode:
1003 // Offset = getOffset based on Coord, ImageRange,ImagePitch.
1004 // Read values in the appropriate format based on ImgChannelOrder and
1005 // ImgChannelType.
1006 // Convert to DataT as per conversion rules in section 8.3 in OpenCL Spec.
1007 //
1008 // TODO: Add additional check for half datatype read.
1009 // Based on OpenCL spec 2.0:
1010 // "The read_imageh calls that take integer coordinates must use a sampler with
1011 // filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to
1012 // CLK_NORMALIZED_COORDS_FALSE and addressing mode set to
1013 // CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise
1014 // the values returned are undefined."
1015 
1016 template <typename CoordT, typename DataT>
1017 DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,
1018  /*All image information*/ range<3> ImgRange,
1019  id<3> ImgPitch,
1020  image_channel_type ImgChannelType,
1021  image_channel_order ImgChannelOrder,
1022  void *BasePtr, uint8_t ElementSize) {
1023 
1024  coordinate_normalization_mode SmplNormMode =
1025  Smpl.get_coordinate_normalization_mode();
1026  addressing_mode SmplAddrMode = Smpl.get_addressing_mode();
1027  filtering_mode SmplFiltMode = Smpl.get_filtering_mode();
1028 
1029  CoordT Coorduvw;
1030  cl_float4 FloatCoorduvw;
1031  DataT RetData;
1032 
1033  // Step 1:
1034  // switch-case code is used for a better view on value of Coorduvw for all
1035  // combinations of Addressing Modes and Normalization Mode.
1036  switch (SmplNormMode) {
1037  case coordinate_normalization_mode::unnormalized:
1038  switch (SmplAddrMode) {
1039  case addressing_mode::mirrored_repeat:
1040  case addressing_mode::repeat:
1041  throw sycl::feature_not_supported(
1042  "Sampler used with unsupported configuration of "
1043  "mirrored_repeat/repeat filtering mode with unnormalized "
1044  "coordinates. ",
1045  PI_ERROR_INVALID_OPERATION);
1046  case addressing_mode::clamp_to_edge:
1048  case addressing_mode::none:
1049  // Continue with the unnormalized coordinates in Coorduvw.
1050  Coorduvw = Coords;
1051  break;
1052  }
1053  break; // Break for coordinate_normalization_mode::unnormalized.
1054  case coordinate_normalization_mode::normalized:
1055  switch (SmplAddrMode) {
1056  case addressing_mode::mirrored_repeat:
1057  case addressing_mode::repeat:
1058  // Continue with the normalized coordinates in Coorduvw.
1059  // Based on Section 8.2 Normalised coordinates are used to compute pixel
1060  // coordinates for addressing_mode::repeat and mirrored_repeat.
1061  Coorduvw = Coords;
1062  break;
1063  case addressing_mode::clamp_to_edge:
1065  case addressing_mode::none:
1066  // Unnormalize these coordinates.
1067  // Based on Section 8.2 Normalised coordinats are used to compute pixel
1068  // coordinates for addressing_mode::clamp/clamp_to_edge and none.
1069  Coorduvw = UnnormalizeCoordinates(Coords, ImgRange);
1070  break;
1071  }
1072  break; // Break for coordinate_normalization_mode::normalized.
1073  }
1074 
1075  // Step 2 & Step 3:
1076 
1077  // converToFloat4 converts CoordT of any kind - cl_int, cl_int2, cl_int4,
1078  // cl_float, cl_float2 and cl_float4 into Coordinates of kind cl_float4 with
1079  // no loss of precision. For pixel_coordinates already in cl_float4 format,
1080  // the function returns the same values. This conversion is done to enable
1081  // implementation of one common function getPixelCoordXXXMode, for any
1082  // datatype of CoordT passed.
1083  FloatCoorduvw = convertToFloat4(Coorduvw);
1084  switch (SmplFiltMode) {
1085  case filtering_mode::nearest: {
1086  // Get Pixel Coordinates in integers that will be read from in the Image.
1087  cl_int4 PixelCoord =
1088  getPixelCoordNearestFiltMode(FloatCoorduvw, SmplAddrMode, ImgRange);
1089 
1090  // Return Border Color for out-of-range coordinates when Sampler has
1091  // addressing_mode::clamp. For all other cases and for in-range coordinates
1092  // read the color and return in DataT type.
1093  RetData =
1094  getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1095  ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1096  break;
1097  }
1098  case filtering_mode::linear: {
1099  cl_float4 Retabc;
1100  // Get Pixel Coordinates in integers that will be read from in the Image.
1101  // Return i0,j0,k0,0,i1,j1,k1,0 to form 8 coordinates in a 3D image and
1102  // multiplication factors a,b,c
1103  cl_int8 CoordValues = getPixelCoordLinearFiltMode(
1104  FloatCoorduvw, SmplAddrMode, ImgRange, Retabc);
1105 
1106  // Find the 8 coordinates with the values in CoordValues.
1107  // Computes the Color Value to return.
1108  RetData = ReadPixelDataLinearFiltMode<DataT>(
1109  CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1110  ImgChannelOrder, BasePtr, ElementSize);
1111 
1112  break;
1113  }
1114  }
1115 
1116  return RetData;
1117 }
1118 
1119 } // namespace detail
1120 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1121 } // namespace sycl
1122 #endif
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: types.hpp:557
#define __SYCL_INLINE_VER_NAMESPACE(X)
vec< ChannelType, 4 > convertWriteData(const vec< cl_half, 4 > WriteData, const image_channel_type ImageChannelType)
bool isOutOfRange(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode, const range< 3 > ImgRange)
std::enable_if_t< std::is_integral< T >::value, size_t > getImageOffset(const vec< T, 4 > &Coords, const id< 3 > ImgPitch, const uint8_t ElementSize)
vec< ChannelType, 4 > processFloatDataToPixel(vec< cl_float, 4 > WriteData, float MulFactor)
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)
typename is_contained< T, type_list< cl_int, cl_float > >::type IsValidCoordType
vec< T, 4 > readPixel(T *Ptr, const image_channel_order ChannelOrder, const image_channel_type ChannelType)
std::enable_if_t< IsValidCoordType< T >::value, cl_float4 > convertToFloat4(vec< T, 4 > Coords)
cl_int4 getPixelCoordNearestFiltMode(cl_float4, const addressing_mode, const range< 3 >)
cl_float4 getBorderColor(const image_channel_order ImgChannelOrder)
DataT getColor(const cl_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 convertReadData(const vec< ChannelType, 4 > PixelData, const image_channel_type ImageChannelType, vec< cl_half, 4 > &RetData)
DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl, range< 3 > ImgRange, id< 3 > ImgPitch, image_channel_type ImgChannelType, image_channel_order ImgChannelOrder, void *BasePtr, uint8_t ElementSize)
DataT ReadPixelDataLinearFiltMode(const cl_int8 CoordValues, const cl_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)
cl_int8 getPixelCoordLinearFiltMode(cl_float4, const addressing_mode, const range< 3 >, cl_float4 &)
std::enable_if_t< IsValidCoordType< T >::value, vec< T, 4 > > UnnormalizeCoordinates(const vec< T, 4 > &Coords, const range< 3 > &Range)
DataT ReadPixelData(const cl_int4 PixelCoord, const id< 3 > ImgPitch, const image_channel_type ImageChannelType, const image_channel_order ImageChannelOrder, void *BasePtr, const uint8_t ElementSize)
filtering_mode
Definition: sampler.hpp:28
std::int8_t cl_char
Definition: aliases.hpp:79
half cl_half
Definition: aliases.hpp:87
std::int16_t cl_short
Definition: aliases.hpp:81
std::uint8_t cl_uchar
Definition: aliases.hpp:80
coordinate_normalization_mode
Definition: sampler.hpp:33
std::uint16_t cl_ushort
Definition: aliases.hpp:82
std::int32_t cl_int
Definition: aliases.hpp:83
std::uint32_t cl_uint
Definition: aliases.hpp:84
addressing_mode
Definition: sampler.hpp:20
image_channel_order
Definition: image.hpp:27
float cl_float
Definition: aliases.hpp:88
image_channel_type
Definition: image.hpp:45
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
simd< _Tp, _Abi > clamp(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &)