potok/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-07-31 00:03:03 +04:00
Code Packages Releases Wiki Activity

Merge remote-tracking branch 'origin/2.4'

Browse Source 
Conflicts:
	3rdparty/ffmpeg/ffmpeg_version.cmake
	cmake/OpenCVFindLibsGrfmt.cmake
	cmake/templates/cvconfig.h.cmake
	modules/bioinspired/doc/retina/index.rst
	modules/calib3d/doc/camera_calibration_and_3d_reconstruction.rst
	modules/calib3d/src/precomp.hpp
	modules/contrib/src/inputoutput.cpp
	modules/contrib/src/precomp.hpp
	modules/core/include/opencv2/core/internal.hpp
	modules/core/include/opencv2/core/types_c.h
	modules/core/src/drawing.cpp
	modules/core/src/precomp.hpp
	modules/core/src/system.cpp
	modules/features2d/doc/common_interfaces_of_descriptor_matchers.rst
	modules/features2d/doc/common_interfaces_of_feature_detectors.rst
	modules/features2d/include/opencv2/features2d/features2d.hpp
	modules/features2d/src/precomp.hpp
	modules/flann/src/precomp.hpp
	modules/gpu/doc/camera_calibration_and_3d_reconstruction.rst
	modules/gpu/doc/image_filtering.rst
	modules/gpu/doc/image_processing.rst
	modules/gpu/doc/video.rst
	modules/gpu/perf/perf_imgproc.cpp
	modules/gpu/perf4au/main.cpp
	modules/gpu/src/imgproc.cpp
	modules/gpu/src/precomp.hpp
	modules/gpu/test/test_imgproc.cpp
	modules/highgui/CMakeLists.txt
	modules/highgui/test/test_precomp.hpp
	modules/imgproc/doc/structural_analysis_and_shape_descriptors.rst
	modules/imgproc/src/precomp.hpp
	modules/java/generator/src/cpp/Mat.cpp
	modules/legacy/src/precomp.hpp
	modules/ml/doc/k_nearest_neighbors.rst
	modules/ml/src/precomp.hpp
	modules/nonfree/doc/feature_detection.rst
	modules/nonfree/src/precomp.hpp
	modules/objdetect/include/opencv2/objdetect/objdetect.hpp
	modules/objdetect/src/cascadedetect.cpp
	modules/objdetect/src/hog.cpp
	modules/objdetect/src/precomp.hpp
	modules/objdetect/test/test_latentsvmdetector.cpp
	modules/ocl/src/hog.cpp
	modules/ocl/src/opencl/objdetect_hog.cl
	modules/ocl/src/precomp.hpp
	modules/photo/src/precomp.hpp
	modules/stitching/src/precomp.hpp
	modules/superres/perf/perf_precomp.hpp
	modules/superres/src/optical_flow.cpp
	modules/superres/src/precomp.hpp
	modules/superres/test/test_precomp.hpp
	modules/ts/include/opencv2/ts.hpp
	modules/video/src/precomp.hpp
	modules/videostab/src/precomp.hpp
	modules/world/src/precomp.hpp
