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 <CL/sycl/builtins.hpp>
18 #include <CL/sycl/image.hpp>
19 #include <CL/sycl/sampler.hpp>
20 #include <CL/sycl/types.hpp>
21 
22 #include <cmath>
23 #include <iostream>
24 
26 namespace sycl {
27 namespace detail {
28 
29 template <typename T>
30 using IsValidCoordType =
32 
33 // The formula for unnormalization coordinates:
34 // NormalizedCoords = [UnnormalizedCoords[i] * Range[i] for i in range(0, 3)]
35 template <typename T>
37 UnnormalizeCoordinates(const T &Coords, const range<3> &Range) {
38  return Coords * Range[0];
39 }
40 
41 template <typename T>
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>
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>
60 convertToFloat4(T Coords) {
61  return {static_cast<float>(Coords), 0.5f, 0.5f, 0.f};
62 }
63 
64 template <typename T>
67  return {static_cast<float>(Coords.x()), static_cast<float>(Coords.y()), 0.5f,
68  0.f};
69 }
70 
71 template <typename T>
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>
83 getImageOffset(const T &Coords, const id<3>, const uint8_t ElementSize) {
84  return Coords * ElementSize;
85 }
86 
87 template <typename 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>
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 cl::sycl::invalid_parameter_error(
299  "Datatype of read data - cl_uint4 is incompatible with the "
300  "image_channel_type of the image.",
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 cl::sycl::invalid_parameter_error(
321  "Datatype of read data - cl_int4 is incompatible with the "
322  "image_channel_type of the image.",
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 = cl::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 = cl::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 cl::sycl::invalid_parameter_error(
403  "Datatype of read data - cl_float4 is incompatible with the "
404  "image_channel_type of the image.",
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 = cl::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 = cl::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 cl::sycl::feature_not_supported(
446  "Currently unsupported datatype conversion from image_channel_type "
447  "to cl_half4.",
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 cl::sycl::invalid_parameter_error(
459  "Datatype to read- cl_half4 is incompatible with the "
460  "image_channel_type of the image.",
462  case image_channel_type::fp16:
463  RetData = PixelData.template convert<cl_half>();
464  return;
465  case image_channel_type::fp32:
466  throw cl::sycl::invalid_parameter_error(
467  "Datatype to read - cl_half4 is incompatible with the "
468  "image_channel_type of the image.",
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 = cl::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 = cl::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 cl::sycl::invalid_parameter_error(
507  "Datatype of data to write - cl_uint4 is incompatible with the "
508  "image_channel_type of the image.",
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 = cl::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 = cl::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 cl::sycl::invalid_parameter_error(
540  "Datatype of data to write - cl_int4 is incompatible with the "
541  "image_channel_type of the image.",
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  cl::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 cl::sycl::feature_not_supported(
579  "Currently unsupported datatype conversion from image_channel_type "
580  "to cl_float4.",
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 = cl::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 = cl::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 cl::sycl::invalid_parameter_error(
622  "Datatype of data to write - cl_float4 is incompatible with the "
623  "image_channel_type of the image.",
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 cl::sycl::feature_not_supported(
657  "Currently unsupported datatype conversion from image_channel_type "
658  "to cl_half4.",
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 cl::sycl::invalid_parameter_error(
670  "Datatype of data to write - cl_float4 is incompatible with the "
671  "image_channel_type of the image.",
673  case image_channel_type::fp16:
674  return WriteData.convert<ChannelType>();
675  case image_channel_type::fp32:
676  throw cl::sycl::invalid_parameter_error(
677  "Datatype of data to write - cl_float4 is incompatible with the "
678  "image_channel_type of the image.",
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,
803  void *BasePtr, 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 =
929  [&](cl_int4 V) {
930  DataT Res = getColor<DataT>(V, SmplAddrMode,
931  ImgRange, ImgPitch, ImgChannelType,
932  ImgChannelOrder, BasePtr, ElementSize);
933  return Res.template convert<cl_float>();
934  };
935 
936  // Get Color Values at each Coordinate.
937  cl_float4 Ci0j0k0 = getColorInFloat(cl_int4{i0, j0, k0, 0});
938 
939  cl_float4 Ci1j0k0 = getColorInFloat(cl_int4{i1, j0, k0, 0});
940 
941  cl_float4 Ci0j1k0 = getColorInFloat(cl_int4{i0, j1, k0, 0});
942 
943  cl_float4 Ci1j1k0 = getColorInFloat(cl_int4{i1, j1, k0, 0});
944 
945  cl_float4 Ci0j0k1 = getColorInFloat(cl_int4{i0, j0, k1, 0});
946 
947  cl_float4 Ci1j0k1 = getColorInFloat(cl_int4{i1, j0, k1, 0});
948 
949  cl_float4 Ci0j1k1 = getColorInFloat(cl_int4{i0, j1, k1, 0});
950 
951  cl_float4 Ci1j1k1 = getColorInFloat(cl_int4{i1, j1, k1, 0});
952 
953  cl_float a = abc.x();
954  cl_float b = abc.y();
955  cl_float c = abc.z();
956 
957  Ci0j0k0 = (1 - a) * (1 - b) * (1 - c) * Ci0j0k0;
958  Ci1j0k0 = a * (1 - b) * (1 - c) * Ci1j0k0;
959  Ci0j1k0 = (1 - a) * b * (1 - c) * Ci0j1k0;
960  Ci1j1k0 = a * b * (1 - c) * Ci1j1k0;
961  Ci0j0k1 = (1 - a) * (1 - b) * c * Ci0j0k1;
962  Ci1j0k1 = a * (1 - b) * c * Ci1j0k1;
963  Ci0j1k1 = (1 - a) * b * c * Ci0j1k1;
964  Ci1j1k1 = a * b * c * Ci1j1k1;
965 
966  cl_float4 RetData = Ci0j0k0 + Ci1j0k0 + Ci0j1k0 + Ci1j1k0 + Ci0j0k1 +
967  Ci1j0k1 + Ci0j1k1 + Ci1j1k1;
968 
969  // For 2D image:k0 = 0, k1 = 0, c = 0.5
970  // RetData = (1 – a) * (1 – b) * Ci0j0 + a * (1 – b) * Ci1j0 +
971  // (1 – a) * b * Ci0j1 + a * b * Ci1j1;
972  // For 1D image: j0 = 0, j1 = 0, k0 = 0, k1 = 0, b = 0.5, c = 0.5.
973  // RetData = (1 – a) * Ci0 + a * Ci1;
974  return RetData.convert<typename TryToGetElementType<DataT>::type>();
975 }
976 
977 // imageReadSamplerHostImpl method is called by the read API in image accessors
978 // for host device.
979 // Algorithm used: The Algorithm is based on OpenCL spec section 8.2.
980 // It can be broken down into three major steps:
981 // Step 1.
982 // Check for valid sampler options and Compute u,v,w coordinates:
983 // These coordinates are used to compute the Pixel Coordinates that will be
984 // read from to compute the return values.
985 // u,v,w are normalized for AddrMode:mirror_repeat and repeat.
986 // u,v,w are unnormalized for AddrMode:clamp_to_edge, clamp, none.
987 // Convert normalized into unnormalized coords using image range.
988 // note: When dims=1, u,v,w={u,0,0}
989 // dims=2, u,v,w={u,v,0}
990 // dims=3, u,v,w-{u,v,w}
991 // Step 2.
992 // Process u,v,w, to find the exact Coordinates to read from:
993 // if(Nearest Filtering Mode)
994 // compute i,j,k pixel Coordinates based on AddrMode.
995 // else(Linear Filtering Mode)
996 // compute i0,j0,k0,i1,j1,k1,a,b,c values.
997 // Used to load following number of pixels in Step 3.
998 // 2x2x2 image for Dims=3
999 // 2x2 image for Dims=2
1000 // 1 pixel for Dims=1 // I think same value should be
1001 // returned as nearest case.
1002 // Step 3.
1003 // Load Image Data, Different for Linear and Nearest Mode:
1004 // Offset = getOffset based on Coord, ImageRange,ImagePitch.
1005 // Read values in the appropriate format based on ImgChannelOrder and
1006 // ImgChannelType.
1007 // Convert to DataT as per conversion rules in section 8.3 in OpenCL Spec.
1008 //
1009 // TODO: Add additional check for half datatype read.
1010 // Based on OpenCL spec 2.0:
1011 // "The read_imageh calls that take integer coordinates must use a sampler with
1012 // filter mode set to CLK_FILTER_NEAREST, normalized coordinates set to
1013 // CLK_NORMALIZED_COORDS_FALSE and addressing mode set to
1014 // CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP or CLK_ADDRESS_NONE; otherwise
1015 // the values returned are undefined."
1016 
1017 template <typename CoordT, typename DataT>
1018 DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,
1019  /*All image information*/ range<3> ImgRange,
1020  id<3> ImgPitch,
1021  image_channel_type ImgChannelType,
1022  image_channel_order ImgChannelOrder,
1023  void *BasePtr, uint8_t ElementSize) {
1024 
1025  coordinate_normalization_mode SmplNormMode =
1027  addressing_mode SmplAddrMode = Smpl.get_addressing_mode();
1028  filtering_mode SmplFiltMode = Smpl.get_filtering_mode();
1029 
1030  CoordT Coorduvw;
1031  cl_float4 FloatCoorduvw;
1032  DataT RetData;
1033 
1034  // Step 1:
1035  // switch-case code is used for a better view on value of Coorduvw for all
1036  // combinations of Addressing Modes and Normalization Mode.
1037  switch (SmplNormMode) {
1038  case coordinate_normalization_mode::unnormalized:
1039  switch (SmplAddrMode) {
1040  case addressing_mode::mirrored_repeat:
1041  case addressing_mode::repeat:
1042  throw cl::sycl::feature_not_supported(
1043  "Sampler used with unsupported configuration of "
1044  "mirrored_repeat/repeat filtering mode with unnormalized "
1045  "coordinates. ",
1047  case addressing_mode::clamp_to_edge:
1049  case addressing_mode::none:
1050  // Continue with the unnormalized coordinates in Coorduvw.
1051  Coorduvw = Coords;
1052  break;
1053  }
1054  break; // Break for coordinate_normalization_mode::unnormalized.
1055  case coordinate_normalization_mode::normalized:
1056  switch (SmplAddrMode) {
1057  case addressing_mode::mirrored_repeat:
1058  case addressing_mode::repeat:
1059  // Continue with the normalized coordinates in Coorduvw.
1060  // Based on Section 8.2 Normalised coordinates are used to compute pixel
1061  // coordinates for addressing_mode::repeat and mirrored_repeat.
1062  Coorduvw = Coords;
1063  break;
1064  case addressing_mode::clamp_to_edge:
1066  case addressing_mode::none:
1067  // Unnormalize these coordinates.
1068  // Based on Section 8.2 Normalised coordinats are used to compute pixel
1069  // coordinates for addressing_mode::clamp/clamp_to_edge and none.
1070  Coorduvw = UnnormalizeCoordinates(Coords, ImgRange);
1071  break;
1072  }
1073  break; // Break for coordinate_normalization_mode::normalized.
1074  }
1075 
1076  // Step 2 & Step 3:
1077 
1078  // converToFloat4 converts CoordT of any kind - cl_int, cl_int2, cl_int4,
1079  // cl_float, cl_float2 and cl_float4 into Coordinates of kind cl_float4 with
1080  // no loss of precision. For pixel_coordinates already in cl_float4 format,
1081  // the function returns the same values. This conversion is done to enable
1082  // implementation of one common function getPixelCoordXXXMode, for any
1083  // datatype of CoordT passed.
1084  FloatCoorduvw = convertToFloat4(Coorduvw);
1085  switch (SmplFiltMode) {
1086  case filtering_mode::nearest: {
1087  // Get Pixel Coordinates in integers that will be read from in the Image.
1088  cl_int4 PixelCoord =
1089  getPixelCoordNearestFiltMode(FloatCoorduvw, SmplAddrMode, ImgRange);
1090 
1091  // Return Border Color for out-of-range coordinates when Sampler has
1092  // addressing_mode::clamp. For all other cases and for in-range coordinates
1093  // read the color and return in DataT type.
1094  RetData =
1095  getColor<DataT>(PixelCoord, SmplAddrMode, ImgRange, ImgPitch,
1096  ImgChannelType, ImgChannelOrder, BasePtr, ElementSize);
1097  break;
1098  }
1099  case filtering_mode::linear: {
1100  cl_float4 Retabc;
1101  // Get Pixel Coordinates in integers that will be read from in the Image.
1102  // Return i0,j0,k0,0,i1,j1,k1,0 to form 8 coordinates in a 3D image and
1103  // multiplication factors a,b,c
1104  cl_int8 CoordValues = getPixelCoordLinearFiltMode(
1105  FloatCoorduvw, SmplAddrMode, ImgRange, Retabc);
1106 
1107  // Find the 8 coordinates with the values in CoordValues.
1108  // Computes the Color Value to return.
1109  RetData = ReadPixelDataLinearFiltMode<DataT>(
1110  CoordValues, Retabc, SmplAddrMode, ImgRange, ImgPitch, ImgChannelType,
1111  ImgChannelOrder, BasePtr, ElementSize);
1112 
1113  break;
1114  }
1115  }
1116 
1117  return RetData;
1118 }
1119 
1120 } // namespace detail
1121 } // namespace sycl
1122 } // __SYCL_INLINE_NAMESPACE(cl)
1123 #endif
cl::sycl::cl_uchar
std::uint8_t cl_uchar
Definition: aliases.hpp:79
cl::sycl::detail::getColor
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)
Definition: image_accessor_util.hpp:896
cl::sycl::detail::getPixelCoordNearestFiltMode
cl_int4 getPixelCoordNearestFiltMode(cl_float4, const addressing_mode, const range< 3 >)
Definition: image_accessor_util.cpp:18
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:88
cl::sycl::detail::convertReadData
void convertReadData(const vec< ChannelType, 4 > PixelData, const image_channel_type ImageChannelType, vec< cl_half, 4 > &RetData)
Definition: image_accessor_util.hpp:418
cl::sycl::detail::UnnormalizeCoordinates
detail::enable_if_t< IsValidCoordType< T >::value, vec< T, 4 > > UnnormalizeCoordinates(const vec< T, 4 > &Coords, const range< 3 > &Range)
Definition: image_accessor_util.hpp:49
cl::sycl::id< 3 >
cl::sycl::detail::convertToFloat4
detail::enable_if_t< IsValidCoordType< T >::value, cl_float4 > convertToFloat4(vec< T, 4 > Coords)
Definition: image_accessor_util.hpp:73
cl::sycl::detail::readPixel
vec< T, 4 > readPixel(T *Ptr, const image_channel_order ChannelOrder, const image_channel_type ChannelType)
Definition: image_accessor_util.hpp:129
cl::sycl::sampler
Encapsulates a configuration for sampling an image accessor.
Definition: sampler.hpp:65
cl::sycl::detail::convertWriteData
vec< ChannelType, 4 > convertWriteData(const vec< cl_half, 4 > WriteData, const image_channel_type ImageChannelType)
Definition: image_accessor_util.hpp:636
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
cl::sycl::detail::IsValidCoordType
typename is_contained< T, type_list< cl_int, cl_float > >::type IsValidCoordType
Definition: image_accessor_util.hpp:31
sycl
Definition: invoke_simd.hpp:68
cl::sycl::range< 3 >
cl::sycl::detail::isOutOfRange
bool isOutOfRange(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode, const range< 3 > ImgRange)
Definition: image_accessor_util.cpp:156
cl::sycl::detail::getImageOffset
detail::enable_if_t< std::is_integral< T >::value, size_t > getImageOffset(const vec< T, 4 > &Coords, const id< 3 > ImgPitch, const uint8_t ElementSize)
Definition: image_accessor_util.hpp:96
cl::sycl::detail::TryToGetElementType::type
decltype(check(T())) type
Definition: generic_type_traits.hpp:308
cl::sycl::detail::ReadPixelData
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)
Definition: image_accessor_util.hpp:800
export.hpp
cl::sycl::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:212
cl::sycl::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:698
cl::sycl::detail::getPixelCoordLinearFiltMode
cl_int8 getPixelCoordLinearFiltMode(cl_float4, const addressing_mode, const range< 3 >, cl_float4 &)
Definition: image_accessor_util.cpp:93
cl::sycl::sampler::get_addressing_mode
addressing_mode get_addressing_mode() const
Definition: sampler.cpp:25
cl::sycl::cl_ushort
std::uint16_t cl_ushort
Definition: aliases.hpp:81
cl::sycl::detail::ReadPixelDataLinearFiltMode
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)
Definition: image_accessor_util.hpp:918
cl::sycl::detail::half_impl::half
Definition: half_type.hpp:329
generic_type_traits.hpp
cl::sycl::fmax
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmax(T x, T y) __NOEXC
Definition: builtins.hpp:203
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::cl_char
std::int8_t cl_char
Definition: aliases.hpp:78
cl::sycl::cl_float
float cl_float
Definition: aliases.hpp:87
image.hpp
clamp
simd< _Tp, _Abi > clamp(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &)
cl::sycl::image_channel_order::a
@ a
builtins.hpp
cl::sycl::detail::imageReadSamplerHostImpl
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)
Definition: image_accessor_util.hpp:1018
PI_INVALID_VALUE
@ PI_INVALID_VALUE
Definition: pi.h:91
cl::sycl::detail::getBorderColor
cl_float4 getBorderColor(const image_channel_order ImgChannelOrder)
Definition: image_accessor_util.cpp:174
cl::sycl::coordinate_normalization_mode
coordinate_normalization_mode
Definition: sampler.hpp:32
cl::sycl::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:18
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:82
cl::sycl::addressing_mode
addressing_mode
Definition: sampler.hpp:19
cl::sycl::detail::processFloatDataToPixel
vec< ChannelType, 4 > processFloatDataToPixel(vec< cl_float, 4 > WriteData, float MulFactor)
Definition: image_accessor_util.hpp:547
sampler.hpp
cl::sycl::sampler::get_coordinate_normalization_mode
coordinate_normalization_mode get_coordinate_normalization_mode() const
Definition: sampler.cpp:34
cl::sycl::cl_uint
std::uint32_t cl_uint
Definition: aliases.hpp:83
cl::sycl::filtering_mode
filtering_mode
Definition: sampler.hpp:27
cl::sycl::image_channel_type
image_channel_type
Definition: image.hpp:41
cl::sycl::cl_short
std::int16_t cl_short
Definition: aliases.hpp:80
types.hpp
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::sampler::get_filtering_mode
filtering_mode get_filtering_mode() const
Definition: sampler.cpp:29
cl::sycl::image_channel_order
image_channel_order
Definition: image.hpp:23
cl::sycl::clamp
detail::enable_if_t< detail::is_genfloat< T >::value, T > clamp(T x, T minval, T maxval) __NOEXC
Definition: builtins.hpp:506
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12