DPC++ Runtime
Runtime libraries for oneAPI DPC++
image_accessor_util.cpp
Go to the documentation of this file.
1 //==----------- image_accessor_util.cpp ------------------------------------==//
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 
9 #include <CL/sycl/accessor.hpp>
10 #include <CL/sycl/builtins.hpp>
11 
13 namespace sycl {
14 namespace detail {
15 
16 // For Nearest Filtering mode, process cl_float4 Coordinates and return the
17 // appropriate Pixel Coordinates based on Addressing Mode.
18 cl_int4 getPixelCoordNearestFiltMode(cl_float4 Coorduvw,
19  const addressing_mode SmplAddrMode,
20  const range<3> ImgRange) {
21  cl_int4 Coordijk(0);
22  cl_int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0);
23  switch (SmplAddrMode) {
24  case addressing_mode::mirrored_repeat: {
25  cl_float4 Tempuvw(0);
26  Tempuvw = 2.0f * cl::sycl::rint(0.5f * Coorduvw);
27  Tempuvw = cl::sycl::fabs(Coorduvw - Tempuvw);
28  Tempuvw = Tempuvw * (Rangewhd.convert<cl_float>());
29  Tempuvw = (cl::sycl::floor(Tempuvw));
30  Coordijk = Tempuvw.convert<cl_int>();
31  Coordijk = cl::sycl::min(Coordijk, (Rangewhd - 1));
32  // Eg:
33  // u,v,w = {2.3,1.7,0.5} // normalized coordinates.
34  // w,h,d = {9,9,9}
35  // u1=2*rint(1.15)=2
36  // v1=2*rint(0.85)=2
37  // w1=2*rint(0.5)=0
38  // u1=fabs(2.3-2)=.3
39  // v1=fabs(1.7-2)=.3
40  // w1=fabs(0.5-0)=.5
41  // u1=0.3*9=2.7
42  // v1=0.3*9=2.7
43  // w1=0.5*9=4.5
44  // i,j,k = {2,2,4}
45 
46  } break;
47  case addressing_mode::repeat: {
48 
49  cl_float4 Tempuvw(0);
50  Tempuvw =
51  (Coorduvw - cl::sycl::floor(Coorduvw)) * Rangewhd.convert<cl_float>();
52  Coordijk = (cl::sycl::floor(Tempuvw)).convert<cl_int>();
53  cl_int4 GreaterThanEqual = (Coordijk >= Rangewhd);
54  Coordijk =
55  cl::sycl::select(Coordijk, (Coordijk - Rangewhd), GreaterThanEqual);
56  // Eg:
57  // u = 2.3; v = 1.5; w = 0.5; // normalized coordinates.
58  // w,h,d = {9,9,9};
59  // u1= 0.3*w;
60  // v1= 0.5*d;
61  // w1= 0.5*h;
62  // i = floor(2.7);
63  // j = floor(4.5);
64  // k = floor(4.5);
65  // if (i/j/k > w/h/d-1)
66  // // Condition is not satisfied.
67  // (This condition I think will only be satisfied if the floating point
68  // arithmetic of multiplication
69  // gave a value in u1/v1/w1 as > w/h/d)
70  // i = 2; j = 4; k = 4;
71  } break;
72  case addressing_mode::clamp_to_edge:
73  Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
74  Coordijk = cl::sycl::clamp(Coordijk, cl_int4(0), (Rangewhd - 1));
75  break;
77  Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
78  Coordijk = cl::sycl::clamp(Coordijk, cl_int4(-1), Rangewhd);
79  break;
80  case addressing_mode::none:
81  Coordijk = (cl::sycl::floor(Coorduvw)).convert<cl_int>();
82  break;
83  }
84  return Coordijk;
85 }
86 
87 // For Linear Filtering mode, process Coordinates-Coord_uvw and return
88 // coordinate indexes based on Addressing Mode.
89 // The value returned contains (i0,j0,k0,0,i1,j1,k1,0).
90 // Retabc contains the values of (a,b,c,0)
91 // The caller of this function should use these values to create the 8 pixel
92 // coordinates and multiplication coefficients.
93 cl_int8 getPixelCoordLinearFiltMode(cl_float4 Coorduvw,
94  const addressing_mode SmplAddrMode,
95  const range<3> ImgRange, cl_float4 &Retabc) {
96  cl_int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0);
97  cl_int4 Ci0j0k0(0);
98  cl_int4 Ci1j1k1(0);
99  cl_int4 Int_uvwsubhalf = cl::sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
100 
101  switch (SmplAddrMode) {
102  case addressing_mode::mirrored_repeat: {
103  cl_float4 Temp;
104  Temp = (cl::sycl::rint(Coorduvw * 0.5f)) * 2.0f;
105  Temp = cl::sycl::fabs(Coorduvw - Temp);
106  Coorduvw = Temp * Rangewhd.convert<cl_float>();
107  Int_uvwsubhalf = cl::sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
108 
109  Ci0j0k0 = Int_uvwsubhalf;
110  Ci1j1k1 = Ci0j0k0 + 1;
111 
112  Ci0j0k0 = cl::sycl::max(Ci0j0k0, 0);
113  Ci1j1k1 = cl::sycl::min(Ci1j1k1, (Rangewhd - 1));
114  } break;
115  case addressing_mode::repeat: {
116 
117  Coorduvw =
118  (Coorduvw - cl::sycl::floor(Coorduvw)) * Rangewhd.convert<cl_float>();
119  Int_uvwsubhalf = cl::sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
120 
121  Ci0j0k0 = Int_uvwsubhalf;
122  Ci1j1k1 = Ci0j0k0 + 1;
123 
124  Ci0j0k0 =
125  cl::sycl::select(Ci0j0k0, (Ci0j0k0 + Rangewhd), Ci0j0k0 < cl_int4(0));
126  Ci1j1k1 =
127  cl::sycl::select(Ci1j1k1, (Ci1j1k1 - Rangewhd), Ci1j1k1 >= Rangewhd);
128 
129  } break;
130  case addressing_mode::clamp_to_edge: {
131  Ci0j0k0 = cl::sycl::clamp(Int_uvwsubhalf, cl_int4(0), (Rangewhd - 1));
132  Ci1j1k1 = cl::sycl::clamp((Int_uvwsubhalf + 1), cl_int4(0), (Rangewhd - 1));
133  break;
134  }
135  case addressing_mode::clamp: {
136  Ci0j0k0 = cl::sycl::clamp(Int_uvwsubhalf, cl_int4(-1), Rangewhd);
137  Ci1j1k1 = cl::sycl::clamp((Int_uvwsubhalf + 1), cl_int4(-1), Rangewhd);
138  break;
139  }
140  case addressing_mode::none: {
141  Ci0j0k0 = Int_uvwsubhalf;
142  Ci1j1k1 = Ci0j0k0 + 1;
143  break;
144  }
145  }
146  Retabc = (Coorduvw - 0.5f) - (Int_uvwsubhalf.convert<cl_float>());
147  Retabc.w() = 0.0f;
148  return cl_int8(Ci0j0k0, Ci1j1k1);
149 }
150 
151 // Function returns true when PixelCoord is out of image's range.
152 // It is only valid for addressing_mode::clamp/none.
153 // Note: For addressing_mode::none , spec says outofrange access is not defined.
154 // This function handles this addressing_mode to avoid accessing out of bound
155 // memories on host.
156 bool isOutOfRange(const cl_int4 PixelCoord, const addressing_mode SmplAddrMode,
157  const range<3> ImgRange) {
158 
159  if (SmplAddrMode != addressing_mode::clamp &&
160  SmplAddrMode != addressing_mode::none)
161  return false;
162 
163  auto CheckOutOfRange = [](cl_int Coord, cl_int Range) {
164  return ((Coord < 0) || (Coord >= Range));
165  };
166 
167  bool CheckWidth = CheckOutOfRange(PixelCoord.x(), ImgRange[0]);
168  bool CheckHeight = CheckOutOfRange(PixelCoord.y(), ImgRange[1]);
169  bool CheckDepth = CheckOutOfRange(PixelCoord.z(), ImgRange[2]);
170 
171  return (CheckWidth || CheckHeight || CheckDepth);
172 }
173 
174 cl_float4 getBorderColor(const image_channel_order ImgChannelOrder) {
175 
176  cl_float4 BorderColor(0.0f);
177  switch (ImgChannelOrder) {
178  case image_channel_order::r:
179  case image_channel_order::rg:
180  case image_channel_order::rgb:
181  case image_channel_order::luminance:
182  BorderColor.w() = 1.0f;
183  break;
184  default:
185  break;
186  }
187  return BorderColor;
188 }
189 
190 } // namespace detail
191 } // namespace sycl
192 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::select
detail::enable_if_t< detail::is_sgentype< T >::value, T > select(T a, T b, bool c) __NOEXC
Definition: builtins.hpp:1318
cl::sycl::detail::getPixelCoordNearestFiltMode
cl_int4 getPixelCoordNearestFiltMode(cl_float4, const addressing_mode, const range< 3 >)
Definition: image_accessor_util.cpp:18
cl::sycl::floor
detail::enable_if_t< detail::is_genfloat< T >::value, T > floor(T x) __NOEXC
Definition: builtins.hpp:190
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::getPixelCoordLinearFiltMode
cl_int8 getPixelCoordLinearFiltMode(cl_float4, const addressing_mode, const range< 3 >, cl_float4 &)
Definition: image_accessor_util.cpp:93
cl::sycl::fabs
detail::enable_if_t< detail::is_genfloat< T >::value, T > fabs(T x) __NOEXC
Definition: builtins.hpp:178
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::cl_float
float cl_float
Definition: aliases.hpp:88
clamp
simd< _Tp, _Abi > clamp(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &)
accessor.hpp
builtins.hpp
cl::sycl::detail::getBorderColor
cl_float4 getBorderColor(const image_channel_order ImgChannelOrder)
Definition: image_accessor_util.cpp:174
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:83
cl::sycl::addressing_mode
addressing_mode
Definition: sampler.hpp:20
cl::sycl::rint
detail::enable_if_t< detail::is_genfloat< T >::value, T > rint(T x) __NOEXC
Definition: builtins.hpp:415
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