This commit is contained in:
Roman Donchenko
2013-08-06 13:56:49 +04:00
parent 4c03afe006 b2d1d87ed1
commit dcaf923517
119 changed files with 3713 additions and 2281 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/gpuimgproc/doc/color.rst
+3
View File
@@ -123,6 +123,9 @@ Composites two images using alpha opacity values contained in each image.
:param stream: Stream for the asynchronous version.
.. Sample code::
* : An example demonstrating the use of alphaComp can be found at opencv_source_code/samples/gpu/alpha_comp.cpp
.. [MHT2011] Pascal Getreuer, Malvar-He-Cutler Linear Image Demosaicking, Image Processing On Line, 2011
modules/gpuimgproc/doc/hough.rst
+3
View File
@@ -108,6 +108,9 @@ Base class for line segments detector algorithm. ::
virtual int getMaxLines() const = 0;
};
.. Sample code::
* : An example using the Hough segment detector can be found at opencv_source_code/samples/gpu/houghlines.cpp
gpu::HoughSegmentDetector::detect
modules/gpuimgproc/src/cuda/hist.cu
+80
View File
@@ -109,6 +109,86 @@ namespace hist
/////////////////////////////////////////////////////////////////////////
namespace hist
{
__device__ __forceinline__ void histEvenInc(int* shist, uint data, int binSize, int lowerLevel, int upperLevel)
{
if (data >= lowerLevel && data <= upperLevel)
{
const uint ind = (data - lowerLevel) / binSize;
Emulation::smem::atomicAdd(shist + ind, 1);
}
}
__global__ void histEven8u(const uchar* src, const size_t step, const int rows, const int cols,
int* hist, const int binCount, const int binSize, const int lowerLevel, const int upperLevel)
{
extern __shared__ int shist[];
const int y = blockIdx.x * blockDim.y + threadIdx.y;
const int tid = threadIdx.y * blockDim.x + threadIdx.x;
if (tid < binCount)
shist[tid] = 0;
__syncthreads();
if (y < rows)
{
const uchar* rowPtr = src + y * step;
const uint* rowPtr4 = (uint*) rowPtr;
const int cols_4 = cols / 4;
for (int x = threadIdx.x; x < cols_4; x += blockDim.x)
{
const uint data = rowPtr4[x];
histEvenInc(shist, (data >> 0) & 0xFFU, binSize, lowerLevel, upperLevel);
histEvenInc(shist, (data >> 8) & 0xFFU, binSize, lowerLevel, upperLevel);
histEvenInc(shist, (data >> 16) & 0xFFU, binSize, lowerLevel, upperLevel);
histEvenInc(shist, (data >> 24) & 0xFFU, binSize, lowerLevel, upperLevel);
}
if (cols % 4 != 0 && threadIdx.x == 0)
{
for (int x = cols_4 * 4; x < cols; ++x)
{
const uchar data = rowPtr[x];
histEvenInc(shist, data, binSize, lowerLevel, upperLevel);
}
}
}
__syncthreads();
if (tid < binCount)
{
const int histVal = shist[tid];
if (histVal > 0)
::atomicAdd(hist + tid, histVal);
}
}
void histEven8u(PtrStepSzb src, int* hist, int binCount, int lowerLevel, int upperLevel, cudaStream_t stream)
{
const dim3 block(32, 8);
const dim3 grid(divUp(src.rows, block.y));
const int binSize = divUp(upperLevel - lowerLevel, binCount);
const size_t smem_size = binCount * sizeof(int);
histEven8u<<<grid, block, smem_size, stream>>>(src.data, src.step, src.rows, src.cols, hist, binCount, binSize, lowerLevel, upperLevel);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
/////////////////////////////////////////////////////////////////////////
namespace hist
{
__constant__ int c_lut[256];
modules/gpuimgproc/src/histogram.cpp
+21
View File
@@ -478,6 +478,21 @@ void cv::gpu::evenLevels(OutputArray _levels, int nLevels, int lowerLevel, int u
_levels.getGpuMatRef().upload(host_levels);
}
namespace hist
{
void histEven8u(PtrStepSzb src, int* hist, int binCount, int lowerLevel, int upperLevel, cudaStream_t stream);
}
namespace
{
void histEven8u(const GpuMat& src, GpuMat& hist, int histSize, int lowerLevel, int upperLevel, cudaStream_t stream)
{
hist.create(1, histSize, CV_32S);
cudaSafeCall( cudaMemsetAsync(hist.data, 0, histSize * sizeof(int), stream) );
hist::histEven8u(src, hist.ptr<int>(), histSize, lowerLevel, upperLevel, stream);
}
}
void cv::gpu::histEven(InputArray _src, OutputArray hist, InputOutputArray buf, int histSize, int lowerLevel, int upperLevel, Stream& stream)
{
typedef void (*hist_t)(const GpuMat& src, OutputArray hist, InputOutputArray buf, int levels, int lowerLevel, int upperLevel, cudaStream_t stream);
@@ -491,6 +506,12 @@ void cv::gpu::histEven(InputArray _src, OutputArray hist, InputOutputArray buf,
GpuMat src = _src.getGpuMat();
if (src.depth() == CV_8U && deviceSupports(FEATURE_SET_COMPUTE_30))
{
histEven8u(src, hist.getGpuMatRef(), histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));
return;
}
CV_Assert( src.type() == CV_8UC1 || src.type() == CV_16UC1 || src.type() == CV_16SC1 );
hist_callers[src.depth()](src, hist, buf, histSize, lowerLevel, upperLevel, StreamAccessor::getStream(stream));
modules/gpuimgproc/test/test_histogram.cpp
+24 -35
View File
@@ -49,13 +49,16 @@ using namespace cvtest;
///////////////////////////////////////////////////////////////////////////////////////////////////////
// HistEven
struct HistEven : testing::TestWithParam<cv::gpu::DeviceInfo>
PARAM_TEST_CASE(HistEven, cv::gpu::DeviceInfo, cv::Size)
{
cv::gpu::DeviceInfo devInfo;
cv::Size size;
virtual void SetUp()
{
devInfo = GetParam();
devInfo = GET_PARAM(0);
size = GET_PARAM(1);
cv::gpu::setDevice(devInfo.deviceID());
}
@@ -63,57 +66,34 @@ struct HistEven : testing::TestWithParam<cv::gpu::DeviceInfo>
GPU_TEST_P(HistEven, Accuracy)
{
cv::Mat img = readImage("stereobm/aloe-L.png");
ASSERT_FALSE(img.empty());
cv::Mat hsv;
cv::cvtColor(img, hsv, cv::COLOR_BGR2HSV);
cv::Mat src = randomMat(size, CV_8UC1);
int hbins = 30;
float hranges[] = {0.0f, 180.0f};
std::vector<cv::Mat> srcs;
cv::split(hsv, srcs);
float hranges[] = {50.0f, 200.0f};
cv::gpu::GpuMat hist;
cv::gpu::histEven(loadMat(srcs[0]), hist, hbins, (int)hranges[0], (int)hranges[1]);
cv::gpu::histEven(loadMat(src), hist, hbins, (int) hranges[0], (int) hranges[1]);
cv::Mat hist_gold;
cv::MatND histnd;
int histSize[] = {hbins};
const float* ranges[] = {hranges};
int channels[] = {0};
cv::calcHist(&hsv, 1, channels, cv::Mat(), histnd, 1, histSize, ranges);
cv::calcHist(&src, 1, channels, cv::Mat(), hist_gold, 1, histSize, ranges);
cv::Mat hist_gold = histnd;
hist_gold = hist_gold.t();
hist_gold.convertTo(hist_gold, CV_32S);
EXPECT_MAT_NEAR(hist_gold, hist, 0.0);
}
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, HistEven, ALL_DEVICES);
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, HistEven, testing::Combine(
ALL_DEVICES,
DIFFERENT_SIZES));
///////////////////////////////////////////////////////////////////////////////////////////////////////
// CalcHist
namespace
{
void calcHistGold(const cv::Mat& src, cv::Mat& hist)
{
hist.create(1, 256, CV_32SC1);
hist.setTo(cv::Scalar::all(0));
int* hist_row = hist.ptr<int>();
for (int y = 0; y < src.rows; ++y)
{
const uchar* src_row = src.ptr(y);
for (int x = 0; x < src.cols; ++x)
++hist_row[src_row[x]];
}
}
}
PARAM_TEST_CASE(CalcHist, cv::gpu::DeviceInfo, cv::Size)
{
cv::gpu::DeviceInfo devInfo;
@@ -137,7 +117,16 @@ GPU_TEST_P(CalcHist, Accuracy)
cv::gpu::calcHist(loadMat(src), hist);
cv::Mat hist_gold;
calcHistGold(src, hist_gold);
const int hbins = 256;
const float hranges[] = {0.0f, 256.0f};
const int histSize[] = {hbins};
const float* ranges[] = {hranges};
const int channels[] = {0};
cv::calcHist(&src, 1, channels, cv::Mat(), hist_gold, 1, histSize, ranges);
hist_gold = hist_gold.reshape(1, 1);
hist_gold.convertTo(hist_gold, CV_32S);
EXPECT_MAT_NEAR(hist_gold, hist, 0.0);
}