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