← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/cudaimgproc/src/cuda/clahe.cu

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/cudaimgproc/src/cuda/clahe.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
// 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
#if !defined CUDA_DISABLER
 
44
 
 
45
#include "opencv2/core/cuda/common.hpp"
 
46
#include "opencv2/core/cuda/functional.hpp"
 
47
#include "opencv2/core/cuda/emulation.hpp"
 
48
#include "opencv2/core/cuda/scan.hpp"
 
49
#include "opencv2/core/cuda/reduce.hpp"
 
50
#include "opencv2/core/cuda/saturate_cast.hpp"
 
51
 
 
52
using namespace cv::cuda;
 
53
using namespace cv::cuda::device;
 
54
 
 
55
namespace clahe
 
56
{
 
57
    __global__ void calcLutKernel(const PtrStepb src, PtrStepb lut,
 
58
                                  const int2 tileSize, const int tilesX,
 
59
                                  const int clipLimit, const float lutScale)
 
60
    {
 
61
        __shared__ int smem[512];
 
62
 
 
63
        const int tx = blockIdx.x;
 
64
        const int ty = blockIdx.y;
 
65
        const unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x;
 
66
 
 
67
        smem[tid] = 0;
 
68
        __syncthreads();
 
69
 
 
70
        for (int i = threadIdx.y; i < tileSize.y; i += blockDim.y)
 
71
        {
 
72
            const uchar* srcPtr = src.ptr(ty * tileSize.y + i) + tx * tileSize.x;
 
73
            for (int j = threadIdx.x; j < tileSize.x; j += blockDim.x)
 
74
            {
 
75
                const int data = srcPtr[j];
 
76
                Emulation::smem::atomicAdd(&smem[data], 1);
 
77
            }
 
78
        }
 
79
 
 
80
        __syncthreads();
 
81
 
 
82
        int tHistVal = smem[tid];
 
83
 
 
84
        __syncthreads();
 
85
 
 
86
        if (clipLimit > 0)
 
87
        {
 
88
            // clip histogram bar
 
89
 
 
90
            int clipped = 0;
 
91
            if (tHistVal > clipLimit)
 
92
            {
 
93
                clipped = tHistVal - clipLimit;
 
94
                tHistVal = clipLimit;
 
95
            }
 
96
 
 
97
            // find number of overall clipped samples
 
98
 
 
99
            reduce<256>(smem, clipped, tid, plus<int>());
 
100
 
 
101
            // broadcast evaluated value
 
102
 
 
103
            __shared__ int totalClipped;
 
104
 
 
105
            if (tid == 0)
 
106
                totalClipped = clipped;
 
107
            __syncthreads();
 
108
 
 
109
            // redistribute clipped samples evenly
 
110
 
 
111
            int redistBatch = totalClipped / 256;
 
112
            tHistVal += redistBatch;
 
113
 
 
114
            int residual = totalClipped - redistBatch * 256;
 
115
            if (tid < residual)
 
116
                ++tHistVal;
 
117
        }
 
118
 
 
119
        const int lutVal = blockScanInclusive<256>(tHistVal, smem, tid);
 
120
 
 
121
        lut(ty * tilesX + tx, tid) = saturate_cast<uchar>(__float2int_rn(lutScale * lutVal));
 
122
    }
 
123
 
 
124
    void calcLut(PtrStepSzb src, PtrStepb lut, int tilesX, int tilesY, int2 tileSize, int clipLimit, float lutScale, cudaStream_t stream)
 
125
    {
 
126
        const dim3 block(32, 8);
 
127
        const dim3 grid(tilesX, tilesY);
 
128
 
 
129
        calcLutKernel<<<grid, block, 0, stream>>>(src, lut, tileSize, tilesX, clipLimit, lutScale);
 
130
 
 
131
        cudaSafeCall( cudaGetLastError() );
 
132
 
 
133
        if (stream == 0)
 
134
            cudaSafeCall( cudaDeviceSynchronize() );
 
135
    }
 
136
 
 
137
    __global__ void tranformKernel(const PtrStepSzb src, PtrStepb dst, const PtrStepb lut, const int2 tileSize, const int tilesX, const int tilesY)
 
138
    {
 
139
        const int x = blockIdx.x * blockDim.x + threadIdx.x;
 
140
        const int y = blockIdx.y * blockDim.y + threadIdx.y;
 
141
 
 
142
        if (x >= src.cols || y >= src.rows)
 
143
            return;
 
144
 
 
145
        const float tyf = (static_cast<float>(y) / tileSize.y) - 0.5f;
 
146
        int ty1 = __float2int_rd(tyf);
 
147
        int ty2 = ty1 + 1;
 
148
        const float ya = tyf - ty1;
 
149
        ty1 = ::max(ty1, 0);
 
150
        ty2 = ::min(ty2, tilesY - 1);
 
151
 
 
152
        const float txf = (static_cast<float>(x) / tileSize.x) - 0.5f;
 
153
        int tx1 = __float2int_rd(txf);
 
154
        int tx2 = tx1 + 1;
 
155
        const float xa = txf - tx1;
 
156
        tx1 = ::max(tx1, 0);
 
157
        tx2 = ::min(tx2, tilesX - 1);
 
158
 
 
159
        const int srcVal = src(y, x);
 
160
 
 
161
        float res = 0;
 
162
 
 
163
        res += lut(ty1 * tilesX + tx1, srcVal) * ((1.0f - xa) * (1.0f - ya));
 
164
        res += lut(ty1 * tilesX + tx2, srcVal) * ((xa) * (1.0f - ya));
 
165
        res += lut(ty2 * tilesX + tx1, srcVal) * ((1.0f - xa) * (ya));
 
166
        res += lut(ty2 * tilesX + tx2, srcVal) * ((xa) * (ya));
 
167
 
 
168
        dst(y, x) = saturate_cast<uchar>(res);
 
169
    }
 
170
 
 
171
    void transform(PtrStepSzb src, PtrStepSzb dst, PtrStepb lut, int tilesX, int tilesY, int2 tileSize, cudaStream_t stream)
 
172
    {
 
173
        const dim3 block(32, 8);
 
174
        const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
 
175
 
 
176
        cudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) );
 
177
 
 
178
        tranformKernel<<<grid, block, 0, stream>>>(src, dst, lut, tileSize, tilesX, tilesY);
 
179
        cudaSafeCall( cudaGetLastError() );
 
180
 
 
181
        if (stream == 0)
 
182
            cudaSafeCall( cudaDeviceSynchronize() );
 
183
    }
 
184
}
 
185
 
 
186
#endif // CUDA_DISABLER