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 // Copyright (C) 2013, OpenCV Foundation, all rights reserved. 16 // Third party copyrights are property of their respective owners. 17 // 18 // Redistribution and use in source and binary forms, with or without modification, 19 // are permitted provided that the following conditions are met: 20 // 21 // * Redistribution's of source code must retain the above copyright notice, 22 // this list of conditions and the following disclaimer. 23 // 24 // * Redistribution's in binary form must reproduce the above copyright notice, 25 // this list of conditions and the following disclaimer in the documentation 26 // and/or other materials provided with the distribution. 27 // 28 // * The name of the copyright holders may not be used to endorse or promote products 29 // derived from this software without specific prior written permission. 30 // 31 // This software is provided by the copyright holders and contributors "as is" and 32 // any express or implied warranties, including, but not limited to, the implied 33 // warranties of merchantability and fitness for a particular purpose are disclaimed. 34 // In no event shall the Intel Corporation or contributors be liable for any direct, 35 // indirect, incidental, special, exemplary, or consequential damages 36 // (including, but not limited to, procurement of substitute goods or services; 37 // loss of use, data, or profits; or business interruption) however caused 38 // and on any theory of liability, whether in contract, strict liability, 39 // or tort (including negligence or otherwise) arising in any way out of 40 // the use of this software, even if advised of the possibility of such damage. 41 // 42 //M*/ 43 44 #pragma once 45 46 #ifndef __OPENCV_CUDEV_GRID_COPY_DETAIL_HPP__ 47 #define __OPENCV_CUDEV_GRID_COPY_DETAIL_HPP__ 48 49 #include "../../common.hpp" 50 #include "../../util/tuple.hpp" 51 #include "../../util/saturate_cast.hpp" 52 #include "../../ptr2d/glob.hpp" 53 #include "../../ptr2d/traits.hpp" 54 55 namespace cv { namespace cudev { 56 57 namespace grid_copy_detail 58 { 59 template <class SrcPtr, typename DstType, class MaskPtr> 60 __global__ void copy(const SrcPtr src, GlobPtr<DstType> dst, const MaskPtr mask, const int rows, const int cols) 61 { 62 const int x = blockIdx.x * blockDim.x + threadIdx.x; 63 const int y = blockIdx.y * blockDim.y + threadIdx.y; 64 65 if (x >= cols || y >= rows || !mask(y, x)) 66 return; 67 68 dst(y, x) = saturate_cast<DstType>(src(y, x)); 69 } 70 71 template <class Policy, class SrcPtr, typename DstType, class MaskPtr> 72 __host__ void copy(const SrcPtr& src, const GlobPtr<DstType>& dst, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) 73 { 74 const dim3 block(Policy::block_size_x, Policy::block_size_y); 75 const dim3 grid(divUp(cols, block.x), divUp(rows, block.y)); 76 77 copy<<<grid, block, 0, stream>>>(src, dst, mask, rows, cols); 78 CV_CUDEV_SAFE_CALL( cudaGetLastError() ); 79 80 if (stream == 0) 81 CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); 82 } 83 84 template <int count> struct Unroll 85 { 86 template <class SrcPtrTuple, class DstPtrTuple> 87 __device__ static void copy(const SrcPtrTuple& src, DstPtrTuple& dst, const int y, const int x) 88 { 89 typedef typename tuple_element<count - 1, DstPtrTuple>::type dst_ptr_type; 90 typedef typename PtrTraits<dst_ptr_type>::value_type dst_type; 91 92 get<count - 1>(dst)(y, x) = saturate_cast<dst_type>(get<count - 1>(src)(y, x)); 93 Unroll<count - 1>::copy(src, dst, y, x); 94 } 95 }; 96 template <> struct Unroll<0> 97 { 98 template <class SrcPtrTuple, class DstPtrTuple> 99 __device__ __forceinline__ static void copy(const SrcPtrTuple&, DstPtrTuple&, const int, const int) 100 { 101 } 102 }; 103 104 template <class SrcPtrTuple, class DstPtrTuple, class MaskPtr> 105 __global__ void copy_tuple(const SrcPtrTuple src, DstPtrTuple dst, const MaskPtr mask, const int rows, const int cols) 106 { 107 const int x = blockIdx.x * blockDim.x + threadIdx.x; 108 const int y = blockIdx.y * blockDim.y + threadIdx.y; 109 110 if (x >= cols || y >= rows || !mask(y, x)) 111 return; 112 113 Unroll<tuple_size<SrcPtrTuple>::value>::copy(src, dst, y, x); 114 } 115 116 template <class Policy, class SrcPtrTuple, class DstPtrTuple, class MaskPtr> 117 __host__ void copy_tuple(const SrcPtrTuple& src, const DstPtrTuple& dst, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) 118 { 119 const dim3 block(Policy::block_size_x, Policy::block_size_y); 120 const dim3 grid(divUp(cols, block.x), divUp(rows, block.y)); 121 122 copy_tuple<<<grid, block, 0, stream>>>(src, dst, mask, rows, cols); 123 CV_CUDEV_SAFE_CALL( cudaGetLastError() ); 124 125 if (stream == 0) 126 CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); 127 } 128 } 129 130 }} 131 132 #endif 133
