1793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler/*M/////////////////////////////////////////////////////////////////////////////////////// 2793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 3793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 5793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// By downloading, copying, installing or using the software you agree to this license. 6793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// If you do not agree to this license, do not download, install, 7793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// copy or use the software. 8793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 9793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 10793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// License Agreement 11793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// For Open Source Computer Vision Library 12793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 13793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// Third party copyrights are property of their respective owners. 16793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 17793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// Redistribution and use in source and binary forms, with or without modification, 18793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// are permitted provided that the following conditions are met: 19793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 20793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// * Redistribution's of source code must retain the above copyright notice, 21793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// this list of conditions and the following disclaimer. 22793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 23793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// * Redistribution's in binary form must reproduce the above copyright notice, 24793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// this list of conditions and the following disclaimer in the documentation 25793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// and/or other materials provided with the distribution. 26793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 27793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// * The name of the copyright holders may not be used to endorse or promote products 28793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// derived from this software without specific prior written permission. 29793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 30793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// This software is provided by the copyright holders and contributors "as is" and 31793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// any express or implied warranties, including, but not limited to, the implied 32793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// warranties of merchantability and fitness for a particular purpose are disclaimed. 33793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// In no event shall the Intel Corporation or contributors be liable for any direct, 34793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// indirect, incidental, special, exemplary, or consequential damages 35793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// (including, but not limited to, procurement of substitute goods or services; 36793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// loss of use, data, or profits; or business interruption) however caused 37793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// and on any theory of liability, whether in contract, strict liability, 38793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// or tort (including negligence or otherwise) arising in any way out of 39793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// the use of this software, even if advised of the possibility of such damage. 40793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 41793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler//M*/ 42793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 43793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler#include <iostream> 44793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler#include <vector> 45793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 46793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler#include "opencv2/cudalegacy/NCV.hpp" 47793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 48793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler//=================================================================== 49793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 50793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// Operations with rectangles 51793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler// 52793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler//=================================================================== 53793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 54793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 55793ee12c6df9cad3806238d32528c49a3ff9331dNoah Preslerconst Ncv32u NUMTHREADS_DRAWRECTS = 32; 56793ee12c6df9cad3806238d32528c49a3ff9331dNoah Preslerconst Ncv32u NUMTHREADS_DRAWRECTS_LOG2 = 5; 57793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 58793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 59793ee12c6df9cad3806238d32528c49a3ff9331dNoah Preslertemplate <class T> 60793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler__global__ void drawRects(T *d_dst, 61793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstStride, 62793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstWidth, 63793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstHeight, 64793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvRect32u *d_rects, 65793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numRects, 66793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler T color) 67793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler{ 68793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u blockId = blockIdx.y * 65535 + blockIdx.x; 69793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (blockId > numRects * 4) 70793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 71793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler return; 72793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 73793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 74793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvRect32u curRect = d_rects[blockId >> 2]; 75793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvBool bVertical = blockId & 0x1; 76793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvBool bTopLeft = blockId & 0x2; 77793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 78793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u pt0x, pt0y; 79793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (bVertical) 80793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 81793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numChunks = (curRect.height + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2; 82793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 83793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler pt0x = bTopLeft ? curRect.x : curRect.x + curRect.width - 1; 84793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler pt0y = curRect.y; 85793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 86793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (pt0x < dstWidth) 87793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 88793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++) 89793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 90793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u ptY = pt0y + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x; 91793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (ptY < pt0y + curRect.height && ptY < dstHeight) 92793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 93793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler d_dst[ptY * dstStride + pt0x] = color; 94793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 95793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 96793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 97793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 98793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler else 99793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 100793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numChunks = (curRect.width + NUMTHREADS_DRAWRECTS - 1) >> NUMTHREADS_DRAWRECTS_LOG2; 101793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 102793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler pt0x = curRect.x; 103793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler pt0y = bTopLeft ? curRect.y : curRect.y + curRect.height - 1; 104793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 105793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (pt0y < dstHeight) 106793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 107793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler for (Ncv32u chunkId = 0; chunkId < numChunks; chunkId++) 108793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 109793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u ptX = pt0x + chunkId * NUMTHREADS_DRAWRECTS + threadIdx.x; 110793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (ptX < pt0x + curRect.width && ptX < dstWidth) 111793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 112793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler d_dst[pt0y * dstStride + ptX] = color; 113793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 114793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 115793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 116793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 117793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler} 118793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 119793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 120793ee12c6df9cad3806238d32528c49a3ff9331dNoah Preslertemplate <class T> 121793ee12c6df9cad3806238d32528c49a3ff9331dNoah Preslerstatic NCVStatus drawRectsWrapperDevice(T *d_dst, 122793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstStride, 123793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstWidth, 124793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstHeight, 125793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvRect32u *d_rects, 126793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numRects, 127793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler T color, 128793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler cudaStream_t cuStream) 129793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler{ 130793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler (void)cuStream; 131793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler ncvAssertReturn(d_dst != NULL && d_rects != NULL, NCV_NULL_PTR); 132793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler ncvAssertReturn(dstWidth > 0 && dstHeight > 0, NCV_DIMENSIONS_INVALID); 133793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler ncvAssertReturn(dstStride >= dstWidth, NCV_INVALID_STEP); 134793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler ncvAssertReturn(numRects <= dstWidth * dstHeight, NCV_DIMENSIONS_INVALID); 135793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 136793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (numRects == 0) 137793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 138793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler return NCV_SUCCESS; 139793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 140793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 141793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler dim3 grid(numRects * 4); 142793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler dim3 block(NUMTHREADS_DRAWRECTS); 143793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler if (grid.x > 65535) 144793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler { 145793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler grid.y = (grid.x + 65534) / 65535; 146793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler grid.x = 65535; 147793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler } 148793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 149793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler drawRects<T><<<grid, block>>>(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color); 150793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 151793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler ncvAssertCUDALastErrorReturn(NCV_CUDA_ERROR); 152793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 153793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler return NCV_SUCCESS; 154793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler} 155793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 156793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 157793ee12c6df9cad3806238d32528c49a3ff9331dNoah PreslerNCVStatus ncvDrawRects_8u_device(Ncv8u *d_dst, 158793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstStride, 159793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstWidth, 160793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstHeight, 161793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvRect32u *d_rects, 162793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numRects, 163793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv8u color, 164793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler cudaStream_t cuStream) 165793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler{ 166793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream); 167793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler} 168793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 169793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler 170793ee12c6df9cad3806238d32528c49a3ff9331dNoah PreslerNCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, 171793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstStride, 172793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstWidth, 173793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u dstHeight, 174793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler NcvRect32u *d_rects, 175793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u numRects, 176793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler Ncv32u color, 177793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler cudaStream_t cuStream) 178793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler{ 179793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler return drawRectsWrapperDevice(d_dst, dstStride, dstWidth, dstHeight, d_rects, numRects, color, cuStream); 180793ee12c6df9cad3806238d32528c49a3ff9331dNoah Presler} 181