pyr_down.hpp 8.0 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201
  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. #pragma once
  44. #ifndef OPENCV_CUDEV_GRID_PYR_DOWN_DETAIL_HPP
  45. #define OPENCV_CUDEV_GRID_PYR_DOWN_DETAIL_HPP
  46. #include "../../common.hpp"
  47. #include "../../util/vec_traits.hpp"
  48. #include "../../util/saturate_cast.hpp"
  49. #include "../../util/type_traits.hpp"
  50. #include "../../ptr2d/glob.hpp"
  51. #include "../../ptr2d/traits.hpp"
  52. namespace cv { namespace cudev {
  53. namespace pyramids_detail
  54. {
  55. template <class Brd, class SrcPtr, typename DstType>
  56. __global__ void pyrDown(const SrcPtr src, GlobPtr<DstType> dst, const int src_rows, const int src_cols, const int dst_cols)
  57. {
  58. typedef typename PtrTraits<SrcPtr>::value_type src_type;
  59. typedef typename VecTraits<src_type>::elem_type src_elem_type;
  60. typedef typename LargerType<float, src_elem_type>::type work_elem_type;
  61. typedef typename MakeVec<work_elem_type, VecTraits<src_type>::cn>::type work_type;
  62. __shared__ work_type smem[256 + 4];
  63. const int x = blockIdx.x * blockDim.x + threadIdx.x;
  64. const int y = blockIdx.y;
  65. const int src_y = 2 * y;
  66. if (src_y >= 2 && src_y < src_rows - 2 && x >= 2 && x < src_cols - 2)
  67. {
  68. {
  69. work_type sum;
  70. sum = 0.0625f * src(src_y - 2, x);
  71. sum = sum + 0.25f * src(src_y - 1, x);
  72. sum = sum + 0.375f * src(src_y , x);
  73. sum = sum + 0.25f * src(src_y + 1, x);
  74. sum = sum + 0.0625f * src(src_y + 2, x);
  75. smem[2 + threadIdx.x] = sum;
  76. }
  77. if (threadIdx.x < 2)
  78. {
  79. const int left_x = x - 2;
  80. work_type sum;
  81. sum = 0.0625f * src(src_y - 2, left_x);
  82. sum = sum + 0.25f * src(src_y - 1, left_x);
  83. sum = sum + 0.375f * src(src_y , left_x);
  84. sum = sum + 0.25f * src(src_y + 1, left_x);
  85. sum = sum + 0.0625f * src(src_y + 2, left_x);
  86. smem[threadIdx.x] = sum;
  87. }
  88. if (threadIdx.x > 253)
  89. {
  90. const int right_x = x + 2;
  91. work_type sum;
  92. sum = 0.0625f * src(src_y - 2, right_x);
  93. sum = sum + 0.25f * src(src_y - 1, right_x);
  94. sum = sum + 0.375f * src(src_y , right_x);
  95. sum = sum + 0.25f * src(src_y + 1, right_x);
  96. sum = sum + 0.0625f * src(src_y + 2, right_x);
  97. smem[4 + threadIdx.x] = sum;
  98. }
  99. }
  100. else
  101. {
  102. {
  103. work_type sum;
  104. sum = 0.0625f * src(Brd::idx_low(src_y - 2, src_rows) , Brd::idx_high(x, src_cols));
  105. sum = sum + 0.25f * src(Brd::idx_low(src_y - 1, src_rows) , Brd::idx_high(x, src_cols));
  106. sum = sum + 0.375f * src(src_y , Brd::idx_high(x, src_cols));
  107. sum = sum + 0.25f * src(Brd::idx_high(src_y + 1, src_rows), Brd::idx_high(x, src_cols));
  108. sum = sum + 0.0625f * src(Brd::idx_high(src_y + 2, src_rows), Brd::idx_high(x, src_cols));
  109. smem[2 + threadIdx.x] = sum;
  110. }
  111. if (threadIdx.x < 2)
  112. {
  113. const int left_x = x - 2;
  114. work_type sum;
  115. sum = 0.0625f * src(Brd::idx_low(src_y - 2, src_rows) , Brd::idx_low(Brd::idx_high(left_x, src_cols), src_cols));
  116. sum = sum + 0.25f * src(Brd::idx_low(src_y - 1, src_rows) , Brd::idx_low(Brd::idx_high(left_x, src_cols), src_cols));
  117. sum = sum + 0.375f * src(src_y , Brd::idx_low(Brd::idx_high(left_x, src_cols), src_cols));
  118. sum = sum + 0.25f * src(Brd::idx_high(src_y + 1, src_rows), Brd::idx_low(Brd::idx_high(left_x, src_cols), src_cols));
  119. sum = sum + 0.0625f * src(Brd::idx_high(src_y + 2, src_rows), Brd::idx_low(Brd::idx_high(left_x, src_cols), src_cols));
  120. smem[threadIdx.x] = sum;
  121. }
  122. if (threadIdx.x > 253)
  123. {
  124. const int right_x = x + 2;
  125. work_type sum;
  126. sum = 0.0625f * src(Brd::idx_low(src_y - 2, src_rows) , Brd::idx_high(right_x, src_cols));
  127. sum = sum + 0.25f * src(Brd::idx_low(src_y - 1, src_rows) , Brd::idx_high(right_x, src_cols));
  128. sum = sum + 0.375f * src(src_y , Brd::idx_high(right_x, src_cols));
  129. sum = sum + 0.25f * src(Brd::idx_high(src_y + 1, src_rows), Brd::idx_high(right_x, src_cols));
  130. sum = sum + 0.0625f * src(Brd::idx_high(src_y + 2, src_rows), Brd::idx_high(right_x, src_cols));
  131. smem[4 + threadIdx.x] = sum;
  132. }
  133. }
  134. __syncthreads();
  135. if (threadIdx.x < 128)
  136. {
  137. const int tid2 = threadIdx.x * 2;
  138. work_type sum;
  139. sum = 0.0625f * smem[2 + tid2 - 2];
  140. sum = sum + 0.25f * smem[2 + tid2 - 1];
  141. sum = sum + 0.375f * smem[2 + tid2 ];
  142. sum = sum + 0.25f * smem[2 + tid2 + 1];
  143. sum = sum + 0.0625f * smem[2 + tid2 + 2];
  144. const int dst_x = (blockIdx.x * blockDim.x + tid2) / 2;
  145. if (dst_x < dst_cols)
  146. dst(y, dst_x) = saturate_cast<DstType>(sum);
  147. }
  148. }
  149. template <class Brd, class SrcPtr, typename DstType>
  150. __host__ void pyrDown(const SrcPtr& src, const GlobPtr<DstType>& dst, int src_rows, int src_cols, int dst_rows, int dst_cols, cudaStream_t stream)
  151. {
  152. const dim3 block(256);
  153. const dim3 grid(divUp(src_cols, block.x), dst_rows);
  154. pyrDown<Brd><<<grid, block, 0, stream>>>(src, dst, src_rows, src_cols, dst_cols);
  155. CV_CUDEV_SAFE_CALL( cudaGetLastError() );
  156. if (stream == 0)
  157. CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
  158. }
  159. }
  160. }}
  161. #endif