Renesas GR-PEACH OpenCV Development
/
gr-peach-opencv-project-sd-card_update
Important changes to repositories hosted on mbed.com
Mbed hosted mercurial repositories are deprecated and are due to be permanently deleted in July 2026.
To keep a copy of this software download the repository Zip archive or clone locally using Mercurial.
It is also possible to export all your personal repositories from the account settings page.
Fork of
gr-peach-opencv-project-sd-card
by 
the do
Embed:
(wiki syntax)
Show/hide line numbers
border_interpolate.hpp
Go to the documentation of this file.
00001 /*M/////////////////////////////////////////////////////////////////////////////////////// 00002 // 00003 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 00004 // 00005 // By downloading, copying, installing or using the software you agree to this license. 00006 // If you do not agree to this license, do not download, install, 00007 // copy or use the software. 00008 // 00009 // 00010 // License Agreement 00011 // For Open Source Computer Vision Library 00012 // 00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 00015 // Third party copyrights are property of their respective owners. 00016 // 00017 // Redistribution and use in source and binary forms, with or without modification, 00018 // are permitted provided that the following conditions are met: 00019 // 00020 // * Redistribution's of source code must retain the above copyright notice, 00021 // this list of conditions and the following disclaimer. 00022 // 00023 // * Redistribution's in binary form must reproduce the above copyright notice, 00024 // this list of conditions and the following disclaimer in the documentation 00025 // and/or other materials provided with the distribution. 00026 // 00027 // * The name of the copyright holders may not be used to endorse or promote products 00028 // derived from this software without specific prior written permission. 00029 // 00030 // This software is provided by the copyright holders and contributors "as is" and 00031 // any express or implied warranties, including, but not limited to, the implied 00032 // warranties of merchantability and fitness for a particular purpose are disclaimed. 00033 // In no event shall the Intel Corporation or contributors be liable for any direct, 00034 // indirect, incidental, special, exemplary, or consequential damages 00035 // (including, but not limited to, procurement of substitute goods or services; 00036 // loss of use, data, or profits; or business interruption) however caused 00037 // and on any theory of liability, whether in contract, strict liability, 00038 // or tort (including negligence or otherwise) arising in any way out of 00039 // the use of this software, even if advised of the possibility of such damage. 00040 // 00041 //M*/ 00042 00043 #ifndef __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 00044 #define __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 00045 00046 #include "saturate_cast.hpp " 00047 #include "vec_traits.hpp " 00048 #include "vec_math.hpp " 00049 00050 /** @file 00051 * @deprecated Use @ref cudev instead. 00052 */ 00053 00054 //! @cond IGNORED 00055 00056 namespace cv { namespace cuda { namespace device 00057 { 00058 ////////////////////////////////////////////////////////////// 00059 // BrdConstant 00060 00061 template <typename D> struct BrdRowConstant 00062 { 00063 typedef D result_type; 00064 00065 explicit __host__ __device__ __forceinline__ BrdRowConstant(int width_, const D& val_ = VecTraits<D>::all(0)) : width(width_), val(val_) {} 00066 00067 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 00068 { 00069 return x >= 0 ? saturate_cast<D>(data[x]) : val; 00070 } 00071 00072 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 00073 { 00074 return x < width ? saturate_cast<D>(data[x]) : val; 00075 } 00076 00077 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 00078 { 00079 return (x >= 0 && x < width) ? saturate_cast<D>(data[x]) : val; 00080 } 00081 00082 int width; 00083 D val; 00084 }; 00085 00086 template <typename D> struct BrdColConstant 00087 { 00088 typedef D result_type; 00089 00090 explicit __host__ __device__ __forceinline__ BrdColConstant(int height_, const D& val_ = VecTraits<D>::all(0)) : height(height_), val(val_) {} 00091 00092 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 00093 { 00094 return y >= 0 ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 00095 } 00096 00097 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 00098 { 00099 return y < height ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 00100 } 00101 00102 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 00103 { 00104 return (y >= 0 && y < height) ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 00105 } 00106 00107 int height; 00108 D val; 00109 }; 00110 00111 template <typename D> struct BrdConstant 00112 { 00113 typedef D result_type; 00114 00115 __host__ __device__ __forceinline__ BrdConstant(int height_, int width_, const D& val_ = VecTraits<D>::all(0)) : height(height_), width(width_), val(val_) 00116 { 00117 } 00118 00119 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 00120 { 00121 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(((const T*)((const uchar*)data + y * step))[x]) : val; 00122 } 00123 00124 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 00125 { 00126 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val; 00127 } 00128 00129 int height; 00130 int width; 00131 D val; 00132 }; 00133 00134 ////////////////////////////////////////////////////////////// 00135 // BrdReplicate 00136 00137 template <typename D> struct BrdRowReplicate 00138 { 00139 typedef D result_type; 00140 00141 explicit __host__ __device__ __forceinline__ BrdRowReplicate(int width) : last_col(width - 1) {} 00142 template <typename U> __host__ __device__ __forceinline__ BrdRowReplicate(int width, U) : last_col(width - 1) {} 00143 00144 __device__ __forceinline__ int idx_col_low(int x) const 00145 { 00146 return ::max(x, 0); 00147 } 00148 00149 __device__ __forceinline__ int idx_col_high(int x) const 00150 { 00151 return ::min(x, last_col); 00152 } 00153 00154 __device__ __forceinline__ int idx_col(int x) const 00155 { 00156 return idx_col_low(idx_col_high(x)); 00157 } 00158 00159 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 00160 { 00161 return saturate_cast<D>(data[idx_col_low(x)]); 00162 } 00163 00164 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 00165 { 00166 return saturate_cast<D>(data[idx_col_high(x)]); 00167 } 00168 00169 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 00170 { 00171 return saturate_cast<D>(data[idx_col(x)]); 00172 } 00173 00174 int last_col; 00175 }; 00176 00177 template <typename D> struct BrdColReplicate 00178 { 00179 typedef D result_type; 00180 00181 explicit __host__ __device__ __forceinline__ BrdColReplicate(int height) : last_row(height - 1) {} 00182 template <typename U> __host__ __device__ __forceinline__ BrdColReplicate(int height, U) : last_row(height - 1) {} 00183 00184 __device__ __forceinline__ int idx_row_low(int y) const 00185 { 00186 return ::max(y, 0); 00187 } 00188 00189 __device__ __forceinline__ int idx_row_high(int y) const 00190 { 00191 return ::min(y, last_row); 00192 } 00193 00194 __device__ __forceinline__ int idx_row(int y) const 00195 { 00196 return idx_row_low(idx_row_high(y)); 00197 } 00198 00199 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 00200 { 00201 return saturate_cast<D>(*(const T*)((const char*)data + idx_row_low(y) * step)); 00202 } 00203 00204 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 00205 { 00206 return saturate_cast<D>(*(const T*)((const char*)data + idx_row_high(y) * step)); 00207 } 00208 00209 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 00210 { 00211 return saturate_cast<D>(*(const T*)((const char*)data + idx_row(y) * step)); 00212 } 00213 00214 int last_row; 00215 }; 00216 00217 template <typename D> struct BrdReplicate 00218 { 00219 typedef D result_type; 00220 00221 __host__ __device__ __forceinline__ BrdReplicate(int height, int width) : last_row(height - 1), last_col(width - 1) {} 00222 template <typename U> __host__ __device__ __forceinline__ BrdReplicate(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 00223 00224 __device__ __forceinline__ int idx_row_low(int y) const 00225 { 00226 return ::max(y, 0); 00227 } 00228 00229 __device__ __forceinline__ int idx_row_high(int y) const 00230 { 00231 return ::min(y, last_row); 00232 } 00233 00234 __device__ __forceinline__ int idx_row(int y) const 00235 { 00236 return idx_row_low(idx_row_high(y)); 00237 } 00238 00239 __device__ __forceinline__ int idx_col_low(int x) const 00240 { 00241 return ::max(x, 0); 00242 } 00243 00244 __device__ __forceinline__ int idx_col_high(int x) const 00245 { 00246 return ::min(x, last_col); 00247 } 00248 00249 __device__ __forceinline__ int idx_col(int x) const 00250 { 00251 return idx_col_low(idx_col_high(x)); 00252 } 00253 00254 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 00255 { 00256 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 00257 } 00258 00259 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 00260 { 00261 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 00262 } 00263 00264 int last_row; 00265 int last_col; 00266 }; 00267 00268 ////////////////////////////////////////////////////////////// 00269 // BrdReflect101 00270 00271 template <typename D> struct BrdRowReflect101 00272 { 00273 typedef D result_type; 00274 00275 explicit __host__ __device__ __forceinline__ BrdRowReflect101(int width) : last_col(width - 1) {} 00276 template <typename U> __host__ __device__ __forceinline__ BrdRowReflect101(int width, U) : last_col(width - 1) {} 00277 00278 __device__ __forceinline__ int idx_col_low(int x) const 00279 { 00280 return ::abs(x) % (last_col + 1); 00281 } 00282 00283 __device__ __forceinline__ int idx_col_high(int x) const 00284 { 00285 return ::abs(last_col - ::abs(last_col - x)) % (last_col + 1); 00286 } 00287 00288 __device__ __forceinline__ int idx_col(int x) const 00289 { 00290 return idx_col_low(idx_col_high(x)); 00291 } 00292 00293 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 00294 { 00295 return saturate_cast<D>(data[idx_col_low(x)]); 00296 } 00297 00298 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 00299 { 00300 return saturate_cast<D>(data[idx_col_high(x)]); 00301 } 00302 00303 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 00304 { 00305 return saturate_cast<D>(data[idx_col(x)]); 00306 } 00307 00308 int last_col; 00309 }; 00310 00311 template <typename D> struct BrdColReflect101 00312 { 00313 typedef D result_type; 00314 00315 explicit __host__ __device__ __forceinline__ BrdColReflect101(int height) : last_row(height - 1) {} 00316 template <typename U> __host__ __device__ __forceinline__ BrdColReflect101(int height, U) : last_row(height - 1) {} 00317 00318 __device__ __forceinline__ int idx_row_low(int y) const 00319 { 00320 return ::abs(y) % (last_row + 1); 00321 } 00322 00323 __device__ __forceinline__ int idx_row_high(int y) const 00324 { 00325 return ::abs(last_row - ::abs(last_row - y)) % (last_row + 1); 00326 } 00327 00328 __device__ __forceinline__ int idx_row(int y) const 00329 { 00330 return idx_row_low(idx_row_high(y)); 00331 } 00332 00333 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 00334 { 00335 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 00336 } 00337 00338 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 00339 { 00340 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 00341 } 00342 00343 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 00344 { 00345 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 00346 } 00347 00348 int last_row; 00349 }; 00350 00351 template <typename D> struct BrdReflect101 00352 { 00353 typedef D result_type; 00354 00355 __host__ __device__ __forceinline__ BrdReflect101(int height, int width) : last_row(height - 1), last_col(width - 1) {} 00356 template <typename U> __host__ __device__ __forceinline__ BrdReflect101(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 00357 00358 __device__ __forceinline__ int idx_row_low(int y) const 00359 { 00360 return ::abs(y) % (last_row + 1); 00361 } 00362 00363 __device__ __forceinline__ int idx_row_high(int y) const 00364 { 00365 return ::abs(last_row - ::abs(last_row - y)) % (last_row + 1); 00366 } 00367 00368 __device__ __forceinline__ int idx_row(int y) const 00369 { 00370 return idx_row_low(idx_row_high(y)); 00371 } 00372 00373 __device__ __forceinline__ int idx_col_low(int x) const 00374 { 00375 return ::abs(x) % (last_col + 1); 00376 } 00377 00378 __device__ __forceinline__ int idx_col_high(int x) const 00379 { 00380 return ::abs(last_col - ::abs(last_col - x)) % (last_col + 1); 00381 } 00382 00383 __device__ __forceinline__ int idx_col(int x) const 00384 { 00385 return idx_col_low(idx_col_high(x)); 00386 } 00387 00388 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 00389 { 00390 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 00391 } 00392 00393 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 00394 { 00395 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 00396 } 00397 00398 int last_row; 00399 int last_col; 00400 }; 00401 00402 ////////////////////////////////////////////////////////////// 00403 // BrdReflect 00404 00405 template <typename D> struct BrdRowReflect 00406 { 00407 typedef D result_type; 00408 00409 explicit __host__ __device__ __forceinline__ BrdRowReflect(int width) : last_col(width - 1) {} 00410 template <typename U> __host__ __device__ __forceinline__ BrdRowReflect(int width, U) : last_col(width - 1) {} 00411 00412 __device__ __forceinline__ int idx_col_low(int x) const 00413 { 00414 return (::abs(x) - (x < 0)) % (last_col + 1); 00415 } 00416 00417 __device__ __forceinline__ int idx_col_high(int x) const 00418 { 00419 return ::abs(last_col - ::abs(last_col - x) + (x > last_col)) % (last_col + 1); 00420 } 00421 00422 __device__ __forceinline__ int idx_col(int x) const 00423 { 00424 return idx_col_high(::abs(x) - (x < 0)); 00425 } 00426 00427 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 00428 { 00429 return saturate_cast<D>(data[idx_col_low(x)]); 00430 } 00431 00432 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 00433 { 00434 return saturate_cast<D>(data[idx_col_high(x)]); 00435 } 00436 00437 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 00438 { 00439 return saturate_cast<D>(data[idx_col(x)]); 00440 } 00441 00442 int last_col; 00443 }; 00444 00445 template <typename D> struct BrdColReflect 00446 { 00447 typedef D result_type; 00448 00449 explicit __host__ __device__ __forceinline__ BrdColReflect(int height) : last_row(height - 1) {} 00450 template <typename U> __host__ __device__ __forceinline__ BrdColReflect(int height, U) : last_row(height - 1) {} 00451 00452 __device__ __forceinline__ int idx_row_low(int y) const 00453 { 00454 return (::abs(y) - (y < 0)) % (last_row + 1); 00455 } 00456 00457 __device__ __forceinline__ int idx_row_high(int y) const 00458 { 00459 return ::abs(last_row - ::abs(last_row - y) + (y > last_row)) % (last_row + 1); 00460 } 00461 00462 __device__ __forceinline__ int idx_row(int y) const 00463 { 00464 return idx_row_high(::abs(y) - (y < 0)); 00465 } 00466 00467 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 00468 { 00469 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 00470 } 00471 00472 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 00473 { 00474 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 00475 } 00476 00477 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 00478 { 00479 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 00480 } 00481 00482 int last_row; 00483 }; 00484 00485 template <typename D> struct BrdReflect 00486 { 00487 typedef D result_type; 00488 00489 __host__ __device__ __forceinline__ BrdReflect(int height, int width) : last_row(height - 1), last_col(width - 1) {} 00490 template <typename U> __host__ __device__ __forceinline__ BrdReflect(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 00491 00492 __device__ __forceinline__ int idx_row_low(int y) const 00493 { 00494 return (::abs(y) - (y < 0)) % (last_row + 1); 00495 } 00496 00497 __device__ __forceinline__ int idx_row_high(int y) const 00498 { 00499 return /*::abs*/(last_row - ::abs(last_row - y) + (y > last_row)) /*% (last_row + 1)*/; 00500 } 00501 00502 __device__ __forceinline__ int idx_row(int y) const 00503 { 00504 return idx_row_low(idx_row_high(y)); 00505 } 00506 00507 __device__ __forceinline__ int idx_col_low(int x) const 00508 { 00509 return (::abs(x) - (x < 0)) % (last_col + 1); 00510 } 00511 00512 __device__ __forceinline__ int idx_col_high(int x) const 00513 { 00514 return (last_col - ::abs(last_col - x) + (x > last_col)); 00515 } 00516 00517 __device__ __forceinline__ int idx_col(int x) const 00518 { 00519 return idx_col_low(idx_col_high(x)); 00520 } 00521 00522 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 00523 { 00524 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 00525 } 00526 00527 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 00528 { 00529 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 00530 } 00531 00532 int last_row; 00533 int last_col; 00534 }; 00535 00536 ////////////////////////////////////////////////////////////// 00537 // BrdWrap 00538 00539 template <typename D> struct BrdRowWrap 00540 { 00541 typedef D result_type; 00542 00543 explicit __host__ __device__ __forceinline__ BrdRowWrap(int width_) : width(width_) {} 00544 template <typename U> __host__ __device__ __forceinline__ BrdRowWrap(int width_, U) : width(width_) {} 00545 00546 __device__ __forceinline__ int idx_col_low(int x) const 00547 { 00548 return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width); 00549 } 00550 00551 __device__ __forceinline__ int idx_col_high(int x) const 00552 { 00553 return (x < width) * x + (x >= width) * (x % width); 00554 } 00555 00556 __device__ __forceinline__ int idx_col(int x) const 00557 { 00558 return idx_col_high(idx_col_low(x)); 00559 } 00560 00561 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 00562 { 00563 return saturate_cast<D>(data[idx_col_low(x)]); 00564 } 00565 00566 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 00567 { 00568 return saturate_cast<D>(data[idx_col_high(x)]); 00569 } 00570 00571 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 00572 { 00573 return saturate_cast<D>(data[idx_col(x)]); 00574 } 00575 00576 int width; 00577 }; 00578 00579 template <typename D> struct BrdColWrap 00580 { 00581 typedef D result_type; 00582 00583 explicit __host__ __device__ __forceinline__ BrdColWrap(int height_) : height(height_) {} 00584 template <typename U> __host__ __device__ __forceinline__ BrdColWrap(int height_, U) : height(height_) {} 00585 00586 __device__ __forceinline__ int idx_row_low(int y) const 00587 { 00588 return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height); 00589 } 00590 00591 __device__ __forceinline__ int idx_row_high(int y) const 00592 { 00593 return (y < height) * y + (y >= height) * (y % height); 00594 } 00595 00596 __device__ __forceinline__ int idx_row(int y) const 00597 { 00598 return idx_row_high(idx_row_low(y)); 00599 } 00600 00601 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 00602 { 00603 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 00604 } 00605 00606 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 00607 { 00608 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 00609 } 00610 00611 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 00612 { 00613 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 00614 } 00615 00616 int height; 00617 }; 00618 00619 template <typename D> struct BrdWrap 00620 { 00621 typedef D result_type; 00622 00623 __host__ __device__ __forceinline__ BrdWrap(int height_, int width_) : 00624 height(height_), width(width_) 00625 { 00626 } 00627 template <typename U> 00628 __host__ __device__ __forceinline__ BrdWrap(int height_, int width_, U) : 00629 height(height_), width(width_) 00630 { 00631 } 00632 00633 __device__ __forceinline__ int idx_row_low(int y) const 00634 { 00635 return (y >= 0) ? y : (y - ((y - height + 1) / height) * height); 00636 } 00637 00638 __device__ __forceinline__ int idx_row_high(int y) const 00639 { 00640 return (y < height) ? y : (y % height); 00641 } 00642 00643 __device__ __forceinline__ int idx_row(int y) const 00644 { 00645 return idx_row_high(idx_row_low(y)); 00646 } 00647 00648 __device__ __forceinline__ int idx_col_low(int x) const 00649 { 00650 return (x >= 0) ? x : (x - ((x - width + 1) / width) * width); 00651 } 00652 00653 __device__ __forceinline__ int idx_col_high(int x) const 00654 { 00655 return (x < width) ? x : (x % width); 00656 } 00657 00658 __device__ __forceinline__ int idx_col(int x) const 00659 { 00660 return idx_col_high(idx_col_low(x)); 00661 } 00662 00663 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 00664 { 00665 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 00666 } 00667 00668 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 00669 { 00670 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 00671 } 00672 00673 int height; 00674 int width; 00675 }; 00676 00677 ////////////////////////////////////////////////////////////// 00678 // BorderReader 00679 00680 template <typename Ptr2D, typename B> struct BorderReader 00681 { 00682 typedef typename B::result_type elem_type; 00683 typedef typename Ptr2D::index_type index_type; 00684 00685 __host__ __device__ __forceinline__ BorderReader(const Ptr2D& ptr_, const B& b_) : ptr(ptr_), b(b_) {} 00686 00687 __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const 00688 { 00689 return b.at(y, x, ptr); 00690 } 00691 00692 Ptr2D ptr; 00693 B b; 00694 }; 00695 00696 // under win32 there is some bug with templated types that passed as kernel parameters 00697 // with this specialization all works fine 00698 template <typename Ptr2D, typename D> struct BorderReader< Ptr2D, BrdConstant<D> > 00699 { 00700 typedef typename BrdConstant<D>::result_type elem_type; 00701 typedef typename Ptr2D::index_type index_type; 00702 00703 __host__ __device__ __forceinline__ BorderReader(const Ptr2D& src_, const BrdConstant<D>& b) : 00704 src(src_), height(b.height), width(b.width), val(b.val) 00705 { 00706 } 00707 00708 __device__ __forceinline__ D operator ()(index_type y, index_type x) const 00709 { 00710 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val; 00711 } 00712 00713 Ptr2D src; 00714 int height; 00715 int width; 00716 D val; 00717 }; 00718 }}} // namespace cv { namespace cuda { namespace cudev 00719 00720 //! @endcond 00721 00722 #endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 00723
Generated on Tue Jul 12 2022 14:46:03 by
1.7.2