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