modules/gpu/src/cuda/hist.cu

   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) 1993-2011, NVIDIA Corporation, 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 bpied warranties, including, but not limited to, the bpied
  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 #if !defined CUDA_DISABLER
  45
  46 #include "opencv2/gpu/device/common.hpp"
  47 #include "opencv2/gpu/device/functional.hpp"
  48 #include "opencv2/gpu/device/emulation.hpp"
  49 #include "opencv2/gpu/device/transform.hpp"
  50
  51 using namespace cv::gpu;
  52 using namespace cv::gpu::device;
  53
  54 namespace hist
  55 {
  56     __global__ void histogram256Kernel(const uchar* src, int cols, int rows, size_t step, int* hist)
  57     {
  58         __shared__ int shist[256];
  59
  60         const int y = blockIdx.x * blockDim.y + threadIdx.y;
  61         const int tid = threadIdx.y * blockDim.x + threadIdx.x;
  62
  63         shist[tid] = 0;
  64         __syncthreads();
  65
  66         if (y < rows)
  67         {
  68             const unsigned int* rowPtr = (const unsigned int*) (src + y * step);
  69
  70             const int cols_4 = cols / 4;
  71             for (int x = threadIdx.x; x < cols_4; x += blockDim.x)
  72             {
  73                 unsigned int data = rowPtr[x];
  74
  75                 Emulation::smem::atomicAdd(&shist[(data >>  0) & 0xFFU], 1);
  76                 Emulation::smem::atomicAdd(&shist[(data >>  8) & 0xFFU], 1);
  77                 Emulation::smem::atomicAdd(&shist[(data >> 16) & 0xFFU], 1);
  78                 Emulation::smem::atomicAdd(&shist[(data >> 24) & 0xFFU], 1);
  79             }
  80
  81             if (cols % 4 != 0 && threadIdx.x == 0)
  82             {
  83                 for (int x = cols_4 * 4; x < cols; ++x)
  84                 {
  85                     unsigned int data = ((const uchar*)rowPtr)[x];
  86                     Emulation::smem::atomicAdd(&shist[data], 1);
  87                 }
  88             }
  89         }
  90
  91         __syncthreads();
  92
  93         const int histVal = shist[tid];
  94         if (histVal > 0)
  95             ::atomicAdd(hist + tid, histVal);
  96     }
  97
  98     void histogram256(PtrStepSzb src, int* hist, cudaStream_t stream)
  99     {
 100         const dim3 block(32, 8);
 101         const dim3 grid(divUp(src.rows, block.y));
 102
 103         histogram256Kernel<<<grid, block, 0, stream>>>(src.data, src.cols, src.rows, src.step, hist);
 104         cudaSafeCall( cudaGetLastError() );
 105
 106         if (stream == 0)
 107             cudaSafeCall( cudaDeviceSynchronize() );
 108     }
 109 }
 110
 111 /////////////////////////////////////////////////////////////////////////
 112
 113 namespace hist
 114 {
 115     __constant__ int c_lut[256];
 116
 117     struct EqualizeHist : unary_function<uchar, uchar>
 118     {
 119         float scale;
 120
 121         __host__ EqualizeHist(float _scale) : scale(_scale) {}
 122
 123         __device__ __forceinline__ uchar operator ()(uchar val) const
 124         {
 125             const int lut = c_lut[val];
 126             return __float2int_rn(scale * lut);
 127         }
 128     };
 129 }
 130
 131 namespace cv { namespace gpu { namespace device
 132 {
 133     template <> struct TransformFunctorTraits<hist::EqualizeHist> : DefaultTransformFunctorTraits<hist::EqualizeHist>
 134     {
 135         enum { smart_shift = 4 };
 136     };
 137 }}}
 138
 139 namespace hist
 140 {
 141     void equalizeHist(PtrStepSzb src, PtrStepSzb dst, const int* lut, cudaStream_t stream)
 142     {
 143         if (stream == 0)
 144             cudaSafeCall( cudaMemcpyToSymbol(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice) );
 145         else
 146             cudaSafeCall( cudaMemcpyToSymbolAsync(c_lut, lut, 256 * sizeof(int), 0, cudaMemcpyDeviceToDevice, stream) );
 147
 148         const float scale = 255.0f / (src.cols * src.rows);
 149
 150         transform(src, dst, EqualizeHist(scale), WithOutMask(), stream);
 151     }
 152 }
 153
 154 #endif /* CUDA_DISABLER */