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 #include "precomp.hpp" 44 45 using namespace cv; 46 using namespace cv::cuda; 47 48 #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) 49 50 void cv::cuda::cvtColor(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); } 51 52 void cv::cuda::demosaicing(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); } 53 54 void cv::cuda::swapChannels(InputOutputArray, const int[], Stream&) { throw_no_cuda(); } 55 56 void cv::cuda::gammaCorrection(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); } 57 58 void cv::cuda::alphaComp(InputArray, InputArray, OutputArray, int, Stream&) { throw_no_cuda(); } 59 60 61 #else /* !defined (HAVE_CUDA) */ 62 63 #include "cvt_color_internal.h" 64 65 namespace cv { namespace cuda { 66 namespace device 67 { 68 template <int cn> 69 void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); 70 template <int cn> 71 void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); 72 73 template <int cn> 74 void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream); 75 } 76 }} 77 78 using namespace ::cv::cuda::device; 79 80 namespace 81 { 82 typedef void (*gpu_func_t)(const GpuMat& _src, GpuMat& _dst, Stream& stream); 83 84 void BGR_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream) 85 { 86 using namespace cv::cuda::device; 87 static const gpu_func_t funcs[] = {BGR_to_RGB_8u, 0, BGR_to_RGB_16u, 0, 0, BGR_to_RGB_32f}; 88 89 GpuMat src = _src.getGpuMat(); 90 91 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 92 CV_Assert( src.channels() == 3 ); 93 94 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3)); 95 GpuMat dst = _dst.getGpuMat(); 96 97 funcs[src.depth()](src, dst, stream); 98 } 99 100 void BGR_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream) 101 { 102 using namespace cv::cuda::device; 103 static const gpu_func_t funcs[] = {BGR_to_BGRA_8u, 0, BGR_to_BGRA_16u, 0, 0, BGR_to_BGRA_32f}; 104 105 GpuMat src = _src.getGpuMat(); 106 107 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 108 CV_Assert( src.channels() == 3 ); 109 110 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4)); 111 GpuMat dst = _dst.getGpuMat(); 112 113 funcs[src.depth()](src, dst, stream); 114 } 115 116 void BGR_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream) 117 { 118 using namespace cv::cuda::device; 119 static const gpu_func_t funcs[] = {BGR_to_RGBA_8u, 0, BGR_to_RGBA_16u, 0, 0, BGR_to_RGBA_32f}; 120 121 GpuMat src = _src.getGpuMat(); 122 123 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 124 CV_Assert( src.channels() == 3 ); 125 126 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4)); 127 GpuMat dst = _dst.getGpuMat(); 128 129 funcs[src.depth()](src, dst, stream); 130 } 131 132 void BGRA_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream) 133 { 134 using namespace cv::cuda::device; 135 static const gpu_func_t funcs[] = {BGRA_to_BGR_8u, 0, BGRA_to_BGR_16u, 0, 0, BGRA_to_BGR_32f}; 136 137 GpuMat src = _src.getGpuMat(); 138 139 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 140 CV_Assert( src.channels() == 4 ); 141 142 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3)); 143 GpuMat dst = _dst.getGpuMat(); 144 145 funcs[src.depth()](src, dst, stream); 146 } 147 148 void BGRA_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream) 149 { 150 using namespace cv::cuda::device; 151 static const gpu_func_t funcs[] = {BGRA_to_RGB_8u, 0, BGRA_to_RGB_16u, 0, 0, BGRA_to_RGB_32f}; 152 153 GpuMat src = _src.getGpuMat(); 154 155 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 156 CV_Assert( src.channels() == 4 ); 157 158 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3)); 159 GpuMat dst = _dst.getGpuMat(); 160 161 funcs[src.depth()](src, dst, stream); 162 } 163 164 void BGRA_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream) 165 { 166 using namespace cv::cuda::device; 167 static const gpu_func_t funcs[] = {BGRA_to_RGBA_8u, 0, BGRA_to_RGBA_16u, 0, 0, BGRA_to_RGBA_32f}; 168 169 GpuMat src = _src.getGpuMat(); 170 171 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 172 CV_Assert( src.channels() == 4 ); 173 174 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4)); 175 GpuMat dst = _dst.getGpuMat(); 176 177 funcs[src.depth()](src, dst, stream); 178 } 179 180 void BGR_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream) 181 { 182 GpuMat src = _src.getGpuMat(); 183 184 CV_Assert( src.depth() == CV_8U ); 185 CV_Assert( src.channels() == 3 ); 186 187 _dst.create(src.size(), CV_8UC2); 188 GpuMat dst = _dst.getGpuMat(); 189 190 cv::cuda::device::BGR_to_BGR555(src, dst, stream); 191 } 192 193 void BGR_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream) 194 { 195 GpuMat src = _src.getGpuMat(); 196 197 CV_Assert( src.depth() == CV_8U ); 198 CV_Assert( src.channels() == 3 ); 199 200 _dst.create(src.size(), CV_8UC2); 201 GpuMat dst = _dst.getGpuMat(); 202 203 cv::cuda::device::BGR_to_BGR565(src, dst, stream); 204 } 205 206 void RGB_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream) 207 { 208 GpuMat src = _src.getGpuMat(); 209 210 CV_Assert( src.depth() == CV_8U ); 211 CV_Assert( src.channels() == 3 ); 212 213 _dst.create(src.size(), CV_8UC2); 214 GpuMat dst = _dst.getGpuMat(); 215 216 cv::cuda::device::RGB_to_BGR555(src, dst, stream); 217 } 218 219 void RGB_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream) 220 { 221 GpuMat src = _src.getGpuMat(); 222 223 CV_Assert( src.depth() == CV_8U ); 224 CV_Assert( src.channels() == 3 ); 225 226 _dst.create(src.size(), CV_8UC2); 227 GpuMat dst = _dst.getGpuMat(); 228 229 cv::cuda::device::RGB_to_BGR565(src, dst, stream); 230 } 231 232 void BGRA_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream) 233 { 234 GpuMat src = _src.getGpuMat(); 235 236 CV_Assert( src.depth() == CV_8U ); 237 CV_Assert( src.channels() == 4 ); 238 239 _dst.create(src.size(), CV_8UC2); 240 GpuMat dst = _dst.getGpuMat(); 241 242 cv::cuda::device::BGRA_to_BGR555(src, dst, stream); 243 } 244 245 void BGRA_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream) 246 { 247 GpuMat src = _src.getGpuMat(); 248 249 CV_Assert( src.depth() == CV_8U ); 250 CV_Assert( src.channels() == 4 ); 251 252 _dst.create(src.size(), CV_8UC2); 253 GpuMat dst = _dst.getGpuMat(); 254 255 cv::cuda::device::BGRA_to_BGR565(src, dst, stream); 256 } 257 258 void RGBA_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream) 259 { 260 GpuMat src = _src.getGpuMat(); 261 262 CV_Assert( src.depth() == CV_8U ); 263 CV_Assert( src.channels() == 4 ); 264 265 _dst.create(src.size(), CV_8UC2); 266 GpuMat dst = _dst.getGpuMat(); 267 268 cv::cuda::device::RGBA_to_BGR555(src, dst, stream); 269 } 270 271 void RGBA_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream) 272 { 273 GpuMat src = _src.getGpuMat(); 274 275 CV_Assert( src.depth() == CV_8U ); 276 CV_Assert( src.channels() == 4 ); 277 278 _dst.create(src.size(), CV_8UC2); 279 GpuMat dst = _dst.getGpuMat(); 280 281 cv::cuda::device::RGBA_to_BGR565(src, dst, stream); 282 } 283 284 void BGR555_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream) 285 { 286 GpuMat src = _src.getGpuMat(); 287 288 CV_Assert( src.depth() == CV_8U ); 289 CV_Assert( src.channels() == 2 ); 290 291 _dst.create(src.size(), CV_8UC3); 292 GpuMat dst = _dst.getGpuMat(); 293 294 cv::cuda::device::BGR555_to_RGB(src, dst, stream); 295 } 296 297 void BGR565_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream) 298 { 299 GpuMat src = _src.getGpuMat(); 300 301 CV_Assert( src.depth() == CV_8U ); 302 CV_Assert( src.channels() == 2 ); 303 304 _dst.create(src.size(), CV_8UC3); 305 GpuMat dst = _dst.getGpuMat(); 306 307 cv::cuda::device::BGR565_to_RGB(src, dst, stream); 308 } 309 310 void BGR555_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream) 311 { 312 GpuMat src = _src.getGpuMat(); 313 314 CV_Assert( src.depth() == CV_8U ); 315 CV_Assert( src.channels() == 2 ); 316 317 _dst.create(src.size(), CV_8UC3); 318 GpuMat dst = _dst.getGpuMat(); 319 320 cv::cuda::device::BGR555_to_BGR(src, dst, stream); 321 } 322 323 void BGR565_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream) 324 { 325 GpuMat src = _src.getGpuMat(); 326 327 CV_Assert( src.depth() == CV_8U ); 328 CV_Assert( src.channels() == 2 ); 329 330 _dst.create(src.size(), CV_8UC3); 331 GpuMat dst = _dst.getGpuMat(); 332 333 cv::cuda::device::BGR565_to_BGR(src, dst, stream); 334 } 335 336 void BGR555_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream) 337 { 338 GpuMat src = _src.getGpuMat(); 339 340 CV_Assert( src.depth() == CV_8U ); 341 CV_Assert( src.channels() == 2 ); 342 343 _dst.create(src.size(), CV_8UC4); 344 GpuMat dst = _dst.getGpuMat(); 345 346 cv::cuda::device::BGR555_to_RGBA(src, dst, stream); 347 } 348 349 void BGR565_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream) 350 { 351 GpuMat src = _src.getGpuMat(); 352 353 CV_Assert( src.depth() == CV_8U ); 354 CV_Assert( src.channels() == 2 ); 355 356 _dst.create(src.size(), CV_8UC4); 357 GpuMat dst = _dst.getGpuMat(); 358 359 cv::cuda::device::BGR565_to_RGBA(src, dst, stream); 360 } 361 362 void BGR555_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream) 363 { 364 GpuMat src = _src.getGpuMat(); 365 366 CV_Assert( src.depth() == CV_8U ); 367 CV_Assert( src.channels() == 2 ); 368 369 _dst.create(src.size(), CV_8UC4); 370 GpuMat dst = _dst.getGpuMat(); 371 372 cv::cuda::device::BGR555_to_BGRA(src, dst, stream); 373 } 374 375 void BGR565_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream) 376 { 377 GpuMat src = _src.getGpuMat(); 378 379 CV_Assert( src.depth() == CV_8U ); 380 CV_Assert( src.channels() == 2 ); 381 382 _dst.create(src.size(), CV_8UC4); 383 GpuMat dst = _dst.getGpuMat(); 384 385 cv::cuda::device::BGR565_to_BGRA(src, dst, stream); 386 } 387 388 void GRAY_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream) 389 { 390 using namespace cv::cuda::device; 391 static const gpu_func_t funcs[] = {GRAY_to_BGR_8u, 0, GRAY_to_BGR_16u, 0, 0, GRAY_to_BGR_32f}; 392 393 GpuMat src = _src.getGpuMat(); 394 395 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 396 CV_Assert( src.channels() == 1 ); 397 398 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3)); 399 GpuMat dst = _dst.getGpuMat(); 400 401 funcs[src.depth()](src, dst, stream); 402 } 403 404 void GRAY_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream) 405 { 406 using namespace cv::cuda::device; 407 static const gpu_func_t funcs[] = {GRAY_to_BGRA_8u, 0, GRAY_to_BGRA_16u, 0, 0, GRAY_to_BGRA_32f}; 408 409 GpuMat src = _src.getGpuMat(); 410 411 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 412 CV_Assert( src.channels() == 1 ); 413 414 _dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); 415 GpuMat dst = _dst.getGpuMat(); 416 417 funcs[src.depth()](src, dst, stream); 418 } 419 420 void GRAY_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream) 421 { 422 GpuMat src = _src.getGpuMat(); 423 424 CV_Assert( src.depth() == CV_8U ); 425 CV_Assert( src.channels() == 1 ); 426 427 _dst.create(src.size(), CV_8UC2); 428 GpuMat dst = _dst.getGpuMat(); 429 430 cv::cuda::device::GRAY_to_BGR555(src, dst, stream); 431 } 432 433 void GRAY_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream) 434 { 435 GpuMat src = _src.getGpuMat(); 436 437 CV_Assert( src.depth() == CV_8U ); 438 CV_Assert( src.channels() == 1 ); 439 440 _dst.create(src.size(), CV_8UC2); 441 GpuMat dst = _dst.getGpuMat(); 442 443 cv::cuda::device::GRAY_to_BGR565(src, dst, stream); 444 } 445 446 void BGR555_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 447 { 448 GpuMat src = _src.getGpuMat(); 449 450 CV_Assert( src.depth() == CV_8U ); 451 CV_Assert( src.channels() == 2 ); 452 453 _dst.create(src.size(), CV_8UC1); 454 GpuMat dst = _dst.getGpuMat(); 455 456 cv::cuda::device::BGR555_to_GRAY(src, dst, stream); 457 } 458 459 void BGR565_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 460 { 461 GpuMat src = _src.getGpuMat(); 462 463 CV_Assert( src.depth() == CV_8U ); 464 CV_Assert( src.channels() == 2 ); 465 466 _dst.create(src.size(), CV_8UC1); 467 GpuMat dst = _dst.getGpuMat(); 468 469 cv::cuda::device::BGR565_to_GRAY(src, dst, stream); 470 } 471 472 void RGB_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 473 { 474 using namespace cv::cuda::device; 475 static const gpu_func_t funcs[] = {RGB_to_GRAY_8u, 0, RGB_to_GRAY_16u, 0, 0, RGB_to_GRAY_32f}; 476 477 GpuMat src = _src.getGpuMat(); 478 479 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 480 CV_Assert( src.channels() == 3 ); 481 482 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1)); 483 GpuMat dst = _dst.getGpuMat(); 484 485 funcs[src.depth()](src, dst, stream); 486 } 487 488 void BGR_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 489 { 490 using namespace cv::cuda::device; 491 static const gpu_func_t funcs[] = {BGR_to_GRAY_8u, 0, BGR_to_GRAY_16u, 0, 0, BGR_to_GRAY_32f}; 492 493 GpuMat src = _src.getGpuMat(); 494 495 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 496 CV_Assert( src.channels() == 3 ); 497 498 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1)); 499 GpuMat dst = _dst.getGpuMat(); 500 501 funcs[src.depth()](src, dst, stream); 502 } 503 504 void RGBA_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 505 { 506 using namespace cv::cuda::device; 507 static const gpu_func_t funcs[] = {RGBA_to_GRAY_8u, 0, RGBA_to_GRAY_16u, 0, 0, RGBA_to_GRAY_32f}; 508 509 GpuMat src = _src.getGpuMat(); 510 511 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 512 CV_Assert( src.channels() == 4 ); 513 514 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1)); 515 GpuMat dst = _dst.getGpuMat(); 516 517 funcs[src.depth()](src, dst, stream); 518 } 519 520 void BGRA_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream) 521 { 522 using namespace cv::cuda::device; 523 static const gpu_func_t funcs[] = {BGRA_to_GRAY_8u, 0, BGRA_to_GRAY_16u, 0, 0, BGRA_to_GRAY_32f}; 524 525 GpuMat src = _src.getGpuMat(); 526 527 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 528 CV_Assert( src.channels() == 4 ); 529 530 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1)); 531 GpuMat dst = _dst.getGpuMat(); 532 533 funcs[src.depth()](src, dst, stream); 534 } 535 536 void RGB_to_YUV(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 537 { 538 using namespace cv::cuda::device; 539 static const gpu_func_t funcs[2][2][6] = 540 { 541 { 542 {RGB_to_YUV_8u, 0, RGB_to_YUV_16u, 0, 0, RGB_to_YUV_32f}, 543 {RGBA_to_YUV_8u, 0, RGBA_to_YUV_16u, 0, 0, RGBA_to_YUV_32f} 544 }, 545 { 546 {RGB_to_YUV4_8u, 0, RGB_to_YUV4_16u, 0, 0, RGB_to_YUV4_32f}, 547 {RGBA_to_YUV4_8u, 0, RGBA_to_YUV4_16u, 0, 0, RGBA_to_YUV4_32f} 548 } 549 }; 550 551 if (dcn <= 0) dcn = 3; 552 553 GpuMat src = _src.getGpuMat(); 554 555 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 556 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 557 CV_Assert( dcn == 3 || dcn == 4 ); 558 559 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 560 GpuMat dst = _dst.getGpuMat(); 561 562 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 563 } 564 565 void BGR_to_YUV(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 566 { 567 using namespace cv::cuda::device; 568 static const gpu_func_t funcs[2][2][6] = 569 { 570 { 571 {BGR_to_YUV_8u, 0, BGR_to_YUV_16u, 0, 0, BGR_to_YUV_32f}, 572 {BGRA_to_YUV_8u, 0, BGRA_to_YUV_16u, 0, 0, BGRA_to_YUV_32f} 573 }, 574 { 575 {BGR_to_YUV4_8u, 0, BGR_to_YUV4_16u, 0, 0, BGR_to_YUV4_32f}, 576 {BGRA_to_YUV4_8u, 0, BGRA_to_YUV4_16u, 0, 0, BGRA_to_YUV4_32f} 577 } 578 }; 579 580 if (dcn <= 0) dcn = 3; 581 582 GpuMat src = _src.getGpuMat(); 583 584 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 585 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 586 CV_Assert( dcn == 3 || dcn == 4 ); 587 588 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 589 GpuMat dst = _dst.getGpuMat(); 590 591 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 592 } 593 594 void YUV_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 595 { 596 using namespace cv::cuda::device; 597 static const gpu_func_t funcs[2][2][6] = 598 { 599 { 600 {YUV_to_RGB_8u, 0, YUV_to_RGB_16u, 0, 0, YUV_to_RGB_32f}, 601 {YUV4_to_RGB_8u, 0, YUV4_to_RGB_16u, 0, 0, YUV4_to_RGB_32f} 602 }, 603 { 604 {YUV_to_RGBA_8u, 0, YUV_to_RGBA_16u, 0, 0, YUV_to_RGBA_32f}, 605 {YUV4_to_RGBA_8u, 0, YUV4_to_RGBA_16u, 0, 0, YUV4_to_RGBA_32f} 606 } 607 }; 608 609 if (dcn <= 0) dcn = 3; 610 611 GpuMat src = _src.getGpuMat(); 612 613 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 614 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 615 CV_Assert( dcn == 3 || dcn == 4 ); 616 617 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 618 GpuMat dst = _dst.getGpuMat(); 619 620 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 621 } 622 623 void YUV_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 624 { 625 using namespace cv::cuda::device; 626 static const gpu_func_t funcs[2][2][6] = 627 { 628 { 629 {YUV_to_BGR_8u, 0, YUV_to_BGR_16u, 0, 0, YUV_to_BGR_32f}, 630 {YUV4_to_BGR_8u, 0, YUV4_to_BGR_16u, 0, 0, YUV4_to_BGR_32f} 631 }, 632 { 633 {YUV_to_BGRA_8u, 0, YUV_to_BGRA_16u, 0, 0, YUV_to_BGRA_32f}, 634 {YUV4_to_BGRA_8u, 0, YUV4_to_BGRA_16u, 0, 0, YUV4_to_BGRA_32f} 635 } 636 }; 637 638 if (dcn <= 0) dcn = 3; 639 640 GpuMat src = _src.getGpuMat(); 641 642 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 643 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 644 CV_Assert( dcn == 3 || dcn == 4 ); 645 646 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 647 GpuMat dst = _dst.getGpuMat(); 648 649 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 650 } 651 652 void RGB_to_YCrCb(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 653 { 654 using namespace cv::cuda::device; 655 static const gpu_func_t funcs[2][2][6] = 656 { 657 { 658 {RGB_to_YCrCb_8u, 0, RGB_to_YCrCb_16u, 0, 0, RGB_to_YCrCb_32f}, 659 {RGBA_to_YCrCb_8u, 0, RGBA_to_YCrCb_16u, 0, 0, RGBA_to_YCrCb_32f} 660 }, 661 { 662 {RGB_to_YCrCb4_8u, 0, RGB_to_YCrCb4_16u, 0, 0, RGB_to_YCrCb4_32f}, 663 {RGBA_to_YCrCb4_8u, 0, RGBA_to_YCrCb4_16u, 0, 0, RGBA_to_YCrCb4_32f} 664 } 665 }; 666 667 if (dcn <= 0) dcn = 3; 668 669 GpuMat src = _src.getGpuMat(); 670 671 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 672 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 673 CV_Assert( dcn == 3 || dcn == 4 ); 674 675 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 676 GpuMat dst = _dst.getGpuMat(); 677 678 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 679 } 680 681 void BGR_to_YCrCb(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 682 { 683 using namespace cv::cuda::device; 684 static const gpu_func_t funcs[2][2][6] = 685 { 686 { 687 {BGR_to_YCrCb_8u, 0, BGR_to_YCrCb_16u, 0, 0, BGR_to_YCrCb_32f}, 688 {BGRA_to_YCrCb_8u, 0, BGRA_to_YCrCb_16u, 0, 0, BGRA_to_YCrCb_32f} 689 }, 690 { 691 {BGR_to_YCrCb4_8u, 0, BGR_to_YCrCb4_16u, 0, 0, BGR_to_YCrCb4_32f}, 692 {BGRA_to_YCrCb4_8u, 0, BGRA_to_YCrCb4_16u, 0, 0, BGRA_to_YCrCb4_32f} 693 } 694 }; 695 696 if (dcn <= 0) dcn = 3; 697 698 GpuMat src = _src.getGpuMat(); 699 700 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 701 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 702 CV_Assert( dcn == 3 || dcn == 4 ); 703 704 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 705 GpuMat dst = _dst.getGpuMat(); 706 707 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 708 } 709 710 void YCrCb_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 711 { 712 using namespace cv::cuda::device; 713 static const gpu_func_t funcs[2][2][6] = 714 { 715 { 716 {YCrCb_to_RGB_8u, 0, YCrCb_to_RGB_16u, 0, 0, YCrCb_to_RGB_32f}, 717 {YCrCb4_to_RGB_8u, 0, YCrCb4_to_RGB_16u, 0, 0, YCrCb4_to_RGB_32f} 718 }, 719 { 720 {YCrCb_to_RGBA_8u, 0, YCrCb_to_RGBA_16u, 0, 0, YCrCb_to_RGBA_32f}, 721 {YCrCb4_to_RGBA_8u, 0, YCrCb4_to_RGBA_16u, 0, 0, YCrCb4_to_RGBA_32f} 722 } 723 }; 724 725 if (dcn <= 0) dcn = 3; 726 727 GpuMat src = _src.getGpuMat(); 728 729 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 730 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 731 CV_Assert( dcn == 3 || dcn == 4 ); 732 733 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 734 GpuMat dst = _dst.getGpuMat(); 735 736 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 737 } 738 739 void YCrCb_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 740 { 741 using namespace cv::cuda::device; 742 static const gpu_func_t funcs[2][2][6] = 743 { 744 { 745 {YCrCb_to_BGR_8u, 0, YCrCb_to_BGR_16u, 0, 0, YCrCb_to_BGR_32f}, 746 {YCrCb4_to_BGR_8u, 0, YCrCb4_to_BGR_16u, 0, 0, YCrCb4_to_BGR_32f} 747 }, 748 { 749 {YCrCb_to_BGRA_8u, 0, YCrCb_to_BGRA_16u, 0, 0, YCrCb_to_BGRA_32f}, 750 {YCrCb4_to_BGRA_8u, 0, YCrCb4_to_BGRA_16u, 0, 0, YCrCb4_to_BGRA_32f} 751 } 752 }; 753 754 if (dcn <= 0) dcn = 3; 755 756 GpuMat src = _src.getGpuMat(); 757 758 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 759 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 760 CV_Assert( dcn == 3 || dcn == 4 ); 761 762 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 763 GpuMat dst = _dst.getGpuMat(); 764 765 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 766 } 767 768 void RGB_to_XYZ(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 769 { 770 using namespace cv::cuda::device; 771 static const gpu_func_t funcs[2][2][6] = 772 { 773 { 774 {RGB_to_XYZ_8u, 0, RGB_to_XYZ_16u, 0, 0, RGB_to_XYZ_32f}, 775 {RGBA_to_XYZ_8u, 0, RGBA_to_XYZ_16u, 0, 0, RGBA_to_XYZ_32f} 776 }, 777 { 778 {RGB_to_XYZ4_8u, 0, RGB_to_XYZ4_16u, 0, 0, RGB_to_XYZ4_32f}, 779 {RGBA_to_XYZ4_8u, 0, RGBA_to_XYZ4_16u, 0, 0, RGBA_to_XYZ4_32f} 780 } 781 }; 782 783 if (dcn <= 0) dcn = 3; 784 785 GpuMat src = _src.getGpuMat(); 786 787 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 788 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 789 CV_Assert( dcn == 3 || dcn == 4 ); 790 791 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 792 GpuMat dst = _dst.getGpuMat(); 793 794 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 795 } 796 797 void BGR_to_XYZ(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 798 { 799 using namespace cv::cuda::device; 800 static const gpu_func_t funcs[2][2][6] = 801 { 802 { 803 {BGR_to_XYZ_8u, 0, BGR_to_XYZ_16u, 0, 0, BGR_to_XYZ_32f}, 804 {BGRA_to_XYZ_8u, 0, BGRA_to_XYZ_16u, 0, 0, BGRA_to_XYZ_32f} 805 }, 806 { 807 {BGR_to_XYZ4_8u, 0, BGR_to_XYZ4_16u, 0, 0, BGR_to_XYZ4_32f}, 808 {BGRA_to_XYZ4_8u, 0, BGRA_to_XYZ4_16u, 0, 0, BGRA_to_XYZ4_32f} 809 } 810 }; 811 812 if (dcn <= 0) dcn = 3; 813 814 GpuMat src = _src.getGpuMat(); 815 816 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 817 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 818 CV_Assert( dcn == 3 || dcn == 4 ); 819 820 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 821 GpuMat dst = _dst.getGpuMat(); 822 823 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 824 } 825 826 void XYZ_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 827 { 828 using namespace cv::cuda::device; 829 static const gpu_func_t funcs[2][2][6] = 830 { 831 { 832 {XYZ_to_RGB_8u, 0, XYZ_to_RGB_16u, 0, 0, XYZ_to_RGB_32f}, 833 {XYZ4_to_RGB_8u, 0, XYZ4_to_RGB_16u, 0, 0, XYZ4_to_RGB_32f} 834 }, 835 { 836 {XYZ_to_RGBA_8u, 0, XYZ_to_RGBA_16u, 0, 0, XYZ_to_RGBA_32f}, 837 {XYZ4_to_RGBA_8u, 0, XYZ4_to_RGBA_16u, 0, 0, XYZ4_to_RGBA_32f} 838 } 839 }; 840 841 if (dcn <= 0) dcn = 3; 842 843 GpuMat src = _src.getGpuMat(); 844 845 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 846 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 847 CV_Assert( dcn == 3 || dcn == 4 ); 848 849 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 850 GpuMat dst = _dst.getGpuMat(); 851 852 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 853 } 854 855 void XYZ_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 856 { 857 using namespace cv::cuda::device; 858 static const gpu_func_t funcs[2][2][6] = 859 { 860 { 861 {XYZ_to_BGR_8u, 0, XYZ_to_BGR_16u, 0, 0, XYZ_to_BGR_32f}, 862 {XYZ4_to_BGR_8u, 0, XYZ4_to_BGR_16u, 0, 0, XYZ4_to_BGR_32f} 863 }, 864 { 865 {XYZ_to_BGRA_8u, 0, XYZ_to_BGRA_16u, 0, 0, XYZ_to_BGRA_32f}, 866 {XYZ4_to_BGRA_8u, 0, XYZ4_to_BGRA_16u, 0, 0, XYZ4_to_BGRA_32f} 867 } 868 }; 869 870 if (dcn <= 0) dcn = 3; 871 872 GpuMat src = _src.getGpuMat(); 873 874 CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F ); 875 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 876 CV_Assert( dcn == 3 || dcn == 4 ); 877 878 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 879 GpuMat dst = _dst.getGpuMat(); 880 881 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 882 } 883 884 void RGB_to_HSV(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 885 { 886 using namespace cv::cuda::device; 887 static const gpu_func_t funcs[2][2][6] = 888 { 889 { 890 {RGB_to_HSV_8u, 0, 0, 0, 0, RGB_to_HSV_32f}, 891 {RGBA_to_HSV_8u, 0, 0, 0, 0, RGBA_to_HSV_32f}, 892 }, 893 { 894 {RGB_to_HSV4_8u, 0, 0, 0, 0, RGB_to_HSV4_32f}, 895 {RGBA_to_HSV4_8u, 0, 0, 0, 0, RGBA_to_HSV4_32f}, 896 } 897 }; 898 899 if (dcn <= 0) dcn = 3; 900 901 GpuMat src = _src.getGpuMat(); 902 903 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 904 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 905 CV_Assert( dcn == 3 || dcn == 4 ); 906 907 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 908 GpuMat dst = _dst.getGpuMat(); 909 910 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 911 } 912 913 void BGR_to_HSV(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 914 { 915 using namespace cv::cuda::device; 916 static const gpu_func_t funcs[2][2][6] = 917 { 918 { 919 {BGR_to_HSV_8u, 0, 0, 0, 0, BGR_to_HSV_32f}, 920 {BGRA_to_HSV_8u, 0, 0, 0, 0, BGRA_to_HSV_32f} 921 }, 922 { 923 {BGR_to_HSV4_8u, 0, 0, 0, 0, BGR_to_HSV4_32f}, 924 {BGRA_to_HSV4_8u, 0, 0, 0, 0, BGRA_to_HSV4_32f} 925 } 926 }; 927 928 if (dcn <= 0) dcn = 3; 929 930 GpuMat src = _src.getGpuMat(); 931 932 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 933 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 934 CV_Assert( dcn == 3 || dcn == 4 ); 935 936 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 937 GpuMat dst = _dst.getGpuMat(); 938 939 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 940 } 941 942 void HSV_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 943 { 944 using namespace cv::cuda::device; 945 static const gpu_func_t funcs[2][2][6] = 946 { 947 { 948 {HSV_to_RGB_8u, 0, 0, 0, 0, HSV_to_RGB_32f}, 949 {HSV4_to_RGB_8u, 0, 0, 0, 0, HSV4_to_RGB_32f} 950 }, 951 { 952 {HSV_to_RGBA_8u, 0, 0, 0, 0, HSV_to_RGBA_32f}, 953 {HSV4_to_RGBA_8u, 0, 0, 0, 0, HSV4_to_RGBA_32f} 954 } 955 }; 956 957 if (dcn <= 0) dcn = 3; 958 959 GpuMat src = _src.getGpuMat(); 960 961 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 962 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 963 CV_Assert( dcn == 3 || dcn == 4 ); 964 965 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 966 GpuMat dst = _dst.getGpuMat(); 967 968 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 969 } 970 971 void HSV_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 972 { 973 using namespace cv::cuda::device; 974 static const gpu_func_t funcs[2][2][6] = 975 { 976 { 977 {HSV_to_BGR_8u, 0, 0, 0, 0, HSV_to_BGR_32f}, 978 {HSV4_to_BGR_8u, 0, 0, 0, 0, HSV4_to_BGR_32f} 979 }, 980 { 981 {HSV_to_BGRA_8u, 0, 0, 0, 0, HSV_to_BGRA_32f}, 982 {HSV4_to_BGRA_8u, 0, 0, 0, 0, HSV4_to_BGRA_32f} 983 } 984 }; 985 986 if (dcn <= 0) dcn = 3; 987 988 GpuMat src = _src.getGpuMat(); 989 990 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 991 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 992 CV_Assert( dcn == 3 || dcn == 4 ); 993 994 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 995 GpuMat dst = _dst.getGpuMat(); 996 997 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 998 } 999 1000 void RGB_to_HLS(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1001 { 1002 using namespace cv::cuda::device; 1003 static const gpu_func_t funcs[2][2][6] = 1004 { 1005 { 1006 {RGB_to_HLS_8u, 0, 0, 0, 0, RGB_to_HLS_32f}, 1007 {RGBA_to_HLS_8u, 0, 0, 0, 0, RGBA_to_HLS_32f}, 1008 }, 1009 { 1010 {RGB_to_HLS4_8u, 0, 0, 0, 0, RGB_to_HLS4_32f}, 1011 {RGBA_to_HLS4_8u, 0, 0, 0, 0, RGBA_to_HLS4_32f}, 1012 } 1013 }; 1014 1015 if (dcn <= 0) dcn = 3; 1016 1017 GpuMat src = _src.getGpuMat(); 1018 1019 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1020 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1021 CV_Assert( dcn == 3 || dcn == 4 ); 1022 1023 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1024 GpuMat dst = _dst.getGpuMat(); 1025 1026 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1027 } 1028 1029 void BGR_to_HLS(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1030 { 1031 using namespace cv::cuda::device; 1032 static const gpu_func_t funcs[2][2][6] = 1033 { 1034 { 1035 {BGR_to_HLS_8u, 0, 0, 0, 0, BGR_to_HLS_32f}, 1036 {BGRA_to_HLS_8u, 0, 0, 0, 0, BGRA_to_HLS_32f} 1037 }, 1038 { 1039 {BGR_to_HLS4_8u, 0, 0, 0, 0, BGR_to_HLS4_32f}, 1040 {BGRA_to_HLS4_8u, 0, 0, 0, 0, BGRA_to_HLS4_32f} 1041 } 1042 }; 1043 1044 if (dcn <= 0) dcn = 3; 1045 1046 GpuMat src = _src.getGpuMat(); 1047 1048 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1049 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1050 CV_Assert( dcn == 3 || dcn == 4 ); 1051 1052 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1053 GpuMat dst = _dst.getGpuMat(); 1054 1055 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1056 } 1057 1058 void HLS_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1059 { 1060 using namespace cv::cuda::device; 1061 static const gpu_func_t funcs[2][2][6] = 1062 { 1063 { 1064 {HLS_to_RGB_8u, 0, 0, 0, 0, HLS_to_RGB_32f}, 1065 {HLS4_to_RGB_8u, 0, 0, 0, 0, HLS4_to_RGB_32f} 1066 }, 1067 { 1068 {HLS_to_RGBA_8u, 0, 0, 0, 0, HLS_to_RGBA_32f}, 1069 {HLS4_to_RGBA_8u, 0, 0, 0, 0, HLS4_to_RGBA_32f} 1070 } 1071 }; 1072 1073 if (dcn <= 0) dcn = 3; 1074 1075 GpuMat src = _src.getGpuMat(); 1076 1077 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1078 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1079 CV_Assert( dcn == 3 || dcn == 4 ); 1080 1081 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1082 GpuMat dst = _dst.getGpuMat(); 1083 1084 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1085 } 1086 1087 void HLS_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1088 { 1089 using namespace cv::cuda::device; 1090 static const gpu_func_t funcs[2][2][6] = 1091 { 1092 { 1093 {HLS_to_BGR_8u, 0, 0, 0, 0, HLS_to_BGR_32f}, 1094 {HLS4_to_BGR_8u, 0, 0, 0, 0, HLS4_to_BGR_32f} 1095 }, 1096 { 1097 {HLS_to_BGRA_8u, 0, 0, 0, 0, HLS_to_BGRA_32f}, 1098 {HLS4_to_BGRA_8u, 0, 0, 0, 0, HLS4_to_BGRA_32f} 1099 } 1100 }; 1101 1102 if (dcn <= 0) dcn = 3; 1103 1104 GpuMat src = _src.getGpuMat(); 1105 1106 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1107 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1108 CV_Assert( dcn == 3 || dcn == 4 ); 1109 1110 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1111 GpuMat dst = _dst.getGpuMat(); 1112 1113 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1114 } 1115 1116 void RGB_to_HSV_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1117 { 1118 using namespace cv::cuda::device; 1119 static const gpu_func_t funcs[2][2][6] = 1120 { 1121 { 1122 {RGB_to_HSV_FULL_8u, 0, 0, 0, 0, RGB_to_HSV_FULL_32f}, 1123 {RGBA_to_HSV_FULL_8u, 0, 0, 0, 0, RGBA_to_HSV_FULL_32f}, 1124 }, 1125 { 1126 {RGB_to_HSV4_FULL_8u, 0, 0, 0, 0, RGB_to_HSV4_FULL_32f}, 1127 {RGBA_to_HSV4_FULL_8u, 0, 0, 0, 0, RGBA_to_HSV4_FULL_32f}, 1128 } 1129 }; 1130 1131 if (dcn <= 0) dcn = 3; 1132 1133 GpuMat src = _src.getGpuMat(); 1134 1135 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1136 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1137 CV_Assert( dcn == 3 || dcn == 4 ); 1138 1139 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1140 GpuMat dst = _dst.getGpuMat(); 1141 1142 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1143 } 1144 1145 void BGR_to_HSV_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1146 { 1147 using namespace cv::cuda::device; 1148 static const gpu_func_t funcs[2][2][6] = 1149 { 1150 { 1151 {BGR_to_HSV_FULL_8u, 0, 0, 0, 0, BGR_to_HSV_FULL_32f}, 1152 {BGRA_to_HSV_FULL_8u, 0, 0, 0, 0, BGRA_to_HSV_FULL_32f} 1153 }, 1154 { 1155 {BGR_to_HSV4_FULL_8u, 0, 0, 0, 0, BGR_to_HSV4_FULL_32f}, 1156 {BGRA_to_HSV4_FULL_8u, 0, 0, 0, 0, BGRA_to_HSV4_FULL_32f} 1157 } 1158 }; 1159 1160 if (dcn <= 0) dcn = 3; 1161 1162 GpuMat src = _src.getGpuMat(); 1163 1164 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1165 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1166 CV_Assert( dcn == 3 || dcn == 4 ); 1167 1168 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1169 GpuMat dst = _dst.getGpuMat(); 1170 1171 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1172 } 1173 1174 void HSV_to_RGB_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1175 { 1176 using namespace cv::cuda::device; 1177 static const gpu_func_t funcs[2][2][6] = 1178 { 1179 { 1180 {HSV_to_RGB_FULL_8u, 0, 0, 0, 0, HSV_to_RGB_FULL_32f}, 1181 {HSV4_to_RGB_FULL_8u, 0, 0, 0, 0, HSV4_to_RGB_FULL_32f} 1182 }, 1183 { 1184 {HSV_to_RGBA_FULL_8u, 0, 0, 0, 0, HSV_to_RGBA_FULL_32f}, 1185 {HSV4_to_RGBA_FULL_8u, 0, 0, 0, 0, HSV4_to_RGBA_FULL_32f} 1186 } 1187 }; 1188 1189 if (dcn <= 0) dcn = 3; 1190 1191 GpuMat src = _src.getGpuMat(); 1192 1193 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1194 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1195 CV_Assert( dcn == 3 || dcn == 4 ); 1196 1197 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1198 GpuMat dst = _dst.getGpuMat(); 1199 1200 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1201 } 1202 1203 void HSV_to_BGR_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1204 { 1205 using namespace cv::cuda::device; 1206 static const gpu_func_t funcs[2][2][6] = 1207 { 1208 { 1209 {HSV_to_BGR_FULL_8u, 0, 0, 0, 0, HSV_to_BGR_FULL_32f}, 1210 {HSV4_to_BGR_FULL_8u, 0, 0, 0, 0, HSV4_to_BGR_FULL_32f} 1211 }, 1212 { 1213 {HSV_to_BGRA_FULL_8u, 0, 0, 0, 0, HSV_to_BGRA_FULL_32f}, 1214 {HSV4_to_BGRA_FULL_8u, 0, 0, 0, 0, HSV4_to_BGRA_FULL_32f} 1215 } 1216 }; 1217 1218 if (dcn <= 0) dcn = 3; 1219 1220 GpuMat src = _src.getGpuMat(); 1221 1222 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1223 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1224 CV_Assert( dcn == 3 || dcn == 4 ); 1225 1226 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1227 GpuMat dst = _dst.getGpuMat(); 1228 1229 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1230 } 1231 1232 void RGB_to_HLS_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1233 { 1234 using namespace cv::cuda::device; 1235 static const gpu_func_t funcs[2][2][6] = 1236 { 1237 { 1238 {RGB_to_HLS_FULL_8u, 0, 0, 0, 0, RGB_to_HLS_FULL_32f}, 1239 {RGBA_to_HLS_FULL_8u, 0, 0, 0, 0, RGBA_to_HLS_FULL_32f}, 1240 }, 1241 { 1242 {RGB_to_HLS4_FULL_8u, 0, 0, 0, 0, RGB_to_HLS4_FULL_32f}, 1243 {RGBA_to_HLS4_FULL_8u, 0, 0, 0, 0, RGBA_to_HLS4_FULL_32f}, 1244 } 1245 }; 1246 1247 if (dcn <= 0) dcn = 3; 1248 1249 GpuMat src = _src.getGpuMat(); 1250 1251 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1252 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1253 CV_Assert( dcn == 3 || dcn == 4 ); 1254 1255 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1256 GpuMat dst = _dst.getGpuMat(); 1257 1258 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1259 } 1260 1261 void BGR_to_HLS_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1262 { 1263 using namespace cv::cuda::device; 1264 static const gpu_func_t funcs[2][2][6] = 1265 { 1266 { 1267 {BGR_to_HLS_FULL_8u, 0, 0, 0, 0, BGR_to_HLS_FULL_32f}, 1268 {BGRA_to_HLS_FULL_8u, 0, 0, 0, 0, BGRA_to_HLS_FULL_32f} 1269 }, 1270 { 1271 {BGR_to_HLS4_FULL_8u, 0, 0, 0, 0, BGR_to_HLS4_FULL_32f}, 1272 {BGRA_to_HLS4_FULL_8u, 0, 0, 0, 0, BGRA_to_HLS4_FULL_32f} 1273 } 1274 }; 1275 1276 if (dcn <= 0) dcn = 3; 1277 1278 GpuMat src = _src.getGpuMat(); 1279 1280 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1281 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1282 CV_Assert( dcn == 3 || dcn == 4 ); 1283 1284 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1285 GpuMat dst = _dst.getGpuMat(); 1286 1287 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1288 } 1289 1290 void HLS_to_RGB_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1291 { 1292 using namespace cv::cuda::device; 1293 static const gpu_func_t funcs[2][2][6] = 1294 { 1295 { 1296 {HLS_to_RGB_FULL_8u, 0, 0, 0, 0, HLS_to_RGB_FULL_32f}, 1297 {HLS4_to_RGB_FULL_8u, 0, 0, 0, 0, HLS4_to_RGB_FULL_32f} 1298 }, 1299 { 1300 {HLS_to_RGBA_FULL_8u, 0, 0, 0, 0, HLS_to_RGBA_FULL_32f}, 1301 {HLS4_to_RGBA_FULL_8u, 0, 0, 0, 0, HLS4_to_RGBA_FULL_32f} 1302 } 1303 }; 1304 1305 if (dcn <= 0) dcn = 3; 1306 1307 GpuMat src = _src.getGpuMat(); 1308 1309 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1310 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1311 CV_Assert( dcn == 3 || dcn == 4 ); 1312 1313 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1314 GpuMat dst = _dst.getGpuMat(); 1315 1316 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1317 } 1318 1319 void HLS_to_BGR_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1320 { 1321 using namespace cv::cuda::device; 1322 static const gpu_func_t funcs[2][2][6] = 1323 { 1324 { 1325 {HLS_to_BGR_FULL_8u, 0, 0, 0, 0, HLS_to_BGR_FULL_32f}, 1326 {HLS4_to_BGR_FULL_8u, 0, 0, 0, 0, HLS4_to_BGR_FULL_32f} 1327 }, 1328 { 1329 {HLS_to_BGRA_FULL_8u, 0, 0, 0, 0, HLS_to_BGRA_FULL_32f}, 1330 {HLS4_to_BGRA_FULL_8u, 0, 0, 0, 0, HLS4_to_BGRA_FULL_32f} 1331 } 1332 }; 1333 1334 if (dcn <= 0) dcn = 3; 1335 1336 GpuMat src = _src.getGpuMat(); 1337 1338 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1339 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1340 CV_Assert( dcn == 3 || dcn == 4 ); 1341 1342 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1343 GpuMat dst = _dst.getGpuMat(); 1344 1345 funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream); 1346 } 1347 1348 void BGR_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1349 { 1350 using namespace cv::cuda::device; 1351 static const gpu_func_t funcs[2][2][2] = 1352 { 1353 { 1354 {BGR_to_Lab_8u, BGR_to_Lab_32f}, 1355 {BGRA_to_Lab_8u, BGRA_to_Lab_32f} 1356 }, 1357 { 1358 {BGR_to_Lab4_8u, BGR_to_Lab4_32f}, 1359 {BGRA_to_Lab4_8u, BGRA_to_Lab4_32f} 1360 } 1361 }; 1362 1363 if (dcn <= 0) dcn = 3; 1364 1365 GpuMat src = _src.getGpuMat(); 1366 1367 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1368 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1369 CV_Assert( dcn == 3 || dcn == 4 ); 1370 1371 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1372 GpuMat dst = _dst.getGpuMat(); 1373 1374 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1375 } 1376 1377 void RGB_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1378 { 1379 using namespace cv::cuda::device; 1380 static const gpu_func_t funcs[2][2][2] = 1381 { 1382 { 1383 {RGB_to_Lab_8u, RGB_to_Lab_32f}, 1384 {RGBA_to_Lab_8u, RGBA_to_Lab_32f} 1385 }, 1386 { 1387 {RGB_to_Lab4_8u, RGB_to_Lab4_32f}, 1388 {RGBA_to_Lab4_8u, RGBA_to_Lab4_32f} 1389 } 1390 }; 1391 1392 if (dcn <= 0) dcn = 3; 1393 1394 GpuMat src = _src.getGpuMat(); 1395 1396 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1397 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1398 CV_Assert( dcn == 3 || dcn == 4 ); 1399 1400 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1401 GpuMat dst = _dst.getGpuMat(); 1402 1403 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1404 } 1405 1406 void LBGR_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1407 { 1408 using namespace cv::cuda::device; 1409 static const gpu_func_t funcs[2][2][2] = 1410 { 1411 { 1412 {LBGR_to_Lab_8u, LBGR_to_Lab_32f}, 1413 {LBGRA_to_Lab_8u, LBGRA_to_Lab_32f} 1414 }, 1415 { 1416 {LBGR_to_Lab4_8u, LBGR_to_Lab4_32f}, 1417 {LBGRA_to_Lab4_8u, LBGRA_to_Lab4_32f} 1418 } 1419 }; 1420 1421 if (dcn <= 0) dcn = 3; 1422 1423 GpuMat src = _src.getGpuMat(); 1424 1425 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1426 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1427 CV_Assert( dcn == 3 || dcn == 4 ); 1428 1429 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1430 GpuMat dst = _dst.getGpuMat(); 1431 1432 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1433 } 1434 1435 void LRGB_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1436 { 1437 using namespace cv::cuda::device; 1438 static const gpu_func_t funcs[2][2][2] = 1439 { 1440 { 1441 {LRGB_to_Lab_8u, LRGB_to_Lab_32f}, 1442 {LRGBA_to_Lab_8u, LRGBA_to_Lab_32f} 1443 }, 1444 { 1445 {LRGB_to_Lab4_8u, LRGB_to_Lab4_32f}, 1446 {LRGBA_to_Lab4_8u, LRGBA_to_Lab4_32f} 1447 } 1448 }; 1449 1450 if (dcn <= 0) dcn = 3; 1451 1452 GpuMat src = _src.getGpuMat(); 1453 1454 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1455 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1456 CV_Assert( dcn == 3 || dcn == 4 ); 1457 1458 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1459 GpuMat dst = _dst.getGpuMat(); 1460 1461 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1462 } 1463 1464 void Lab_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1465 { 1466 using namespace cv::cuda::device; 1467 static const gpu_func_t funcs[2][2][2] = 1468 { 1469 { 1470 {Lab_to_BGR_8u, Lab_to_BGR_32f}, 1471 {Lab4_to_BGR_8u, Lab4_to_BGR_32f} 1472 }, 1473 { 1474 {Lab_to_BGRA_8u, Lab_to_BGRA_32f}, 1475 {Lab4_to_BGRA_8u, Lab4_to_BGRA_32f} 1476 } 1477 }; 1478 1479 if (dcn <= 0) dcn = 3; 1480 1481 GpuMat src = _src.getGpuMat(); 1482 1483 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1484 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1485 CV_Assert( dcn == 3 || dcn == 4 ); 1486 1487 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1488 GpuMat dst = _dst.getGpuMat(); 1489 1490 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1491 } 1492 1493 void Lab_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1494 { 1495 using namespace cv::cuda::device; 1496 static const gpu_func_t funcs[2][2][2] = 1497 { 1498 { 1499 {Lab_to_RGB_8u, Lab_to_RGB_32f}, 1500 {Lab4_to_RGB_8u, Lab4_to_RGB_32f} 1501 }, 1502 { 1503 {Lab_to_RGBA_8u, Lab_to_RGBA_32f}, 1504 {Lab4_to_RGBA_8u, Lab4_to_RGBA_32f} 1505 } 1506 }; 1507 1508 if (dcn <= 0) dcn = 3; 1509 1510 GpuMat src = _src.getGpuMat(); 1511 1512 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1513 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1514 CV_Assert( dcn == 3 || dcn == 4 ); 1515 1516 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1517 GpuMat dst = _dst.getGpuMat(); 1518 1519 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1520 } 1521 1522 void Lab_to_LBGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1523 { 1524 using namespace cv::cuda::device; 1525 static const gpu_func_t funcs[2][2][2] = 1526 { 1527 { 1528 {Lab_to_LBGR_8u, Lab_to_LBGR_32f}, 1529 {Lab4_to_LBGR_8u, Lab4_to_LBGR_32f} 1530 }, 1531 { 1532 {Lab_to_LBGRA_8u, Lab_to_LBGRA_32f}, 1533 {Lab4_to_LBGRA_8u, Lab4_to_LBGRA_32f} 1534 } 1535 }; 1536 1537 if (dcn <= 0) dcn = 3; 1538 1539 GpuMat src = _src.getGpuMat(); 1540 1541 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1542 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1543 CV_Assert( dcn == 3 || dcn == 4 ); 1544 1545 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1546 GpuMat dst = _dst.getGpuMat(); 1547 1548 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1549 } 1550 1551 void Lab_to_LRGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1552 { 1553 using namespace cv::cuda::device; 1554 static const gpu_func_t funcs[2][2][2] = 1555 { 1556 { 1557 {Lab_to_LRGB_8u, Lab_to_LRGB_32f}, 1558 {Lab4_to_LRGB_8u, Lab4_to_LRGB_32f} 1559 }, 1560 { 1561 {Lab_to_LRGBA_8u, Lab_to_LRGBA_32f}, 1562 {Lab4_to_LRGBA_8u, Lab4_to_LRGBA_32f} 1563 } 1564 }; 1565 1566 if (dcn <= 0) dcn = 3; 1567 1568 GpuMat src = _src.getGpuMat(); 1569 1570 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1571 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1572 CV_Assert( dcn == 3 || dcn == 4 ); 1573 1574 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1575 GpuMat dst = _dst.getGpuMat(); 1576 1577 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1578 } 1579 1580 void BGR_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1581 { 1582 using namespace cv::cuda::device; 1583 static const gpu_func_t funcs[2][2][2] = 1584 { 1585 { 1586 {BGR_to_Luv_8u, BGR_to_Luv_32f}, 1587 {BGRA_to_Luv_8u, BGRA_to_Luv_32f} 1588 }, 1589 { 1590 {BGR_to_Luv4_8u, BGR_to_Luv4_32f}, 1591 {BGRA_to_Luv4_8u, BGRA_to_Luv4_32f} 1592 } 1593 }; 1594 1595 if (dcn <= 0) dcn = 3; 1596 1597 GpuMat src = _src.getGpuMat(); 1598 1599 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1600 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1601 CV_Assert( dcn == 3 || dcn == 4 ); 1602 1603 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1604 GpuMat dst = _dst.getGpuMat(); 1605 1606 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1607 } 1608 1609 void RGB_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1610 { 1611 using namespace cv::cuda::device; 1612 static const gpu_func_t funcs[2][2][2] = 1613 { 1614 { 1615 {RGB_to_Luv_8u, RGB_to_Luv_32f}, 1616 {RGBA_to_Luv_8u, RGBA_to_Luv_32f} 1617 }, 1618 { 1619 {RGB_to_Luv4_8u, RGB_to_Luv4_32f}, 1620 {RGBA_to_Luv4_8u, RGBA_to_Luv4_32f} 1621 } 1622 }; 1623 1624 if (dcn <= 0) dcn = 3; 1625 1626 GpuMat src = _src.getGpuMat(); 1627 1628 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1629 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1630 CV_Assert( dcn == 3 || dcn == 4 ); 1631 1632 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1633 GpuMat dst = _dst.getGpuMat(); 1634 1635 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1636 } 1637 1638 void LBGR_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1639 { 1640 using namespace cv::cuda::device; 1641 static const gpu_func_t funcs[2][2][2] = 1642 { 1643 { 1644 {LBGR_to_Luv_8u, LBGR_to_Luv_32f}, 1645 {LBGRA_to_Luv_8u, LBGRA_to_Luv_32f} 1646 }, 1647 { 1648 {LBGR_to_Luv4_8u, LBGR_to_Luv4_32f}, 1649 {LBGRA_to_Luv4_8u, LBGRA_to_Luv4_32f} 1650 } 1651 }; 1652 1653 if (dcn <= 0) dcn = 3; 1654 1655 GpuMat src = _src.getGpuMat(); 1656 1657 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1658 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1659 CV_Assert( dcn == 3 || dcn == 4 ); 1660 1661 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1662 GpuMat dst = _dst.getGpuMat(); 1663 1664 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1665 } 1666 1667 void LRGB_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1668 { 1669 using namespace cv::cuda::device; 1670 static const gpu_func_t funcs[2][2][2] = 1671 { 1672 { 1673 {LRGB_to_Luv_8u, LRGB_to_Luv_32f}, 1674 {LRGBA_to_Luv_8u, LRGBA_to_Luv_32f} 1675 }, 1676 { 1677 {LRGB_to_Luv4_8u, LRGB_to_Luv4_32f}, 1678 {LRGBA_to_Luv4_8u, LRGBA_to_Luv4_32f} 1679 } 1680 }; 1681 1682 if (dcn <= 0) dcn = 3; 1683 1684 GpuMat src = _src.getGpuMat(); 1685 1686 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1687 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1688 CV_Assert( dcn == 3 || dcn == 4 ); 1689 1690 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1691 GpuMat dst = _dst.getGpuMat(); 1692 1693 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1694 } 1695 1696 void Luv_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1697 { 1698 using namespace cv::cuda::device; 1699 static const gpu_func_t funcs[2][2][2] = 1700 { 1701 { 1702 {Luv_to_BGR_8u, Luv_to_BGR_32f}, 1703 {Luv4_to_BGR_8u, Luv4_to_BGR_32f} 1704 }, 1705 { 1706 {Luv_to_BGRA_8u, Luv_to_BGRA_32f}, 1707 {Luv4_to_BGRA_8u, Luv4_to_BGRA_32f} 1708 } 1709 }; 1710 1711 if (dcn <= 0) dcn = 3; 1712 1713 GpuMat src = _src.getGpuMat(); 1714 1715 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1716 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1717 CV_Assert( dcn == 3 || dcn == 4 ); 1718 1719 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1720 GpuMat dst = _dst.getGpuMat(); 1721 1722 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1723 } 1724 1725 void Luv_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1726 { 1727 using namespace cv::cuda::device; 1728 static const gpu_func_t funcs[2][2][2] = 1729 { 1730 { 1731 {Luv_to_RGB_8u, Luv_to_RGB_32f}, 1732 {Luv4_to_RGB_8u, Luv4_to_RGB_32f} 1733 }, 1734 { 1735 {Luv_to_RGBA_8u, Luv_to_RGBA_32f}, 1736 {Luv4_to_RGBA_8u, Luv4_to_RGBA_32f} 1737 } 1738 }; 1739 1740 if (dcn <= 0) dcn = 3; 1741 1742 GpuMat src = _src.getGpuMat(); 1743 1744 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1745 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1746 CV_Assert( dcn == 3 || dcn == 4 ); 1747 1748 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1749 GpuMat dst = _dst.getGpuMat(); 1750 1751 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1752 } 1753 1754 void Luv_to_LBGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1755 { 1756 using namespace cv::cuda::device; 1757 static const gpu_func_t funcs[2][2][2] = 1758 { 1759 { 1760 {Luv_to_LBGR_8u, Luv_to_LBGR_32f}, 1761 {Luv4_to_LBGR_8u, Luv4_to_LBGR_32f} 1762 }, 1763 { 1764 {Luv_to_LBGRA_8u, Luv_to_LBGRA_32f}, 1765 {Luv4_to_LBGRA_8u, Luv4_to_LBGRA_32f} 1766 } 1767 }; 1768 1769 if (dcn <= 0) dcn = 3; 1770 1771 GpuMat src = _src.getGpuMat(); 1772 1773 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1774 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1775 CV_Assert( dcn == 3 || dcn == 4 ); 1776 1777 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1778 GpuMat dst = _dst.getGpuMat(); 1779 1780 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1781 } 1782 1783 void Luv_to_LRGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream) 1784 { 1785 using namespace cv::cuda::device; 1786 static const gpu_func_t funcs[2][2][2] = 1787 { 1788 { 1789 {Luv_to_LRGB_8u, Luv_to_LRGB_32f}, 1790 {Luv4_to_LRGB_8u, Luv4_to_LRGB_32f} 1791 }, 1792 { 1793 {Luv_to_LRGBA_8u, Luv_to_LRGBA_32f}, 1794 {Luv4_to_LRGBA_8u, Luv4_to_LRGBA_32f} 1795 } 1796 }; 1797 1798 if (dcn <= 0) dcn = 3; 1799 1800 GpuMat src = _src.getGpuMat(); 1801 1802 CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F ); 1803 CV_Assert( src.channels() == 3 || src.channels() == 4 ); 1804 CV_Assert( dcn == 3 || dcn == 4 ); 1805 1806 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1807 GpuMat dst = _dst.getGpuMat(); 1808 1809 funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream); 1810 } 1811 1812 void RGBA_to_mBGRA(InputArray _src, OutputArray _dst, int, Stream& _stream) 1813 { 1814 #if (CUDA_VERSION < 5000) 1815 (void) _src; 1816 (void) _dst; 1817 (void) _stream; 1818 CV_Error( Error::StsBadFlag, "Unknown/unsupported color conversion code" ); 1819 #else 1820 GpuMat src = _src.getGpuMat(); 1821 1822 CV_Assert( src.type() == CV_8UC4 || src.type() == CV_16UC4 ); 1823 1824 _dst.create(src.size(), src.type()); 1825 GpuMat dst = _dst.getGpuMat(); 1826 1827 cudaStream_t stream = StreamAccessor::getStream(_stream); 1828 NppStreamHandler h(stream); 1829 1830 NppiSize oSizeROI; 1831 oSizeROI.width = src.cols; 1832 oSizeROI.height = src.rows; 1833 1834 if (src.depth() == CV_8U) 1835 nppSafeCall( nppiAlphaPremul_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) ); 1836 else 1837 nppSafeCall( nppiAlphaPremul_16u_AC4R(src.ptr<Npp16u>(), static_cast<int>(src.step), dst.ptr<Npp16u>(), static_cast<int>(dst.step), oSizeROI) ); 1838 1839 if (stream == 0) 1840 cudaSafeCall( cudaDeviceSynchronize() ); 1841 #endif 1842 } 1843 1844 void bayer_to_BGR(InputArray _src, OutputArray _dst, int dcn, bool blue_last, bool start_with_green, Stream& stream) 1845 { 1846 typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); 1847 static const func_t funcs[3][4] = 1848 { 1849 {0,0,Bayer2BGR_8u_gpu<3>, Bayer2BGR_8u_gpu<4>}, 1850 {0,0,0,0}, 1851 {0,0,Bayer2BGR_16u_gpu<3>, Bayer2BGR_16u_gpu<4>} 1852 }; 1853 1854 if (dcn <= 0) dcn = 3; 1855 1856 GpuMat src = _src.getGpuMat(); 1857 1858 CV_Assert( src.type() == CV_8UC1 || src.type() == CV_16UC1 ); 1859 CV_Assert( src.rows > 2 && src.cols > 2 ); 1860 CV_Assert( dcn == 3 || dcn == 4 ); 1861 1862 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn)); 1863 GpuMat dst = _dst.getGpuMat(); 1864 1865 funcs[src.depth()][dcn - 1](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream)); 1866 } 1867 void bayerBG_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream) 1868 { 1869 bayer_to_BGR(src, dst, dcn, false, false, stream); 1870 } 1871 void bayeRGB_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream) 1872 { 1873 bayer_to_BGR(src, dst, dcn, false, true, stream); 1874 } 1875 void bayerRG_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream) 1876 { 1877 bayer_to_BGR(src, dst, dcn, true, false, stream); 1878 } 1879 void bayerGR_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream) 1880 { 1881 bayer_to_BGR(src, dst, dcn, true, true, stream); 1882 } 1883 1884 void bayer_to_gray(InputArray _src, OutputArray _dst, bool blue_last, bool start_with_green, Stream& stream) 1885 { 1886 typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); 1887 static const func_t funcs[3] = 1888 { 1889 Bayer2BGR_8u_gpu<1>, 1890 0, 1891 Bayer2BGR_16u_gpu<1>, 1892 }; 1893 1894 GpuMat src = _src.getGpuMat(); 1895 1896 CV_Assert( src.type() == CV_8UC1 || src.type() == CV_16UC1 ); 1897 CV_Assert( src.rows > 2 && src.cols > 2 ); 1898 1899 _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1)); 1900 GpuMat dst = _dst.getGpuMat(); 1901 1902 funcs[src.depth()](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream)); 1903 } 1904 void bayerBG_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream) 1905 { 1906 bayer_to_gray(src, dst, false, false, stream); 1907 } 1908 void bayeRGB_to_GRAY(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream) 1909 { 1910 bayer_to_gray(src, dst, false, true, stream); 1911 } 1912 void bayerRG_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream) 1913 { 1914 bayer_to_gray(src, dst, true, false, stream); 1915 } 1916 void bayerGR_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream) 1917 { 1918 bayer_to_gray(src, dst, true, true, stream); 1919 } 1920 } 1921 1922 //////////////////////////////////////////////////////////////////////// 1923 // cvtColor 1924 1925 void cv::cuda::cvtColor(InputArray src, OutputArray dst, int code, int dcn, Stream& stream) 1926 { 1927 typedef void (*func_t)(InputArray src, OutputArray dst, int dcn, Stream& stream); 1928 static const func_t funcs[] = 1929 { 1930 BGR_to_BGRA, // CV_BGR2BGRA =0 1931 BGRA_to_BGR, // CV_BGRA2BGR =1 1932 BGR_to_RGBA, // CV_BGR2RGBA =2 1933 BGRA_to_RGB, // CV_RGBA2BGR =3 1934 BGR_to_RGB, // CV_BGR2RGB =4 1935 BGRA_to_RGBA, // CV_BGRA2RGBA =5 1936 1937 BGR_to_GRAY, // CV_BGR2GRAY =6 1938 RGB_to_GRAY, // CV_RGB2GRAY =7 1939 GRAY_to_BGR, // CV_GRAY2BGR =8 1940 GRAY_to_BGRA, // CV_GRAY2BGRA =9 1941 BGRA_to_GRAY, // CV_BGRA2GRAY =10 1942 RGBA_to_GRAY, // CV_RGBA2GRAY =11 1943 1944 BGR_to_BGR565, // CV_BGR2BGR565 =12 1945 RGB_to_BGR565, // CV_RGB2BGR565 =13 1946 BGR565_to_BGR, // CV_BGR5652BGR =14 1947 BGR565_to_RGB, // CV_BGR5652RGB =15 1948 BGRA_to_BGR565, // CV_BGRA2BGR565 =16 1949 RGBA_to_BGR565, // CV_RGBA2BGR565 =17 1950 BGR565_to_BGRA, // CV_BGR5652BGRA =18 1951 BGR565_to_RGBA, // CV_BGR5652RGBA =19 1952 1953 GRAY_to_BGR565, // CV_GRAY2BGR565 =20 1954 BGR565_to_GRAY, // CV_BGR5652GRAY =21 1955 1956 BGR_to_BGR555, // CV_BGR2BGR555 =22 1957 RGB_to_BGR555, // CV_RGB2BGR555 =23 1958 BGR555_to_BGR, // CV_BGR5552BGR =24 1959 BGR555_to_RGB, // CV_BGR5552RGB =25 1960 BGRA_to_BGR555, // CV_BGRA2BGR555 =26 1961 RGBA_to_BGR555, // CV_RGBA2BGR555 =27 1962 BGR555_to_BGRA, // CV_BGR5552BGRA =28 1963 BGR555_to_RGBA, // CV_BGR5552RGBA =29 1964 1965 GRAY_to_BGR555, // CV_GRAY2BGR555 =30 1966 BGR555_to_GRAY, // CV_BGR5552GRAY =31 1967 1968 BGR_to_XYZ, // CV_BGR2XYZ =32 1969 RGB_to_XYZ, // CV_RGB2XYZ =33 1970 XYZ_to_BGR, // CV_XYZ2BGR =34 1971 XYZ_to_RGB, // CV_XYZ2RGB =35 1972 1973 BGR_to_YCrCb, // CV_BGR2YCrCb =36 1974 RGB_to_YCrCb, // CV_RGB2YCrCb =37 1975 YCrCb_to_BGR, // CV_YCrCb2BGR =38 1976 YCrCb_to_RGB, // CV_YCrCb2RGB =39 1977 1978 BGR_to_HSV, // CV_BGR2HSV =40 1979 RGB_to_HSV, // CV_RGB2HSV =41 1980 1981 0, // =42 1982 0, // =43 1983 1984 BGR_to_Lab, // CV_BGR2Lab =44 1985 RGB_to_Lab, // CV_RGB2Lab =45 1986 1987 bayerBG_to_BGR, // CV_BayerBG2BGR =46 1988 bayeRGB_to_BGR, // CV_BayeRGB2BGR =47 1989 bayerRG_to_BGR, // CV_BayerRG2BGR =48 1990 bayerGR_to_BGR, // CV_BayerGR2BGR =49 1991 1992 BGR_to_Luv, // CV_BGR2Luv =50 1993 RGB_to_Luv, // CV_RGB2Luv =51 1994 1995 BGR_to_HLS, // CV_BGR2HLS =52 1996 RGB_to_HLS, // CV_RGB2HLS =53 1997 1998 HSV_to_BGR, // CV_HSV2BGR =54 1999 HSV_to_RGB, // CV_HSV2RGB =55 2000 2001 Lab_to_BGR, // CV_Lab2BGR =56 2002 Lab_to_RGB, // CV_Lab2RGB =57 2003 Luv_to_BGR, // CV_Luv2BGR =58 2004 Luv_to_RGB, // CV_Luv2RGB =59 2005 2006 HLS_to_BGR, // CV_HLS2BGR =60 2007 HLS_to_RGB, // CV_HLS2RGB =61 2008 2009 0, // CV_BayerBG2BGR_VNG =62 2010 0, // CV_BayeRGB2BGR_VNG =63 2011 0, // CV_BayerRG2BGR_VNG =64 2012 0, // CV_BayerGR2BGR_VNG =65 2013 2014 BGR_to_HSV_FULL, // CV_BGR2HSV_FULL = 66 2015 RGB_to_HSV_FULL, // CV_RGB2HSV_FULL = 67 2016 BGR_to_HLS_FULL, // CV_BGR2HLS_FULL = 68 2017 RGB_to_HLS_FULL, // CV_RGB2HLS_FULL = 69 2018 2019 HSV_to_BGR_FULL, // CV_HSV2BGR_FULL = 70 2020 HSV_to_RGB_FULL, // CV_HSV2RGB_FULL = 71 2021 HLS_to_BGR_FULL, // CV_HLS2BGR_FULL = 72 2022 HLS_to_RGB_FULL, // CV_HLS2RGB_FULL = 73 2023 2024 LBGR_to_Lab, // CV_LBGR2Lab = 74 2025 LRGB_to_Lab, // CV_LRGB2Lab = 75 2026 LBGR_to_Luv, // CV_LBGR2Luv = 76 2027 LRGB_to_Luv, // CV_LRGB2Luv = 77 2028 2029 Lab_to_LBGR, // CV_Lab2LBGR = 78 2030 Lab_to_LRGB, // CV_Lab2LRGB = 79 2031 Luv_to_LBGR, // CV_Luv2LBGR = 80 2032 Luv_to_LRGB, // CV_Luv2LRGB = 81 2033 2034 BGR_to_YUV, // CV_BGR2YUV = 82 2035 RGB_to_YUV, // CV_RGB2YUV = 83 2036 YUV_to_BGR, // CV_YUV2BGR = 84 2037 YUV_to_RGB, // CV_YUV2RGB = 85 2038 2039 bayerBG_to_gray, // CV_BayerBG2GRAY = 86 2040 bayeRGB_to_GRAY, // CV_BayeRGB2GRAY = 87 2041 bayerRG_to_gray, // CV_BayerRG2GRAY = 88 2042 bayerGR_to_gray, // CV_BayerGR2GRAY = 89 2043 2044 //YUV 4:2:0 formats family 2045 0, // CV_YUV2RGB_NV12 = 90, 2046 0, // CV_YUV2BGR_NV12 = 91, 2047 0, // CV_YUV2RGB_NV21 = 92, 2048 0, // CV_YUV2BGR_NV21 = 93, 2049 2050 0, // CV_YUV2RGBA_NV12 = 94, 2051 0, // CV_YUV2BGRA_NV12 = 95, 2052 0, // CV_YUV2RGBA_NV21 = 96, 2053 0, // CV_YUV2BGRA_NV21 = 97, 2054 2055 0, // CV_YUV2RGB_YV12 = 98, 2056 0, // CV_YUV2BGR_YV12 = 99, 2057 0, // CV_YUV2RGB_IYUV = 100, 2058 0, // CV_YUV2BGR_IYUV = 101, 2059 2060 0, // CV_YUV2RGBA_YV12 = 102, 2061 0, // CV_YUV2BGRA_YV12 = 103, 2062 0, // CV_YUV2RGBA_IYUV = 104, 2063 0, // CV_YUV2BGRA_IYUV = 105, 2064 2065 0, // CV_YUV2GRAY_420 = 106, 2066 2067 //YUV 4:2:2 formats family 2068 0, // CV_YUV2RGB_UYVY = 107, 2069 0, // CV_YUV2BGR_UYVY = 108, 2070 0, // //CV_YUV2RGB_VYUY = 109, 2071 0, // //CV_YUV2BGR_VYUY = 110, 2072 2073 0, // CV_YUV2RGBA_UYVY = 111, 2074 0, // CV_YUV2BGRA_UYVY = 112, 2075 0, // //CV_YUV2RGBA_VYUY = 113, 2076 0, // //CV_YUV2BGRA_VYUY = 114, 2077 2078 0, // CV_YUV2RGB_YUY2 = 115, 2079 0, // CV_YUV2BGR_YUY2 = 116, 2080 0, // CV_YUV2RGB_YVYU = 117, 2081 0, // CV_YUV2BGR_YVYU = 118, 2082 2083 0, // CV_YUV2RGBA_YUY2 = 119, 2084 0, // CV_YUV2BGRA_YUY2 = 120, 2085 0, // CV_YUV2RGBA_YVYU = 121, 2086 0, // CV_YUV2BGRA_YVYU = 122, 2087 2088 0, // CV_YUV2GRAY_UYVY = 123, 2089 0, // CV_YUV2GRAY_YUY2 = 124, 2090 2091 // alpha premultiplication 2092 RGBA_to_mBGRA, // CV_RGBA2mRGBA = 125, 2093 0, // CV_mRGBA2RGBA = 126, 2094 2095 0, // CV_COLORCVT_MAX = 127 2096 }; 2097 2098 CV_Assert( code < 128 ); 2099 2100 func_t func = funcs[code]; 2101 2102 if (func == 0) 2103 CV_Error(Error::StsBadFlag, "Unknown/unsupported color conversion code"); 2104 2105 func(src, dst, dcn, stream); 2106 } 2107 2108 //////////////////////////////////////////////////////////////////////// 2109 // demosaicing 2110 2111 void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn, Stream& stream) 2112 { 2113 CV_Assert( !_src.empty() ); 2114 2115 switch (code) 2116 { 2117 case cv::COLOR_BayerBG2GRAY: case cv::COLOR_BayerGB2GRAY: case cv::COLOR_BayerRG2GRAY: case cv::COLOR_BayerGR2GRAY: 2118 bayer_to_gray(_src, _dst, code == cv::COLOR_BayerBG2GRAY || code == cv::COLOR_BayerGB2GRAY, code == cv::COLOR_BayerGB2GRAY || code == cv::COLOR_BayerGR2GRAY, stream); 2119 break; 2120 2121 case cv::COLOR_BayerBG2BGR: case cv::COLOR_BayerGB2BGR: case cv::COLOR_BayerRG2BGR: case cv::COLOR_BayerGR2BGR: 2122 bayer_to_BGR(_src, _dst, dcn, code == cv::COLOR_BayerBG2BGR || code == cv::COLOR_BayerGB2BGR, code == cv::COLOR_BayerGB2BGR || code == cv::COLOR_BayerGR2BGR, stream); 2123 break; 2124 2125 case COLOR_BayerBG2BGR_MHT: case COLOR_BayerGB2BGR_MHT: case COLOR_BayerRG2BGR_MHT: case COLOR_BayerGR2BGR_MHT: 2126 { 2127 if (dcn <= 0) dcn = 3; 2128 2129 GpuMat src = _src.getGpuMat(); 2130 const int depth = _src.depth(); 2131 2132 CV_Assert( depth == CV_8U ); 2133 CV_Assert( src.channels() == 1 ); 2134 CV_Assert( dcn == 3 || dcn == 4 ); 2135 2136 _dst.create(_src.size(), CV_MAKE_TYPE(depth, dcn)); 2137 GpuMat dst = _dst.getGpuMat(); 2138 2139 dst.setTo(Scalar::all(0), stream); 2140 2141 Size wholeSize; 2142 Point ofs; 2143 src.locateROI(wholeSize, ofs); 2144 PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step); 2145 2146 const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1, 2147 code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1); 2148 2149 if (dcn == 3) 2150 cv::cuda::device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); 2151 else 2152 cv::cuda::device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); 2153 2154 break; 2155 } 2156 2157 case COLOR_BayerBG2GRAY_MHT: case COLOR_BayerGB2GRAY_MHT: case COLOR_BayerRG2GRAY_MHT: case COLOR_BayerGR2GRAY_MHT: 2158 { 2159 GpuMat src = _src.getGpuMat(); 2160 const int depth = _src.depth(); 2161 2162 CV_Assert( depth == CV_8U ); 2163 2164 _dst.create(_src.size(), CV_MAKE_TYPE(depth, 1)); 2165 GpuMat dst = _dst.getGpuMat(); 2166 2167 dst.setTo(Scalar::all(0), stream); 2168 2169 Size wholeSize; 2170 Point ofs; 2171 src.locateROI(wholeSize, ofs); 2172 PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step); 2173 2174 const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1, 2175 code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1); 2176 2177 cv::cuda::device::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); 2178 2179 break; 2180 } 2181 2182 default: 2183 CV_Error(Error::StsBadFlag, "Unknown / unsupported color conversion code"); 2184 } 2185 } 2186 2187 //////////////////////////////////////////////////////////////////////// 2188 // swapChannels 2189 2190 void cv::cuda::swapChannels(InputOutputArray _image, const int dstOrder[4], Stream& _stream) 2191 { 2192 GpuMat image = _image.getGpuMat(); 2193 2194 CV_Assert( image.type() == CV_8UC4 ); 2195 2196 cudaStream_t stream = StreamAccessor::getStream(_stream); 2197 NppStreamHandler h(stream); 2198 2199 NppiSize sz; 2200 sz.width = image.cols; 2201 sz.height = image.rows; 2202 2203 nppSafeCall( nppiSwapChannels_8u_C4IR(image.ptr<Npp8u>(), static_cast<int>(image.step), sz, dstOrder) ); 2204 2205 if (stream == 0) 2206 cudaSafeCall( cudaDeviceSynchronize() ); 2207 } 2208 2209 //////////////////////////////////////////////////////////////////////// 2210 // gammaCorrection 2211 2212 void cv::cuda::gammaCorrection(InputArray _src, OutputArray _dst, bool forward, Stream& stream) 2213 { 2214 #if (CUDA_VERSION < 5000) 2215 (void) _src; 2216 (void) _dst; 2217 (void) forward; 2218 (void) stream; 2219 CV_Error(Error::StsNotImplemented, "This function works only with CUDA 5.0 or higher"); 2220 #else 2221 typedef NppStatus (*func_t)(const Npp8u* pSrc, int nSrcStep, Npp8u* pDst, int nDstStep, NppiSize oSizeROI); 2222 typedef NppStatus (*func_inplace_t)(Npp8u* pSrcDst, int nSrcDstStep, NppiSize oSizeROI); 2223 2224 static const func_t funcs[2][5] = 2225 { 2226 {0, 0, 0, nppiGammaInv_8u_C3R, nppiGammaInv_8u_AC4R}, 2227 {0, 0, 0, nppiGammaFwd_8u_C3R, nppiGammaFwd_8u_AC4R} 2228 }; 2229 static const func_inplace_t funcs_inplace[2][5] = 2230 { 2231 {0, 0, 0, nppiGammaInv_8u_C3IR, nppiGammaInv_8u_AC4IR}, 2232 {0, 0, 0, nppiGammaFwd_8u_C3IR, nppiGammaFwd_8u_AC4IR} 2233 }; 2234 2235 GpuMat src = _src.getGpuMat(); 2236 2237 CV_Assert( src.type() == CV_8UC3 || src.type() == CV_8UC4 ); 2238 2239 _dst.create(src.size(), src.type()); 2240 GpuMat dst = _dst.getGpuMat(); 2241 2242 NppStreamHandler h(StreamAccessor::getStream(stream)); 2243 2244 NppiSize oSizeROI; 2245 oSizeROI.width = src.cols; 2246 oSizeROI.height = src.rows; 2247 2248 if (dst.data == src.data) 2249 funcs_inplace[forward][src.channels()](dst.ptr<Npp8u>(), static_cast<int>(src.step), oSizeROI); 2250 else 2251 funcs[forward][src.channels()](src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI); 2252 2253 #endif 2254 } 2255 2256 //////////////////////////////////////////////////////////////////////// 2257 // alphaComp 2258 2259 namespace 2260 { 2261 template <int DEPTH> struct NppAlphaCompFunc 2262 { 2263 typedef typename NPPTypeTraits<DEPTH>::npp_type npp_t; 2264 2265 typedef NppStatus (*func_t)(const npp_t* pSrc1, int nSrc1Step, const npp_t* pSrc2, int nSrc2Step, npp_t* pDst, int nDstStep, NppiSize oSizeROI, NppiAlphaOp eAlphaOp); 2266 }; 2267 2268 template <int DEPTH, typename NppAlphaCompFunc<DEPTH>::func_t func> struct NppAlphaComp 2269 { 2270 typedef typename NPPTypeTraits<DEPTH>::npp_type npp_t; 2271 2272 static void call(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, NppiAlphaOp eAlphaOp, cudaStream_t stream) 2273 { 2274 NppStreamHandler h(stream); 2275 2276 NppiSize oSizeROI; 2277 oSizeROI.width = img1.cols; 2278 oSizeROI.height = img2.rows; 2279 2280 nppSafeCall( func(img1.ptr<npp_t>(), static_cast<int>(img1.step), img2.ptr<npp_t>(), static_cast<int>(img2.step), 2281 dst.ptr<npp_t>(), static_cast<int>(dst.step), oSizeROI, eAlphaOp) ); 2282 2283 if (stream == 0) 2284 cudaSafeCall( cudaDeviceSynchronize() ); 2285 } 2286 }; 2287 } 2288 2289 void cv::cuda::alphaComp(InputArray _img1, InputArray _img2, OutputArray _dst, int alpha_op, Stream& stream) 2290 { 2291 static const NppiAlphaOp npp_alpha_ops[] = { 2292 NPPI_OP_ALPHA_OVER, 2293 NPPI_OP_ALPHA_IN, 2294 NPPI_OP_ALPHA_OUT, 2295 NPPI_OP_ALPHA_ATOP, 2296 NPPI_OP_ALPHA_XOR, 2297 NPPI_OP_ALPHA_PLUS, 2298 NPPI_OP_ALPHA_OVER_PREMUL, 2299 NPPI_OP_ALPHA_IN_PREMUL, 2300 NPPI_OP_ALPHA_OUT_PREMUL, 2301 NPPI_OP_ALPHA_ATOP_PREMUL, 2302 NPPI_OP_ALPHA_XOR_PREMUL, 2303 NPPI_OP_ALPHA_PLUS_PREMUL, 2304 NPPI_OP_ALPHA_PREMUL 2305 }; 2306 2307 typedef void (*func_t)(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, NppiAlphaOp eAlphaOp, cudaStream_t stream); 2308 static const func_t funcs[] = 2309 { 2310 NppAlphaComp<CV_8U, nppiAlphaComp_8u_AC4R>::call, 2311 0, 2312 NppAlphaComp<CV_16U, nppiAlphaComp_16u_AC4R>::call, 2313 0, 2314 NppAlphaComp<CV_32S, nppiAlphaComp_32s_AC4R>::call, 2315 NppAlphaComp<CV_32F, nppiAlphaComp_32f_AC4R>::call 2316 }; 2317 2318 GpuMat img1 = _img1.getGpuMat(); 2319 GpuMat img2 = _img2.getGpuMat(); 2320 2321 CV_Assert( img1.type() == CV_8UC4 || img1.type() == CV_16UC4 || img1.type() == CV_32SC4 || img1.type() == CV_32FC4 ); 2322 CV_Assert( img1.size() == img2.size() && img1.type() == img2.type() ); 2323 2324 _dst.create(img1.size(), img1.type()); 2325 GpuMat dst = _dst.getGpuMat(); 2326 2327 const func_t func = funcs[img1.depth()]; 2328 2329 func(img1, img2, dst, npp_alpha_ops[alpha_op], StreamAccessor::getStream(stream)); 2330 } 2331 2332 #endif /* !defined (HAVE_CUDA) */ 2333
