1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 5 // By downloading, copying, installing or using the software you agree to this license. 6 // If you do not agree to this license, do not download, install, 7 // copy or use the software. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15 // Third party copyrights are property of their respective owners. 16 // 17 // Redistribution and use in source and binary forms, with or without modification, 18 // are permitted provided that the following conditions are met: 19 // 20 // * Redistribution's of source code must retain the above copyright notice, 21 // this list of conditions and the following disclaimer. 22 // 23 // * Redistribution's in binary form must reproduce the above copyright notice, 24 // this list of conditions and the following disclaimer in the documentation 25 // and/or other materials provided with the distribution. 26 // 27 // * The name of the copyright holders may not be used to endorse or promote products 28 // derived from this software without specific prior written permission. 29 // 30 // This software is provided by the copyright holders and contributors "as is" and 31 // any express or implied warranties, including, but not limited to, the implied 32 // warranties of merchantability and fitness for a particular purpose are disclaimed. 33 // In no event shall the Intel Corporation or contributors be liable for any direct, 34 // indirect, incidental, special, exemplary, or consequential damages 35 // (including, but not limited to, procurement of substitute goods or services; 36 // loss of use, data, or profits; or business interruption) however caused 37 // and on any theory of liability, whether in contract, strict liability, 38 // or tort (including negligence or otherwise) arising in any way out of 39 // the use of this software, even if advised of the possibility of such damage. 40 // 41 //M*/ 42 43 #ifndef __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 44 #define __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 45 46 #include "saturate_cast.hpp" 47 #include "vec_traits.hpp" 48 #include "vec_math.hpp" 49 50 /** @file 51 * @deprecated Use @ref cudev instead. 52 */ 53 54 //! @cond IGNORED 55 56 namespace cv { namespace cuda { namespace device 57 { 58 ////////////////////////////////////////////////////////////// 59 // BrdConstant 60 61 template <typename D> struct BrdRowConstant 62 { 63 typedef D result_type; 64 65 explicit __host__ __device__ __forceinline__ BrdRowConstant(int width_, const D& val_ = VecTraits<D>::all(0)) : width(width_), val(val_) {} 66 67 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 68 { 69 return x >= 0 ? saturate_cast<D>(data[x]) : val; 70 } 71 72 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 73 { 74 return x < width ? saturate_cast<D>(data[x]) : val; 75 } 76 77 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 78 { 79 return (x >= 0 && x < width) ? saturate_cast<D>(data[x]) : val; 80 } 81 82 int width; 83 D val; 84 }; 85 86 template <typename D> struct BrdColConstant 87 { 88 typedef D result_type; 89 90 explicit __host__ __device__ __forceinline__ BrdColConstant(int height_, const D& val_ = VecTraits<D>::all(0)) : height(height_), val(val_) {} 91 92 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 93 { 94 return y >= 0 ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 95 } 96 97 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 98 { 99 return y < height ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 100 } 101 102 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 103 { 104 return (y >= 0 && y < height) ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val; 105 } 106 107 int height; 108 D val; 109 }; 110 111 template <typename D> struct BrdConstant 112 { 113 typedef D result_type; 114 115 __host__ __device__ __forceinline__ BrdConstant(int height_, int width_, const D& val_ = VecTraits<D>::all(0)) : height(height_), width(width_), val(val_) 116 { 117 } 118 119 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 120 { 121 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(((const T*)((const uchar*)data + y * step))[x]) : val; 122 } 123 124 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 125 { 126 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val; 127 } 128 129 int height; 130 int width; 131 D val; 132 }; 133 134 ////////////////////////////////////////////////////////////// 135 // BrdReplicate 136 137 template <typename D> struct BrdRowReplicate 138 { 139 typedef D result_type; 140 141 explicit __host__ __device__ __forceinline__ BrdRowReplicate(int width) : last_col(width - 1) {} 142 template <typename U> __host__ __device__ __forceinline__ BrdRowReplicate(int width, U) : last_col(width - 1) {} 143 144 __device__ __forceinline__ int idx_col_low(int x) const 145 { 146 return ::max(x, 0); 147 } 148 149 __device__ __forceinline__ int idx_col_high(int x) const 150 { 151 return ::min(x, last_col); 152 } 153 154 __device__ __forceinline__ int idx_col(int x) const 155 { 156 return idx_col_low(idx_col_high(x)); 157 } 158 159 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 160 { 161 return saturate_cast<D>(data[idx_col_low(x)]); 162 } 163 164 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 165 { 166 return saturate_cast<D>(data[idx_col_high(x)]); 167 } 168 169 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 170 { 171 return saturate_cast<D>(data[idx_col(x)]); 172 } 173 174 int last_col; 175 }; 176 177 template <typename D> struct BrdColReplicate 178 { 179 typedef D result_type; 180 181 explicit __host__ __device__ __forceinline__ BrdColReplicate(int height) : last_row(height - 1) {} 182 template <typename U> __host__ __device__ __forceinline__ BrdColReplicate(int height, U) : last_row(height - 1) {} 183 184 __device__ __forceinline__ int idx_row_low(int y) const 185 { 186 return ::max(y, 0); 187 } 188 189 __device__ __forceinline__ int idx_row_high(int y) const 190 { 191 return ::min(y, last_row); 192 } 193 194 __device__ __forceinline__ int idx_row(int y) const 195 { 196 return idx_row_low(idx_row_high(y)); 197 } 198 199 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 200 { 201 return saturate_cast<D>(*(const T*)((const char*)data + idx_row_low(y) * step)); 202 } 203 204 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 205 { 206 return saturate_cast<D>(*(const T*)((const char*)data + idx_row_high(y) * step)); 207 } 208 209 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 210 { 211 return saturate_cast<D>(*(const T*)((const char*)data + idx_row(y) * step)); 212 } 213 214 int last_row; 215 }; 216 217 template <typename D> struct BrdReplicate 218 { 219 typedef D result_type; 220 221 __host__ __device__ __forceinline__ BrdReplicate(int height, int width) : last_row(height - 1), last_col(width - 1) {} 222 template <typename U> __host__ __device__ __forceinline__ BrdReplicate(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 223 224 __device__ __forceinline__ int idx_row_low(int y) const 225 { 226 return ::max(y, 0); 227 } 228 229 __device__ __forceinline__ int idx_row_high(int y) const 230 { 231 return ::min(y, last_row); 232 } 233 234 __device__ __forceinline__ int idx_row(int y) const 235 { 236 return idx_row_low(idx_row_high(y)); 237 } 238 239 __device__ __forceinline__ int idx_col_low(int x) const 240 { 241 return ::max(x, 0); 242 } 243 244 __device__ __forceinline__ int idx_col_high(int x) const 245 { 246 return ::min(x, last_col); 247 } 248 249 __device__ __forceinline__ int idx_col(int x) const 250 { 251 return idx_col_low(idx_col_high(x)); 252 } 253 254 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 255 { 256 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 257 } 258 259 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 260 { 261 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 262 } 263 264 int last_row; 265 int last_col; 266 }; 267 268 ////////////////////////////////////////////////////////////// 269 // BrdReflect101 270 271 template <typename D> struct BrdRowReflect101 272 { 273 typedef D result_type; 274 275 explicit __host__ __device__ __forceinline__ BrdRowReflect101(int width) : last_col(width - 1) {} 276 template <typename U> __host__ __device__ __forceinline__ BrdRowReflect101(int width, U) : last_col(width - 1) {} 277 278 __device__ __forceinline__ int idx_col_low(int x) const 279 { 280 return ::abs(x) % (last_col + 1); 281 } 282 283 __device__ __forceinline__ int idx_col_high(int x) const 284 { 285 return ::abs(last_col - ::abs(last_col - x)) % (last_col + 1); 286 } 287 288 __device__ __forceinline__ int idx_col(int x) const 289 { 290 return idx_col_low(idx_col_high(x)); 291 } 292 293 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 294 { 295 return saturate_cast<D>(data[idx_col_low(x)]); 296 } 297 298 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 299 { 300 return saturate_cast<D>(data[idx_col_high(x)]); 301 } 302 303 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 304 { 305 return saturate_cast<D>(data[idx_col(x)]); 306 } 307 308 int last_col; 309 }; 310 311 template <typename D> struct BrdColReflect101 312 { 313 typedef D result_type; 314 315 explicit __host__ __device__ __forceinline__ BrdColReflect101(int height) : last_row(height - 1) {} 316 template <typename U> __host__ __device__ __forceinline__ BrdColReflect101(int height, U) : last_row(height - 1) {} 317 318 __device__ __forceinline__ int idx_row_low(int y) const 319 { 320 return ::abs(y) % (last_row + 1); 321 } 322 323 __device__ __forceinline__ int idx_row_high(int y) const 324 { 325 return ::abs(last_row - ::abs(last_row - y)) % (last_row + 1); 326 } 327 328 __device__ __forceinline__ int idx_row(int y) const 329 { 330 return idx_row_low(idx_row_high(y)); 331 } 332 333 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 334 { 335 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 336 } 337 338 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 339 { 340 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 341 } 342 343 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 344 { 345 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 346 } 347 348 int last_row; 349 }; 350 351 template <typename D> struct BrdReflect101 352 { 353 typedef D result_type; 354 355 __host__ __device__ __forceinline__ BrdReflect101(int height, int width) : last_row(height - 1), last_col(width - 1) {} 356 template <typename U> __host__ __device__ __forceinline__ BrdReflect101(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 357 358 __device__ __forceinline__ int idx_row_low(int y) const 359 { 360 return ::abs(y) % (last_row + 1); 361 } 362 363 __device__ __forceinline__ int idx_row_high(int y) const 364 { 365 return ::abs(last_row - ::abs(last_row - y)) % (last_row + 1); 366 } 367 368 __device__ __forceinline__ int idx_row(int y) const 369 { 370 return idx_row_low(idx_row_high(y)); 371 } 372 373 __device__ __forceinline__ int idx_col_low(int x) const 374 { 375 return ::abs(x) % (last_col + 1); 376 } 377 378 __device__ __forceinline__ int idx_col_high(int x) const 379 { 380 return ::abs(last_col - ::abs(last_col - x)) % (last_col + 1); 381 } 382 383 __device__ __forceinline__ int idx_col(int x) const 384 { 385 return idx_col_low(idx_col_high(x)); 386 } 387 388 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 389 { 390 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 391 } 392 393 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 394 { 395 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 396 } 397 398 int last_row; 399 int last_col; 400 }; 401 402 ////////////////////////////////////////////////////////////// 403 // BrdReflect 404 405 template <typename D> struct BrdRowReflect 406 { 407 typedef D result_type; 408 409 explicit __host__ __device__ __forceinline__ BrdRowReflect(int width) : last_col(width - 1) {} 410 template <typename U> __host__ __device__ __forceinline__ BrdRowReflect(int width, U) : last_col(width - 1) {} 411 412 __device__ __forceinline__ int idx_col_low(int x) const 413 { 414 return (::abs(x) - (x < 0)) % (last_col + 1); 415 } 416 417 __device__ __forceinline__ int idx_col_high(int x) const 418 { 419 return ::abs(last_col - ::abs(last_col - x) + (x > last_col)) % (last_col + 1); 420 } 421 422 __device__ __forceinline__ int idx_col(int x) const 423 { 424 return idx_col_high(::abs(x) - (x < 0)); 425 } 426 427 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 428 { 429 return saturate_cast<D>(data[idx_col_low(x)]); 430 } 431 432 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 433 { 434 return saturate_cast<D>(data[idx_col_high(x)]); 435 } 436 437 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 438 { 439 return saturate_cast<D>(data[idx_col(x)]); 440 } 441 442 int last_col; 443 }; 444 445 template <typename D> struct BrdColReflect 446 { 447 typedef D result_type; 448 449 explicit __host__ __device__ __forceinline__ BrdColReflect(int height) : last_row(height - 1) {} 450 template <typename U> __host__ __device__ __forceinline__ BrdColReflect(int height, U) : last_row(height - 1) {} 451 452 __device__ __forceinline__ int idx_row_low(int y) const 453 { 454 return (::abs(y) - (y < 0)) % (last_row + 1); 455 } 456 457 __device__ __forceinline__ int idx_row_high(int y) const 458 { 459 return ::abs(last_row - ::abs(last_row - y) + (y > last_row)) % (last_row + 1); 460 } 461 462 __device__ __forceinline__ int idx_row(int y) const 463 { 464 return idx_row_high(::abs(y) - (y < 0)); 465 } 466 467 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 468 { 469 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 470 } 471 472 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 473 { 474 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 475 } 476 477 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 478 { 479 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 480 } 481 482 int last_row; 483 }; 484 485 template <typename D> struct BrdReflect 486 { 487 typedef D result_type; 488 489 __host__ __device__ __forceinline__ BrdReflect(int height, int width) : last_row(height - 1), last_col(width - 1) {} 490 template <typename U> __host__ __device__ __forceinline__ BrdReflect(int height, int width, U) : last_row(height - 1), last_col(width - 1) {} 491 492 __device__ __forceinline__ int idx_row_low(int y) const 493 { 494 return (::abs(y) - (y < 0)) % (last_row + 1); 495 } 496 497 __device__ __forceinline__ int idx_row_high(int y) const 498 { 499 return /*::abs*/(last_row - ::abs(last_row - y) + (y > last_row)) /*% (last_row + 1)*/; 500 } 501 502 __device__ __forceinline__ int idx_row(int y) const 503 { 504 return idx_row_low(idx_row_high(y)); 505 } 506 507 __device__ __forceinline__ int idx_col_low(int x) const 508 { 509 return (::abs(x) - (x < 0)) % (last_col + 1); 510 } 511 512 __device__ __forceinline__ int idx_col_high(int x) const 513 { 514 return (last_col - ::abs(last_col - x) + (x > last_col)); 515 } 516 517 __device__ __forceinline__ int idx_col(int x) const 518 { 519 return idx_col_low(idx_col_high(x)); 520 } 521 522 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 523 { 524 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 525 } 526 527 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 528 { 529 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 530 } 531 532 int last_row; 533 int last_col; 534 }; 535 536 ////////////////////////////////////////////////////////////// 537 // BrdWrap 538 539 template <typename D> struct BrdRowWrap 540 { 541 typedef D result_type; 542 543 explicit __host__ __device__ __forceinline__ BrdRowWrap(int width_) : width(width_) {} 544 template <typename U> __host__ __device__ __forceinline__ BrdRowWrap(int width_, U) : width(width_) {} 545 546 __device__ __forceinline__ int idx_col_low(int x) const 547 { 548 return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width); 549 } 550 551 __device__ __forceinline__ int idx_col_high(int x) const 552 { 553 return (x < width) * x + (x >= width) * (x % width); 554 } 555 556 __device__ __forceinline__ int idx_col(int x) const 557 { 558 return idx_col_high(idx_col_low(x)); 559 } 560 561 template <typename T> __device__ __forceinline__ D at_low(int x, const T* data) const 562 { 563 return saturate_cast<D>(data[idx_col_low(x)]); 564 } 565 566 template <typename T> __device__ __forceinline__ D at_high(int x, const T* data) const 567 { 568 return saturate_cast<D>(data[idx_col_high(x)]); 569 } 570 571 template <typename T> __device__ __forceinline__ D at(int x, const T* data) const 572 { 573 return saturate_cast<D>(data[idx_col(x)]); 574 } 575 576 int width; 577 }; 578 579 template <typename D> struct BrdColWrap 580 { 581 typedef D result_type; 582 583 explicit __host__ __device__ __forceinline__ BrdColWrap(int height_) : height(height_) {} 584 template <typename U> __host__ __device__ __forceinline__ BrdColWrap(int height_, U) : height(height_) {} 585 586 __device__ __forceinline__ int idx_row_low(int y) const 587 { 588 return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height); 589 } 590 591 __device__ __forceinline__ int idx_row_high(int y) const 592 { 593 return (y < height) * y + (y >= height) * (y % height); 594 } 595 596 __device__ __forceinline__ int idx_row(int y) const 597 { 598 return idx_row_high(idx_row_low(y)); 599 } 600 601 template <typename T> __device__ __forceinline__ D at_low(int y, const T* data, size_t step) const 602 { 603 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_low(y) * step)); 604 } 605 606 template <typename T> __device__ __forceinline__ D at_high(int y, const T* data, size_t step) const 607 { 608 return saturate_cast<D>(*(const D*)((const char*)data + idx_row_high(y) * step)); 609 } 610 611 template <typename T> __device__ __forceinline__ D at(int y, const T* data, size_t step) const 612 { 613 return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step)); 614 } 615 616 int height; 617 }; 618 619 template <typename D> struct BrdWrap 620 { 621 typedef D result_type; 622 623 __host__ __device__ __forceinline__ BrdWrap(int height_, int width_) : 624 height(height_), width(width_) 625 { 626 } 627 template <typename U> 628 __host__ __device__ __forceinline__ BrdWrap(int height_, int width_, U) : 629 height(height_), width(width_) 630 { 631 } 632 633 __device__ __forceinline__ int idx_row_low(int y) const 634 { 635 return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height); 636 } 637 638 __device__ __forceinline__ int idx_row_high(int y) const 639 { 640 return (y < height) * y + (y >= height) * (y % height); 641 } 642 643 __device__ __forceinline__ int idx_row(int y) const 644 { 645 return idx_row_high(idx_row_low(y)); 646 } 647 648 __device__ __forceinline__ int idx_col_low(int x) const 649 { 650 return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width); 651 } 652 653 __device__ __forceinline__ int idx_col_high(int x) const 654 { 655 return (x < width) * x + (x >= width) * (x % width); 656 } 657 658 __device__ __forceinline__ int idx_col(int x) const 659 { 660 return idx_col_high(idx_col_low(x)); 661 } 662 663 template <typename T> __device__ __forceinline__ D at(int y, int x, const T* data, size_t step) const 664 { 665 return saturate_cast<D>(((const T*)((const char*)data + idx_row(y) * step))[idx_col(x)]); 666 } 667 668 template <typename Ptr2D> __device__ __forceinline__ D at(typename Ptr2D::index_type y, typename Ptr2D::index_type x, const Ptr2D& src) const 669 { 670 return saturate_cast<D>(src(idx_row(y), idx_col(x))); 671 } 672 673 int height; 674 int width; 675 }; 676 677 ////////////////////////////////////////////////////////////// 678 // BorderReader 679 680 template <typename Ptr2D, typename B> struct BorderReader 681 { 682 typedef typename B::result_type elem_type; 683 typedef typename Ptr2D::index_type index_type; 684 685 __host__ __device__ __forceinline__ BorderReader(const Ptr2D& ptr_, const B& b_) : ptr(ptr_), b(b_) {} 686 687 __device__ __forceinline__ elem_type operator ()(index_type y, index_type x) const 688 { 689 return b.at(y, x, ptr); 690 } 691 692 Ptr2D ptr; 693 B b; 694 }; 695 696 // under win32 there is some bug with templated types that passed as kernel parameters 697 // with this specialization all works fine 698 template <typename Ptr2D, typename D> struct BorderReader< Ptr2D, BrdConstant<D> > 699 { 700 typedef typename BrdConstant<D>::result_type elem_type; 701 typedef typename Ptr2D::index_type index_type; 702 703 __host__ __device__ __forceinline__ BorderReader(const Ptr2D& src_, const BrdConstant<D>& b) : 704 src(src_), height(b.height), width(b.width), val(b.val) 705 { 706 } 707 708 __device__ __forceinline__ D operator ()(index_type y, index_type x) const 709 { 710 return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val; 711 } 712 713 Ptr2D src; 714 int height; 715 int width; 716 D val; 717 }; 718 }}} // namespace cv { namespace cuda { namespace cudev 719 720 //! @endcond 721 722 #endif // __OPENCV_CUDA_BORDER_INTERPOLATE_HPP__ 723
