Point Cloud Library (PCL)  1.9.1-dev
NCV.hpp
1 /*
2  * Software License Agreement (BSD License)
3  *
4  * Point Cloud Library (PCL) - www.pointclouds.org
5  * Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.
6  * Third party copyrights are property of their respective owners.
7  *
8  * All rights reserved.
9  *
10  * Redistribution and use in source and binary forms, with or without
11  * modification, are permitted provided that the following conditions
12  * are met:
13  *
14  * * Redistributions of source code must retain the above copyright
15  * notice, this list of conditions and the following disclaimer.
16  * * Redistributions in binary form must reproduce the above
17  * copyright notice, this list of conditions and the following
18  * disclaimer in the documentation and/or other materials provided
19  * with the distribution.
20  * * Neither the name of Willow Garage, Inc. nor the names of its
21  * contributors may be used to endorse or promote products derived
22  * from this software without specific prior written permission.
23  *
24  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
25  * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
26  * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
27  * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
28  * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
29  * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
30  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
31  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
32  * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
33  * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
34  * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
35  * POSSIBILITY OF SUCH DAMAGE.
36  *
37  * $Id: $
38  * Ported to PCL by Koen Buys : Attention Work in progress!
39  */
40 
41 #ifndef PCL_GPU_PEOPLE__NCV_HPP_
42 #define PCL_GPU_PEOPLE__NCV_HPP_
43 
44 #if (defined WIN32 || defined _WIN32 || defined WINCE) && defined CVAPI_EXPORTS
45  #define NCV_EXPORTS __declspec(dllexport)
46 #else
47  #define NCV_EXPORTS
48 #endif
49 
50 #ifdef _WIN32
51  #define WIN32_LEAN_AND_MEAN
52 #endif
53 
54 #include <cuda_runtime.h>
55 #include <sstream>
56 #include <iostream>
57 #include <pcl/console/print.h>
58 
59 //==============================================================================
60 //
61 // Compile-time assert functionality
62 //
63 //==============================================================================
64 
65 /**
66 * Compile-time assert namespace
67 */
68 namespace NcvCTprep
69 {
70  template <bool x>
72 
73  template <>
74  struct CT_ASSERT_FAILURE<true> {};
75 
76  template <int x>
77  struct assertTest{};
78 }
79 
80 #define NCV_CT_PREP_PASTE_AUX(a,b) a##b ///< Concatenation indirection macro
81 #define NCV_CT_PREP_PASTE(a,b) NCV_CT_PREP_PASTE_AUX(a, b) ///< Concatenation macro
82 
83 /**
84 * Performs compile-time assertion of a condition on the file scope
85 */
86 #define NCV_CT_ASSERT(X) \
87  typedef NcvCTprep::assertTest<sizeof(NcvCTprep::CT_ASSERT_FAILURE< (bool)(X) >)> \
88  NCV_CT_PREP_PASTE(__ct_assert_typedef_, __LINE__)
89 
90 //==============================================================================
91 //
92 // Alignment macros
93 //
94 //==============================================================================
95 
96 #if !defined(__align__) && !defined(__CUDACC__)
97  #if defined(_WIN32) || defined(_WIN64)
98  #define __align__(n) __declspec(align(n))
99  #elif defined(__unix__)
100  #define __align__(n) __attribute__((__aligned__(n)))
101  #endif
102 #endif
103 
104 //==============================================================================
105 //
106 // Integral and compound types of guaranteed size
107 //
108 //==============================================================================
109 
110 using NcvBool = bool;
111 using Ncv64s = long long;
112 
113 #if defined(__APPLE__) && !defined(__CUDACC__)
114  using Ncv64u = std::uint64_t;
115 #else
116  using Ncv64u = unsigned long long;
117 #endif
118 
119 using Ncv32s = int;
120 using Ncv32u = unsigned int;
121 using Ncv16s = short;
122 using Ncv16u = unsigned short;
123 using Ncv8s = char;
124 using Ncv8u = unsigned char;
125 using Ncv32f = float;
126 using Ncv64f = double;
127 
128 struct NcvRect8u
129 {
130  Ncv8u x;
131  Ncv8u y;
132  Ncv8u width;
133  Ncv8u height;
134  __host__ __device__ NcvRect8u() : x(0), y(0), width(0), height(0) {};
135  __host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height) : x(x), y(y), width(width), height(height) {}
136 };
137 
139 {
140  Ncv32s x; ///< x-coordinate of upper left corner.
141  Ncv32s y; ///< y-coordinate of upper left corner.
142  Ncv32s width; ///< Rectangle width.
143  Ncv32s height; ///< Rectangle height.
144  __host__ __device__ NcvRect32s() : x(0), y(0), width(0), height(0) {};
145  __host__ __device__ NcvRect32s(Ncv32s x, Ncv32s y, Ncv32s width, Ncv32s height) : x(x), y(y), width(width), height(height) {}
146 };
147 
149 {
150  Ncv32u x; ///< x-coordinate of upper left corner.
151  Ncv32u y; ///< y-coordinate of upper left corner.
152  Ncv32u width; ///< Rectangle width.
153  Ncv32u height; ///< Rectangle height.
154  __host__ __device__ NcvRect32u() : x(0), y(0), width(0), height(0) {};
155  __host__ __device__ NcvRect32u(Ncv32u x, Ncv32u y, Ncv32u width, Ncv32u height) : x(x), y(y), width(width), height(height) {}
156 };
157 
159 {
160  Ncv32s width; ///< Rectangle width.
161  Ncv32s height; ///< Rectangle height.
162  __host__ __device__ NcvSize32s() : width(0), height(0) {};
163  __host__ __device__ NcvSize32s(Ncv32s width, Ncv32s height) : width(width), height(height) {}
164 };
165 
167 {
168  Ncv32u width; ///< Rectangle width.
169  Ncv32u height; ///< Rectangle height.
170  __host__ __device__ NcvSize32u() : width(0), height(0) {};
171  __host__ __device__ NcvSize32u(Ncv32u width, Ncv32u height) : width(width), height(height) {}
172  __host__ __device__ bool operator == (const NcvSize32u &another) const {return this->width == another.width && this->height == another.height;}
173 };
174 
176 {
177  Ncv32s x; ///< Point X.
178  Ncv32s y; ///< Point Y.
179  __host__ __device__ NcvPoint2D32s() : x(0), y(0) {};
180  __host__ __device__ NcvPoint2D32s(Ncv32s x, Ncv32s y) : x(x), y(y) {}
181 };
182 
184 {
185  Ncv32u x; ///< Point X.
186  Ncv32u y; ///< Point Y.
187  __host__ __device__ NcvPoint2D32u() : x(0), y(0) {};
188  __host__ __device__ NcvPoint2D32u(Ncv32u x, Ncv32u y) : x(x), y(y) {}
189 };
190 
191 NCV_CT_ASSERT(sizeof(NcvBool) <= 4);
192 NCV_CT_ASSERT(sizeof(Ncv64s) == 8);
193 NCV_CT_ASSERT(sizeof(Ncv64u) == 8);
194 NCV_CT_ASSERT(sizeof(Ncv32s) == 4);
195 NCV_CT_ASSERT(sizeof(Ncv32u) == 4);
196 NCV_CT_ASSERT(sizeof(Ncv16s) == 2);
197 NCV_CT_ASSERT(sizeof(Ncv16u) == 2);
198 NCV_CT_ASSERT(sizeof(Ncv8s) == 1);
199 NCV_CT_ASSERT(sizeof(Ncv8u) == 1);
200 NCV_CT_ASSERT(sizeof(Ncv32f) == 4);
201 NCV_CT_ASSERT(sizeof(Ncv64f) == 8);
202 NCV_CT_ASSERT(sizeof(NcvRect8u) == sizeof(Ncv32u));
203 NCV_CT_ASSERT(sizeof(NcvRect32s) == 4 * sizeof(Ncv32s));
204 NCV_CT_ASSERT(sizeof(NcvRect32u) == 4 * sizeof(Ncv32u));
205 NCV_CT_ASSERT(sizeof(NcvSize32u) == 2 * sizeof(Ncv32u));
206 NCV_CT_ASSERT(sizeof(NcvPoint2D32u) == 2 * sizeof(Ncv32u));
207 
208 
209 //==============================================================================
210 //
211 // Persistent constants
212 //
213 //==============================================================================
214 
215 const Ncv32u K_WARP_SIZE = 32;
216 const Ncv32u K_LOG2_WARP_SIZE = 5;
217 
218 //==============================================================================
219 //
220 // Error handling
221 //
222 //==============================================================================
223 
224 NCV_EXPORTS void ncvDebugOutput(const std::string &msg);
225 
226 using NCVDebugOutputHandler = void (const std::string &);
227 
228 NCV_EXPORTS void ncvSetDebugOutputHandler(NCVDebugOutputHandler* func);
229 
230 #define ncvAssertPrintCheck(pred, msg) \
231  do \
232  { \
233  if (!(pred)) \
234  { \
235  std::ostringstream oss; \
236  oss << "NCV Assertion Failed: " << msg << ", file=" << __FILE__ << ", line=" << __LINE__ << std::endl; \
237  ncvDebugOutput(oss.str()); \
238  } \
239  } while (0)
240 
241 #define ncvAssertPrintReturn(pred, msg, err) \
242  do \
243  { \
244  ncvAssertPrintCheck(pred, msg); \
245  if (!(pred)) return err; \
246  } while (0)
247 
248 #define ncvAssertReturn(pred, err) \
249  ncvAssertPrintReturn(pred, "retcode=" << (int)err, err)
250 
251 #define ncvAssertReturnNcvStat(ncvOp) \
252  do \
253  { \
254  NCVStatus _ncvStat = ncvOp; \
255  ncvAssertPrintReturn(NCV_SUCCESS==_ncvStat, "NcvStat=" << (int)_ncvStat, _ncvStat); \
256  } while (0)
257 
258 #define ncvAssertCUDAReturn(cudacall, errCode) \
259  do \
260  { \
261  cudaError_t res = cudacall; \
262  ncvAssertPrintReturn(cudaSuccess==res, "cudaError_t=" << res, errCode); \
263  } while (0)
264 
265 #define ncvAssertCUDALastErrorReturn(errCode) \
266  do \
267  { \
268  cudaError_t res = cudaGetLastError(); \
269  ncvAssertPrintReturn(cudaSuccess==res, "cudaError_t=" << res, errCode); \
270  } while (0)
271 
272 /**
273 * \brief Return-codes for status notification, errors and warnings
274 */
275 enum
276 {
277  //NCV statuses
278  NCV_SUCCESS,
279  NCV_UNKNOWN_ERROR,
280 
281  NCV_CUDA_ERROR,
282  NCV_NPP_ERROR,
283  NCV_FILE_ERROR,
284 
285  NCV_NULL_PTR,
286  NCV_INCONSISTENT_INPUT,
287  NCV_TEXTURE_BIND_ERROR,
288  NCV_DIMENSIONS_INVALID,
289 
290  NCV_INVALID_ROI,
291  NCV_INVALID_STEP,
292  NCV_INVALID_SCALE,
293 
294  NCV_ALLOCATOR_NOT_INITIALIZED,
295  NCV_ALLOCATOR_BAD_ALLOC,
296  NCV_ALLOCATOR_BAD_DEALLOC,
297  NCV_ALLOCATOR_INSUFFICIENT_CAPACITY,
298  NCV_ALLOCATOR_DEALLOC_ORDER,
299  NCV_ALLOCATOR_BAD_REUSE,
300 
301  NCV_MEM_COPY_ERROR,
302  NCV_MEM_RESIDENCE_ERROR,
303  NCV_MEM_INSUFFICIENT_CAPACITY,
304 
305  NCV_HAAR_INVALID_PIXEL_STEP,
306  NCV_HAAR_TOO_MANY_FEATURES_IN_CLASSIFIER,
307  NCV_HAAR_TOO_MANY_FEATURES_IN_CASCADE,
308  NCV_HAAR_TOO_LARGE_FEATURES,
309  NCV_HAAR_XML_LOADING_EXCEPTION,
310 
311  NCV_NOIMPL_HAAR_TILTED_FEATURES,
312  NCV_NOT_IMPLEMENTED,
313 
314  NCV_WARNING_HAAR_DETECTIONS_VECTOR_OVERFLOW,
315 
316  //NPP statuses
317  NPPST_SUCCESS = NCV_SUCCESS, ///< Successful operation (same as NPP_NO_ERROR)
318  NPPST_ERROR, ///< Unknown error
319  NPPST_CUDA_KERNEL_EXECUTION_ERROR, ///< CUDA kernel execution error
320  NPPST_NULL_POINTER_ERROR, ///< NULL pointer argument error
321  NPPST_TEXTURE_BIND_ERROR, ///< CUDA texture binding error or non-zero offset returned
322  NPPST_MEMCPY_ERROR, ///< CUDA memory copy error
323  NPPST_MEM_ALLOC_ERR, ///< CUDA memory allocation error
324  NPPST_MEMFREE_ERR, ///< CUDA memory deallocation error
325 
326  //NPPST statuses
327  NPPST_INVALID_ROI, ///< Invalid region of interest argument
328  NPPST_INVALID_STEP, ///< Invalid image lines step argument (check sign, alignment, relation to image width)
329  NPPST_INVALID_SCALE, ///< Invalid scale parameter passed
330  NPPST_MEM_INSUFFICIENT_BUFFER, ///< Insufficient user-allocated buffer
331  NPPST_MEM_RESIDENCE_ERROR, ///< Memory residence error detected (check if pointers should be device or pinned)
332  NPPST_MEM_INTERNAL_ERROR, ///< Internal memory management error
333 
334  NCV_LAST_STATUS ///< Marker to continue error numeration in other files
335 };
336 
337 using NCVStatus = Ncv32u;
338 
339 #define NCV_SET_SKIP_COND(x) \
340  bool __ncv_skip_cond = x
341 
342 #define NCV_RESET_SKIP_COND(x) \
343  __ncv_skip_cond = x
344 
345 #define NCV_SKIP_COND_BEGIN \
346  if (!__ncv_skip_cond) {
347 
348 #define NCV_SKIP_COND_END \
349  }
350 
351 
352 //==============================================================================
353 //
354 // Memory management classes template compound types
355 //
356 //==============================================================================
357 
358 
359 /**
360 * Calculates the aligned top bound value
361 */
362 NCV_EXPORTS Ncv32u alignUp(Ncv32u what, Ncv32u alignment);
363 
364 
365 /**
366 * NCVMemoryType
367 */
368 enum NCVMemoryType
369 {
370  NCVMemoryTypeNone,
371  NCVMemoryTypeHostPageable,
372  NCVMemoryTypeHostPinned,
373  NCVMemoryTypeDevice
374 };
375 
376 
377 /**
378 * NCVMemPtr
379 */
380 struct NCV_EXPORTS NCVMemPtr
381 {
382  void *ptr;
383  NCVMemoryType memtype;
384  void clear();
385 };
386 
387 
388 /**
389 * NCVMemSegment
390 */
391 struct NCV_EXPORTS NCVMemSegment
392 {
394  std::size_t size;
395  void clear();
396 };
397 
398 
399 /**
400 * INCVMemAllocator (Interface)
401 */
402 class NCV_EXPORTS INCVMemAllocator
403 {
404 public:
405  virtual ~INCVMemAllocator() = 0;
406 
407  virtual NCVStatus alloc(NCVMemSegment &seg, std::size_t size) = 0;
408  virtual NCVStatus dealloc(NCVMemSegment &seg) = 0;
409 
410  virtual NcvBool isInitialized() const = 0;
411  virtual NcvBool isCounting() const = 0;
412 
413  virtual NCVMemoryType memType() const = 0;
414  virtual Ncv32u alignment() const = 0;
415  virtual std::size_t maxSize() const = 0;
416 };
417 
419 
420 
421 /**
422 * NCVMemStackAllocator
423 */
424 class NCV_EXPORTS NCVMemStackAllocator : public INCVMemAllocator
425 {
427  NCVMemStackAllocator(const NCVMemStackAllocator &);
428 
429 public:
430 
431  explicit NCVMemStackAllocator(Ncv32u alignment);
432  NCVMemStackAllocator(NCVMemoryType memT, std::size_t capacity, Ncv32u alignment, void *reusePtr=nullptr);
433  virtual ~NCVMemStackAllocator();
434 
435  virtual NCVStatus alloc(NCVMemSegment &seg, std::size_t size);
436  virtual NCVStatus dealloc(NCVMemSegment &seg);
437 
438  virtual NcvBool isInitialized() const;
439  virtual NcvBool isCounting() const;
440 
441  virtual NCVMemoryType memType() const;
442  virtual Ncv32u alignment() const;
443  virtual std::size_t maxSize() const;
444 
445 private:
446 
447  NCVMemoryType _memType;
448  Ncv32u _alignment;
449  Ncv8u *allocBegin;
450  Ncv8u *begin;
451  Ncv8u *end;
452  std::size_t currentSize;
453  std::size_t _maxSize;
454  NcvBool bReusesMemory;
455 };
456 
457 
458 /**
459 * NCVMemNativeAllocator
460 */
461 class NCV_EXPORTS NCVMemNativeAllocator : public INCVMemAllocator
462 {
463 public:
464 
465  NCVMemNativeAllocator(NCVMemoryType memT, Ncv32u alignment);
466  virtual ~NCVMemNativeAllocator();
467 
468  virtual NCVStatus alloc(NCVMemSegment &seg, std::size_t size);
469  virtual NCVStatus dealloc(NCVMemSegment &seg);
470 
471  virtual NcvBool isInitialized() const;
472  virtual NcvBool isCounting() const;
473 
474  virtual NCVMemoryType memType() const;
475  virtual Ncv32u alignment() const;
476  virtual std::size_t maxSize() const;
477 
478 private:
479 
482 
483  NCVMemoryType _memType;
484  Ncv32u _alignment;
485  std::size_t currentSize;
486  std::size_t _maxSize;
487 };
488 
489 
490 /**
491 * Copy dispatchers
492 */
493 NCV_EXPORTS NCVStatus memSegCopyHelper(void *dst, NCVMemoryType dstType,
494  const void *src, NCVMemoryType srcType,
495  std::size_t sz, cudaStream_t cuStream);
496 
497 
498 NCV_EXPORTS NCVStatus memSegCopyHelper2D(void *dst, Ncv32u dstPitch, NCVMemoryType dstType,
499  const void *src, Ncv32u srcPitch, NCVMemoryType srcType,
500  Ncv32u widthbytes, Ncv32u height, cudaStream_t cuStream);
501 
502 
503 /**
504 * NCVVector (1D)
505 */
506 template <class T>
508 {
509  NCVVector(const NCVVector &) = delete;
510 
511 public:
512 
514  {
515  clear();
516  }
517 
518  virtual ~NCVVector() {}
519 
520  void clear()
521  {
522  _ptr = nullptr;
523  _length = 0;
524  _memtype = NCVMemoryTypeNone;
525  }
526 
527  NCVStatus copySolid(NCVVector<T> &dst, cudaStream_t cuStream, std::size_t howMuch=0) const
528  {
529  if (howMuch == 0)
530  {
531  ncvAssertReturn(dst._length == this->_length, NCV_MEM_COPY_ERROR);
532  howMuch = this->_length * sizeof(T);
533  }
534  else
535  {
536  ncvAssertReturn(dst._length * sizeof(T) >= howMuch &&
537  this->_length * sizeof(T) >= howMuch &&
538  howMuch > 0, NCV_MEM_COPY_ERROR);
539  }
540  ncvAssertReturn((this->_ptr != nullptr || this->_memtype == NCVMemoryTypeNone) &&
541  (dst._ptr != nullptr || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
542 
543  NCVStatus ncvStat = NCV_SUCCESS;
544  if (this->_memtype != NCVMemoryTypeNone)
545  {
546  ncvStat = memSegCopyHelper(dst._ptr, dst._memtype,
547  this->_ptr, this->_memtype,
548  howMuch, cuStream);
549  }
550 
551  return ncvStat;
552  }
553 
554  T *ptr() const {return this->_ptr;}
555  std::size_t length() const {return this->_length;}
556  NCVMemoryType memType() const {return this->_memtype;}
557 
558 protected:
559 
560  T *_ptr;
561  std::size_t _length;
562  NCVMemoryType _memtype;
563 };
564 
565 
566 /**
567 * NCVVectorAlloc
568 */
569 template <class T>
570 class NCVVectorAlloc : public NCVVector<T>
571 {
572  NCVVectorAlloc() = delete;
573  NCVVectorAlloc(const NCVVectorAlloc &) = delete;
574  NCVVectorAlloc& operator=(const NCVVectorAlloc<T>&) = delete;
575 
576 public:
577 
578  NCVVectorAlloc(INCVMemAllocator &allocator, Ncv32u length)
579  :
580  allocator(allocator)
581  {
582  NCVStatus ncvStat;
583 
584  this->clear();
585  this->allocatedMem.clear();
586 
587  ncvStat = allocator.alloc(this->allocatedMem, length * sizeof(T));
588  ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "NCVVectorAlloc ctor:: alloc failed", );
589 
590  this->_ptr = (T *)this->allocatedMem.begin.ptr;
591  this->_length = length;
592  this->_memtype = this->allocatedMem.begin.memtype;
593  }
594 
596  {
597  NCVStatus ncvStat;
598 
599  ncvStat = allocator.dealloc(this->allocatedMem);
600  ncvAssertPrintCheck(ncvStat == NCV_SUCCESS, "NCVVectorAlloc dtor:: dealloc failed");
601 
602  this->clear();
603  }
604 
605  NcvBool isMemAllocated() const
606  {
607  return (this->allocatedMem.begin.ptr != nullptr) || (this->allocator.isCounting());
608  }
609 
610  Ncv32u getAllocatorsAlignment() const
611  {
612  return allocator.alignment();
613  }
614 
616  {
617  return allocatedMem;
618  }
619 
620 private:
621  INCVMemAllocator &allocator;
622  NCVMemSegment allocatedMem;
623 };
624 
625 
626 /**
627 * NCVVectorReuse
628 */
629 template <class T>
630 class NCVVectorReuse : public NCVVector<T>
631 {
632  NCVVectorReuse() = delete;
633  NCVVectorReuse(const NCVVectorReuse &) = delete;
634 
635 public:
636 
637  explicit NCVVectorReuse(const NCVMemSegment &memSegment)
638  {
639  this->bReused = false;
640  this->clear();
641 
642  this->_length = memSegment.size / sizeof(T);
643  this->_ptr = (T *)memSegment.begin.ptr;
644  this->_memtype = memSegment.begin.memtype;
645 
646  this->bReused = true;
647  }
648 
649  NCVVectorReuse(const NCVMemSegment &memSegment, Ncv32u length)
650  {
651  this->bReused = false;
652  this->clear();
653 
654  ncvAssertPrintReturn(length * sizeof(T) <= memSegment.size, \
655  "NCVVectorReuse ctor:: memory binding failed due to size mismatch", );
656 
657  this->_length = length;
658  this->_ptr = (T *)memSegment.begin.ptr;
659  this->_memtype = memSegment.begin.memtype;
660 
661  this->bReused = true;
662  }
663 
664  NcvBool isMemReused() const
665  {
666  return this->bReused;
667  }
668 
669 private:
670 
671  NcvBool bReused;
672 };
673 
674 
675 /**
676 * NCVMatrix (2D)
677 */
678 template <class T>
680 {
681  NCVMatrix(const NCVMatrix &) = delete;
682 
683 public:
684 
686  {
687  clear();
688  }
689 
690  virtual ~NCVMatrix() {}
691 
692  void clear()
693  {
694  _ptr = nullptr;
695  _pitch = 0;
696  _width = 0;
697  _height = 0;
698  _memtype = NCVMemoryTypeNone;
699  }
700 
701  Ncv32u stride() const
702  {
703  return _pitch / sizeof(T);
704  }
705 
706  //a side effect of this function is that it copies everything in a single chunk, so the "padding" will be overwritten
707  NCVStatus copySolid(NCVMatrix<T> &dst, cudaStream_t cuStream, std::size_t howMuch=0) const
708  {
709  if (howMuch == 0)
710  {
711  ncvAssertReturn(dst._pitch == this->_pitch &&
712  dst._height == this->_height, NCV_MEM_COPY_ERROR);
713  howMuch = this->_pitch * this->_height;
714  }
715  else
716  {
717  ncvAssertReturn(dst._pitch * dst._height >= howMuch &&
718  this->_pitch * this->_height >= howMuch &&
719  howMuch > 0, NCV_MEM_COPY_ERROR);
720  }
721  ncvAssertReturn((this->_ptr != nullptr || this->_memtype == NCVMemoryTypeNone) &&
722  (dst._ptr != nullptr || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
723 
724  NCVStatus ncvStat = NCV_SUCCESS;
725  if (this->_memtype != NCVMemoryTypeNone)
726  {
727  ncvStat = memSegCopyHelper(dst._ptr, dst._memtype,
728  this->_ptr, this->_memtype,
729  howMuch, cuStream);
730  }
731 
732  return ncvStat;
733  }
734 
735  NCVStatus copy2D(NCVMatrix<T> &dst, NcvSize32u roi, cudaStream_t cuStream) const
736  {
737  ncvAssertReturn(this->width() >= roi.width && this->height() >= roi.height &&
738  dst.width() >= roi.width && dst.height() >= roi.height, NCV_MEM_COPY_ERROR);
739  ncvAssertReturn((this->_ptr != NULL || this->_memtype == NCVMemoryTypeNone) &&
740  (dst._ptr != NULL || dst._memtype == NCVMemoryTypeNone), NCV_NULL_PTR);
741 
742  NCVStatus ncvStat = NCV_SUCCESS;
743  if (this->_memtype != NCVMemoryTypeNone)
744  {
745  ncvStat = memSegCopyHelper2D(dst._ptr, dst._pitch, dst._memtype,
746  this->_ptr, this->_pitch, this->_memtype,
747  roi.width * sizeof(T), roi.height, cuStream);
748  }
749 
750  return ncvStat;
751  }
752 
753  T &at(Ncv32u x, Ncv32u y) const
754  {
755  NcvBool bOutRange = (x >= this->_width || y >= this->_height);
756  ncvAssertPrintCheck(!bOutRange, "Error addressing matrix at [" << x << ", " << y << "]");
757  if (bOutRange)
758  {
759  return *this->_ptr;
760  }
761  return ((T *)((Ncv8u *)this->_ptr + y * this->_pitch))[x];
762  }
763 
764  T *ptr() const {return this->_ptr;}
765  Ncv32u width() const {return this->_width;}
766  Ncv32u height() const {return this->_height;}
767  NcvSize32u size() const {return NcvSize32u(this->_width, this->_height);}
768  Ncv32u pitch() const {return this->_pitch;}
769  NCVMemoryType memType() const {return this->_memtype;}
770 
771 protected:
772 
773  T *_ptr;
774  Ncv32u _width;
775  Ncv32u _height;
776  Ncv32u _pitch;
777  NCVMemoryType _memtype;
778 };
779 
780 
781 /**
782 * NCVMatrixAlloc
783 */
784 template <class T>
785 class NCVMatrixAlloc : public NCVMatrix<T>
786 {
787  NCVMatrixAlloc() = delete;
788  NCVMatrixAlloc(const NCVMatrixAlloc &) = delete;
789  NCVMatrixAlloc& operator=(const NCVMatrixAlloc &) = delete;
790 public:
791 
792  NCVMatrixAlloc(INCVMemAllocator &allocator, Ncv32u width, Ncv32u height, Ncv32u pitch=0)
793  :
794  allocator(allocator)
795  {
796  NCVStatus ncvStat;
797 
798  this->clear();
799  this->allocatedMem.clear();
800 
801  Ncv32u widthBytes = width * sizeof(T);
802  Ncv32u pitchBytes = alignUp(widthBytes, allocator.alignment());
803 
804  if (pitch != 0)
805  {
806  ncvAssertPrintReturn(pitch >= pitchBytes &&
807  (pitch & (allocator.alignment() - 1)) == 0,
808  "NCVMatrixAlloc ctor:: incorrect pitch passed", );
809  pitchBytes = pitch;
810  }
811 
812  Ncv32u requiredAllocSize = pitchBytes * height;
813 
814  ncvStat = allocator.alloc(this->allocatedMem, requiredAllocSize);
815  ncvAssertPrintReturn(ncvStat == NCV_SUCCESS, "NCVMatrixAlloc ctor:: alloc failed", );
816 
817  this->_ptr = (T *)this->allocatedMem.begin.ptr;
818  this->_width = width;
819  this->_height = height;
820  this->_pitch = pitchBytes;
821  this->_memtype = this->allocatedMem.begin.memtype;
822  }
823 
825  {
826  NCVStatus ncvStat;
827 
828  ncvStat = allocator.dealloc(this->allocatedMem);
829  ncvAssertPrintCheck(ncvStat == NCV_SUCCESS, "NCVMatrixAlloc dtor:: dealloc failed");
830 
831  this->clear();
832  }
833 
834  NcvBool isMemAllocated() const
835  {
836  return (this->allocatedMem.begin.ptr != nullptr) || (this->allocator.isCounting());
837  }
838 
839  Ncv32u getAllocatorsAlignment() const
840  {
841  return allocator.alignment();
842  }
843 
845  {
846  return allocatedMem;
847  }
848 
849 private:
850 
851  INCVMemAllocator &allocator;
852  NCVMemSegment allocatedMem;
853 };
854 
855 
856 /**
857 * NCVMatrixReuse
858 */
859 template <class T>
860 class NCVMatrixReuse : public NCVMatrix<T>
861 {
862  NCVMatrixReuse() = delete;
863  NCVMatrixReuse(const NCVMatrixReuse &) = delete;
864 
865 public:
866 
867  NCVMatrixReuse(const NCVMemSegment &memSegment, Ncv32u alignment, Ncv32u width, Ncv32u height, Ncv32u pitch=0, NcvBool bSkipPitchCheck=false)
868  {
869  this->bReused = false;
870  this->clear();
871 
872  Ncv32u widthBytes = width * sizeof(T);
873  Ncv32u pitchBytes = alignUp(widthBytes, alignment);
874 
875  if (pitch != 0)
876  {
877  if (!bSkipPitchCheck)
878  {
879  ncvAssertPrintReturn(pitch >= pitchBytes &&
880  (pitch & (alignment - 1)) == 0,
881  "NCVMatrixReuse ctor:: incorrect pitch passed", );
882  }
883  else
884  {
885  ncvAssertPrintReturn(pitch >= widthBytes, "NCVMatrixReuse ctor:: incorrect pitch passed", );
886  }
887  pitchBytes = pitch;
888  }
889 
890  ncvAssertPrintReturn(pitchBytes * height <= memSegment.size, \
891  "NCVMatrixReuse ctor:: memory binding failed due to size mismatch", );
892 
893  this->_width = width;
894  this->_height = height;
895  this->_pitch = pitchBytes;
896  this->_ptr = (T *)memSegment.begin.ptr;
897  this->_memtype = memSegment.begin.memtype;
898 
899  this->bReused = true;
900  }
901 
903  {
904  this->bReused = false;
905  this->clear();
906 
907  ncvAssertPrintReturn(roi.x < mat.width() && roi.y < mat.height() && \
908  roi.x + roi.width <= mat.width() && roi.y + roi.height <= mat.height(),
909  "NCVMatrixReuse ctor:: memory binding failed due to mismatching ROI and source matrix dims", );
910 
911  this->_width = roi.width;
912  this->_height = roi.height;
913  this->_pitch = mat.pitch();
914  this->_ptr = &mat.at(roi.x, roi.y);
915  this->_memtype = mat.memType();
916 
917  this->bReused = true;
918  }
919 
920  NcvBool isMemReused() const
921  {
922  return this->bReused;
923  }
924 
925 private:
926 
927  NcvBool bReused;
928 };
929 
930 /**
931 * Operations with rectangles
932 */
933 NCV_EXPORTS NCVStatus ncvGroupRectangles_host(NCVVector<NcvRect32u> &hypotheses, Ncv32u &numHypotheses,
934  Ncv32u minNeighbors, Ncv32f intersectEps, NCVVector<Ncv32u> *hypothesesWeights);
935 
936 NCV_EXPORTS NCVStatus ncvDrawRects_8u_host(Ncv8u *h_dst, Ncv32u dstStride, Ncv32u dstWidth, Ncv32u dstHeight,
937  NcvRect32u *h_rects, Ncv32u numRects, Ncv8u color);
938 
939 NCV_EXPORTS NCVStatus ncvDrawRects_32u_host(Ncv32u *h_dst, Ncv32u dstStride, Ncv32u dstWidth, Ncv32u dstHeight,
940  NcvRect32u *h_rects, Ncv32u numRects, Ncv32u color);
941 
942 NCV_EXPORTS NCVStatus ncvDrawRects_8u_device(Ncv8u *d_dst, Ncv32u dstStride, Ncv32u dstWidth, Ncv32u dstHeight,
943  NcvRect32u *d_rects, Ncv32u numRects, Ncv8u color, cudaStream_t cuStream);
944 
945 NCV_EXPORTS NCVStatus ncvDrawRects_32u_device(Ncv32u *d_dst, Ncv32u dstStride, Ncv32u dstWidth, Ncv32u dstHeight,
946  NcvRect32u *d_rects, Ncv32u numRects, Ncv32u color, cudaStream_t cuStream);
947 
948 #define CLAMP(x,a,b) ( (x) > (b) ? (b) : ( (x) < (a) ? (a) : (x) ) )
949 #define CLAMP_TOP(x, a) (((x) > (a)) ? (a) : (x))
950 #define CLAMP_BOTTOM(x, a) (((x) < (a)) ? (a) : (x))
951 #define CLAMP_0_255(x) CLAMP(x,0,255)
952 
953 #define SUB_BEGIN(type, name) struct { __inline type name
954 #define SUB_END(name) } name;
955 #define SUB_CALL(name) name.name
956 
957 #define SQR(x) ((x)*(x))
958 
959 #define ncvSafeMatAlloc(name, type, alloc, width, height, err) \
960  NCVMatrixAlloc<type> name(alloc, width, height); \
961  ncvAssertReturn(name.isMemAllocated(), err);
962 
963 #endif // PCL_GPU_PEOPLE__NCV_HPP_
Ncv32u height() const
Definition: NCV.hpp:766
__host__ __device__ NcvRect32u()
Definition: NCV.hpp:154
NCVMatrixAlloc(INCVMemAllocator &allocator, Ncv32u width, Ncv32u height, Ncv32u pitch=0)
Definition: NCV.hpp:792
__host__ __device__ NcvSize32s(Ncv32s width, Ncv32s height)
Definition: NCV.hpp:163
Ncv32u getAllocatorsAlignment() const
Definition: NCV.hpp:839
Ncv32s y
Point Y.
Definition: NCV.hpp:178
Ncv8u y
Definition: NCV.hpp:131
NcvBool isMemReused() const
Definition: NCV.hpp:664
T & at(Ncv32u x, Ncv32u y) const
Definition: NCV.hpp:753
NCVMemStackAllocator.
Definition: NCV.hpp:424
Compile-time assert namespace.
Definition: NCV.hpp:68
NCVVector()
Definition: NCV.hpp:513
NCVMatrixAlloc.
Definition: NCV.hpp:785
NCVMemoryType _memtype
Definition: NCV.hpp:777
Ncv32s height
Rectangle height.
Definition: NCV.hpp:143
T * _ptr
Definition: NCV.hpp:560
NCVMemoryType _memtype
Definition: NCV.hpp:562
NCVMemoryType memType() const
Definition: NCV.hpp:556
Ncv32s width
Rectangle width.
Definition: NCV.hpp:160
T * ptr() const
Definition: NCV.hpp:554
NCVVectorReuse(const NCVMemSegment &memSegment)
Definition: NCV.hpp:637
Ncv32u height
Rectangle height.
Definition: NCV.hpp:169
NCVMemSegment.
Definition: NCV.hpp:391
Ncv32u x
Point X.
Definition: NCV.hpp:185
NCVStatus copySolid(NCVMatrix< T > &dst, cudaStream_t cuStream, std::size_t howMuch=0) const
Definition: NCV.hpp:707
Ncv8u width
Definition: NCV.hpp:132
NCVMemNativeAllocator.
Definition: NCV.hpp:461
Ncv32s width
Rectangle width.
Definition: NCV.hpp:142
T * _ptr
Definition: NCV.hpp:773
Ncv8u x
Definition: NCV.hpp:130
__host__ __device__ NcvPoint2D32s(Ncv32s x, Ncv32s y)
Definition: NCV.hpp:180
Ncv32u width
Rectangle width.
Definition: NCV.hpp:152
NCVMemSegment getSegment() const
Definition: NCV.hpp:844
std::size_t length() const
Definition: NCV.hpp:555
NCVVectorAlloc.
Definition: NCV.hpp:570
Ncv32s x
Point X.
Definition: NCV.hpp:177
NcvSize32u size() const
Definition: NCV.hpp:767
virtual ~NCVVector()
Definition: NCV.hpp:518
__host__ __device__ NcvSize32s()
Definition: NCV.hpp:162
std::size_t _length
Definition: NCV.hpp:561
NCVStatus copySolid(NCVVector< T > &dst, cudaStream_t cuStream, std::size_t howMuch=0) const
Definition: NCV.hpp:527
Ncv32u _height
Definition: NCV.hpp:775
NCVVector (1D)
Definition: NCV.hpp:507
NCVVectorReuse.
Definition: NCV.hpp:630
void * ptr
Definition: NCV.hpp:382
Ncv32u y
Point Y.
Definition: NCV.hpp:186
__host__ __device__ NcvPoint2D32u(Ncv32u x, Ncv32u y)
Definition: NCV.hpp:188
void clear()
Definition: NCV.hpp:520
NCVMatrixReuse(const NCVMemSegment &memSegment, Ncv32u alignment, Ncv32u width, Ncv32u height, Ncv32u pitch=0, NcvBool bSkipPitchCheck=false)
Definition: NCV.hpp:867
Ncv32u y
y-coordinate of upper left corner.
Definition: NCV.hpp:151
NCVVectorReuse(const NCVMemSegment &memSegment, Ncv32u length)
Definition: NCV.hpp:649
NCVMatrix()
Definition: NCV.hpp:685
Ncv8u height
Definition: NCV.hpp:133
Ncv32u getAllocatorsAlignment() const
Definition: NCV.hpp:610
NCVMemSegment getSegment() const
Definition: NCV.hpp:615
__host__ __device__ NcvPoint2D32u()
Definition: NCV.hpp:187
Ncv32u width() const
Definition: NCV.hpp:765
Ncv32s height
Rectangle height.
Definition: NCV.hpp:161
__host__ __device__ NcvSize32u(Ncv32u width, Ncv32u height)
Definition: NCV.hpp:171
__host__ __device__ NcvPoint2D32s()
Definition: NCV.hpp:179
~NCVVectorAlloc()
Definition: NCV.hpp:595
__host__ __device__ NcvSize32u()
Definition: NCV.hpp:170
NcvBool isMemAllocated() const
Definition: NCV.hpp:834
INCVMemAllocator (Interface)
Definition: NCV.hpp:402
NCVMemPtr.
Definition: NCV.hpp:380
Ncv32u pitch() const
Definition: NCV.hpp:768
Ncv32u height
Rectangle height.
Definition: NCV.hpp:153
NcvBool isMemAllocated() const
Definition: NCV.hpp:605
NCVVectorAlloc(INCVMemAllocator &allocator, Ncv32u length)
Definition: NCV.hpp:578
NCVStatus copy2D(NCVMatrix< T > &dst, NcvSize32u roi, cudaStream_t cuStream) const
Definition: NCV.hpp:735
void clear()
Definition: NCV.hpp:692
virtual ~INCVMemAllocator()=0
Definition: NCV.hpp:418
__host__ __device__ NcvRect32s()
Definition: NCV.hpp:144
T * ptr() const
Definition: NCV.hpp:764
~NCVMatrixAlloc()
Definition: NCV.hpp:824
virtual NCVStatus alloc(NCVMemSegment &seg, std::size_t size)=0
NCVMemoryType memtype
Definition: NCV.hpp:383
Ncv32s y
y-coordinate of upper left corner.
Definition: NCV.hpp:141
NCVMatrix (2D)
Definition: NCV.hpp:679
NCVMemoryType memType() const
Definition: NCV.hpp:769
NcvBool isMemReused() const
Definition: NCV.hpp:920
__host__ __device__ NcvRect8u(Ncv8u x, Ncv8u y, Ncv8u width, Ncv8u height)
Definition: NCV.hpp:135
Ncv32u x
x-coordinate of upper left corner.
Definition: NCV.hpp:150
Ncv32u width
Rectangle width.
Definition: NCV.hpp:168
Ncv32s x
x-coordinate of upper left corner.
Definition: NCV.hpp:140
NCVMatrixReuse.
Definition: NCV.hpp:860
__host__ __device__ NcvRect8u()
Definition: NCV.hpp:134
Ncv32u stride() const
Definition: NCV.hpp:701
std::size_t size
Definition: NCV.hpp:394
virtual ~NCVMatrix()
Definition: NCV.hpp:690
Ncv32u _pitch
Definition: NCV.hpp:776
__host__ __device__ NcvRect32s(Ncv32s x, Ncv32s y, Ncv32s width, Ncv32s height)
Definition: NCV.hpp:145
virtual Ncv32u alignment() const =0
Ncv32u _width
Definition: NCV.hpp:774
NCVMemPtr begin
Definition: NCV.hpp:393
__host__ __device__ NcvRect32u(Ncv32u x, Ncv32u y, Ncv32u width, Ncv32u height)
Definition: NCV.hpp:155
NCVMatrixReuse(const NCVMatrix< T > &mat, NcvRect32u roi)
Definition: NCV.hpp:902