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 {
13 inline namespace _V1 {
14 namespace detail {
15 
16 // For Nearest Filtering mode, process float4 Coordinates and
17 // return the appropriate Pixel Coordinates based on Addressing Mode.
18 int4 getPixelCoordNearestFiltMode(float4 Coorduvw,
19  const addressing_mode SmplAddrMode,
20  const range<3> ImgRange) {
21  int4 Coordijk(0);
22  int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0);
23  switch (SmplAddrMode) {
25  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;
48 
49  float4 Tempuvw(0);
50  Tempuvw = (Coorduvw - sycl::floor(Coorduvw)) * Rangewhd.convert<cl_float>();
51  Coordijk = (sycl::floor(Tempuvw)).convert<cl_int>();
52  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;
71  Coordijk = (sycl::floor(Coorduvw)).convert<cl_int>();
72  Coordijk = sycl::clamp(Coordijk, int4(0), (Rangewhd - 1));
73  break;
75  Coordijk = (sycl::floor(Coorduvw)).convert<cl_int>();
76  Coordijk = sycl::clamp(Coordijk, int4(-1), Rangewhd);
77  break;
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 int8 getPixelCoordLinearFiltMode(float4 Coorduvw,
92  const addressing_mode SmplAddrMode,
93  const range<3> ImgRange, float4 &Retabc) {
94  int4 Rangewhd(ImgRange[0], ImgRange[1], ImgRange[2], 0);
95  int4 Ci0j0k0(0);
96  int4 Ci1j1k1(0);
97  int4 Int_uvwsubhalf = sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
98 
99  switch (SmplAddrMode) {
101  float4 Temp;
102  Temp = (sycl::rint(Coorduvw * 0.5f)) * 2.0f;
103  Temp = sycl::fabs(Coorduvw - Temp);
104  Coorduvw = Temp * Rangewhd.convert<cl_float>();
105  Int_uvwsubhalf = sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
106 
107  Ci0j0k0 = Int_uvwsubhalf;
108  Ci1j1k1 = Ci0j0k0 + 1;
109 
110  Ci0j0k0 = sycl::max(Ci0j0k0, 0);
111  Ci1j1k1 = sycl::min(Ci1j1k1, (Rangewhd - 1));
112  } break;
114 
115  Coorduvw =
116  (Coorduvw - sycl::floor(Coorduvw)) * Rangewhd.convert<cl_float>();
117  Int_uvwsubhalf = sycl::floor(Coorduvw - 0.5f).convert<cl_int>();
118 
119  Ci0j0k0 = Int_uvwsubhalf;
120  Ci1j1k1 = Ci0j0k0 + 1;
121 
122  Ci0j0k0 = sycl::select(Ci0j0k0, (Ci0j0k0 + Rangewhd), Ci0j0k0 < int4(0));
123  Ci1j1k1 = sycl::select(Ci1j1k1, (Ci1j1k1 - Rangewhd), Ci1j1k1 >= Rangewhd);
124 
125  } break;
127  Ci0j0k0 = sycl::clamp(Int_uvwsubhalf, int4(0), (Rangewhd - 1));
128  Ci1j1k1 = sycl::clamp((Int_uvwsubhalf + 1), int4(0), (Rangewhd - 1));
129  break;
130  }
131  case addressing_mode::clamp: {
132  Ci0j0k0 = sycl::clamp(Int_uvwsubhalf, int4(-1), Rangewhd);
133  Ci1j1k1 = sycl::clamp((Int_uvwsubhalf + 1), int4(-1), Rangewhd);
134  break;
135  }
136  case addressing_mode::none: {
137  Ci0j0k0 = Int_uvwsubhalf;
138  Ci1j1k1 = Ci0j0k0 + 1;
139  break;
140  }
141  }
142  Retabc = (Coorduvw - 0.5f) - (Int_uvwsubhalf.convert<cl_float>());
143  Retabc.w() = 0.0f;
144  return int8(Ci0j0k0, Ci1j1k1);
145 }
146 
147 // Function returns true when PixelCoord is out of image's range.
148 // It is only valid for addressing_mode::clamp/none.
149 // Note: For addressing_mode::none , spec says outofrange access is not defined.
150 // This function handles this addressing_mode to avoid accessing out of bound
151 // memories on host.
152 bool isOutOfRange(const int4 PixelCoord, const addressing_mode SmplAddrMode,
153  const range<3> ImgRange) {
154 
155  if (SmplAddrMode != addressing_mode::clamp &&
156  SmplAddrMode != addressing_mode::none)
157  return false;
158 
159  auto CheckOutOfRange = [](cl_int Coord, cl_int Range) {
160  return ((Coord < 0) || (Coord >= Range));
161  };
162 
163  bool CheckWidth = CheckOutOfRange(PixelCoord.x(), ImgRange[0]);
164  bool CheckHeight = CheckOutOfRange(PixelCoord.y(), ImgRange[1]);
165  bool CheckDepth = CheckOutOfRange(PixelCoord.z(), ImgRange[2]);
166 
167  return (CheckWidth || CheckHeight || CheckDepth);
168 }
169 
170 float4 getBorderColor(const image_channel_order ImgChannelOrder) {
171 
172  float4 BorderColor(0.0f);
173  switch (ImgChannelOrder) {
178  BorderColor.w() = 1.0f;
179  break;
180  default:
181  break;
182  }
183  return BorderColor;
184 }
185 
186 } // namespace detail
187 } // namespace _V1
188 } // namespace sycl
The file contains implementations of accessor class.
float4 getBorderColor(const image_channel_order ImgChannelOrder)
int4 getPixelCoordNearestFiltMode(float4, const addressing_mode, const range< 3 >)
bool isOutOfRange(const int4 PixelCoord, const addressing_mode SmplAddrMode, const range< 3 > ImgRange)
int8 getPixelCoordLinearFiltMode(float4, const addressing_mode, const range< 3 >, float4 &)
std::int32_t cl_int
Definition: aliases.hpp:134
std::enable_if_t< detail::is_gentype_v< T >, marray< T, N > > select(marray< T, N > a, marray< T, N > b, marray< bool, N > c)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > rint(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > fabs(T x)
addressing_mode
Definition: sampler.hpp:23
image_channel_order
Definition: image.hpp:56
std::enable_if_t< detail::is_svgenfloat_v< T >, T > floor(T x)
Definition: access.hpp:18