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 <iostream> 44 #include <vector> 45 46 #include "opencv2/cudalegacy/NCV.hpp" 47 48 //=================================================================== 49 // 50 // Operations with rectangles 51 // 52 //=================================================================== 53 54 55 const Ncv32u NUMTHREADS_DRAWRECTS = 32; 56 const Ncv32u NUMTHREADS_DRAWRECTS_LOG2 = 5; 57 58 59 template <class T> 60 __global__ void drawRects(T *d_dst, 61 Ncv32u dstStride, 62 Ncv32u dstWidth, 63 Ncv32u dstHeight, 64 NcvRect32u *d_rects, 65 Ncv32u numRects, 66 T color) 67 { 68 Ncv32u blockId = blockIdx.y * 65535 + blockIdx.x; 69 if (blockId > numRects * 4) 70 { 71 return; 72 } 73 74 NcvRect32u curRect = d_rects[blockId >> 2]; 75 NcvBool bVertical = blockId & 0x1; 76 NcvBool bTopLeft = blockId & 0x2; 77 78 Ncv32u pt0x, pt0y; 79 if (bVertical) 80 { 81 Ncv32u numChunks = (curRect.height + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2; 82 83 pt0x = bTopLeft ? curRect.x : curRect.x + curRect.width - 1; 84 pt0y = curRect.y; 85 86 if (pt0x < dstWidth) 87 { 88 for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++) 89 { 90 Ncv32u ptY = pt0y + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x; 91 if (ptY < pt0y + curRect.height && ptY < dstHeight) 92 { 93 d_dst[ptY * dstStride + pt0x] = color; 94 } 95 } 96 } 97 } 98 else 99 { 100 Ncv32u numChunks = (curRect.width + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2; 101 102 pt0x = curRect.x; 103 pt0y = bTopLeft ? curRect.y : curRect.y + curRect.height - 1; 104 105 if (pt0y < dstHeight) 106 { 107 for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++) 108 { 109 Ncv32u ptX = pt0x + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x; 110 if (ptX < pt0x + curRect.width && ptX < dstWidth) 111 { 112 d_dst[pt0y * dstStride + ptX] = color; 113 } 114 } 115 } 116 } 117 } 118 119 120 template <class T> 121 static NCVStatus drawRectsWrapperDevice(T *d_dst, 122 Ncv32u dstStride, 123 Ncv32u dstWidth, 124 Ncv32u dstHeight, 125 NcvRect32u *d_rects, 126 Ncv32u numRects, 127 T color, 128 cudaStream_t cuStream) 129 { 130 (void)cuStream; 131 ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR); 132 ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID); 133 ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP); 134 ncvAssertReturn(numRects <= dstWidth * dstHeight, NCV_DIMENSIONS_INVALID); 135 136 if (numRects == 0) 137 { 138 return NCV_SUCCESS; 139 } 140 141 dim3 grid(numRects * 4); 142 dim3 block(NUMTHREADS_DRAWRECTS); 143 if (grid.x > 65535) 144 { 145 grid.y = (grid.x + 65534) / 65535; 146 grid.x = 65535; 147 } 148 149 drawRects<T><<<grid, block>>>(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color); 150 151 ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR); 152 153 return NCV_SUCCESS; 154 } 155 156 157 NCVStatus ncvDrawRects_8u_device(Ncv8u *d_dst, 158 Ncv32u dstStride, 159 Ncv32u dstWidth, 160 Ncv32u dstHeight, 161 NcvRect32u *d_rects, 162 Ncv32u numRects, 163 Ncv8u color, 164 cudaStream_t cuStream) 165 { 166 return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream); 167 } 168 169 170 NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, 171 Ncv32u dstStride, 172 Ncv32u dstWidth, 173 Ncv32u dstHeight, 174 NcvRect32u *d_rects, 175 Ncv32u numRects, 176 Ncv32u color, 177 cudaStream_t cuStream) 178 { 179 return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream); 180 } 181
