From: Andrey Kamaev Date: Fri, 12 Apr 2013 11:35:38 +0000 (+0400) Subject: Merge branch '2.4' X-Git-Tag: submit/tizen_ivi/20141117.190038~2^2~1079 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=3b364330ad4603b452b69160fe3793c20fd13b71;p=profile%2Fivi%2Fopencv.git Merge branch '2.4' --- 3b364330ad4603b452b69160fe3793c20fd13b71 diff --cc modules/highgui/src/cap_v4l.cpp index 829d0ab,a788c90..6fc122f --- a/modules/highgui/src/cap_v4l.cpp +++ b/modules/highgui/src/cap_v4l.cpp @@@ -154,6 -154,11 +154,11 @@@ the symptoms were damaged image and 'Co - USE_TEMP_BUFFER fixes the main problem (improper buffer management) and prevents bad images in the first place + 11th patch: April 2, 2013, Forrest Reiling forrest.reiling@gmail.com -Added v4l2 support for getting capture property CV_CAP_PROP_POS_MSEC. ++Added v4l2 support for getting capture property CV_CAP_PROP_POS_MSEC. + Returns the millisecond timestamp of the last frame grabbed or 0 if no frames have been grabbed + Used to successfully synchonize 2 Logitech C310 USB webcams to within 16 ms of one another + make & enjoy! @@@ -1221,6 -1231,9 +1231,9 @@@ static int read_frame_v4l2(CvCaptureCAM if (-1 == ioctl (capture->deviceHandle, VIDIOC_QBUF, &buf)) perror ("VIDIOC_QBUF"); - //set timestamp in capture struct to be timestamp of most recent frame - capture->timestamp = buf.timestamp; ++ //set timestamp in capture struct to be timestamp of most recent frame ++ capture->timestamp = buf.timestamp; + return 1; } @@@ -2308,6 -2321,13 +2321,13 @@@ static double icvGetPropertyCAM_V4L (Cv /* initialize the control structure */ switch (property_id) { + case CV_CAP_PROP_POS_MSEC: + if (capture->FirstCapture) { + return 0; + } else { - return 1000 * capture->timestamp.tv_sec + ((double) capture->timestamp.tv_usec) / 1000; ++ return 1000 * capture->timestamp.tv_sec + ((double) capture->timestamp.tv_usec) / 1000; + } + break; case CV_CAP_PROP_BRIGHTNESS: capture->control.id = V4L2_CID_BRIGHTNESS; break; diff --cc modules/java/generator/rst_parser.py index a6ee3f0,d32323a..ad83585 --- a/modules/java/generator/rst_parser.py +++ b/modules/java/generator/rst_parser.py @@@ -1,7 -1,7 +1,7 @@@ - #/usr/bin/env python + #!/usr/bin/env python import os, sys, re, string, fnmatch -allmodules = ["core", "flann", "imgproc", "ml", "highgui", "video", "features2d", "calib3d", "objdetect", "legacy", "contrib", "gpu", "androidcamera", "java", "python", "stitching", "ts", "photo", "nonfree", "videostab", "ocl", "superres"] +allmodules = ["core", "flann", "imgproc", "ml", "highgui", "video", "features2d", "calib3d", "objdetect", "legacy", "contrib", "gpu", "androidcamera", "java", "python", "stitching", "ts", "photo", "nonfree", "videostab", "ocl", "softcascade", "superres"] verbose = False show_warnings = True show_errors = True diff --cc modules/ocl/src/arithm.cpp index 82208f2,cba6ccf..6e57bb8 --- a/modules/ocl/src/arithm.cpp +++ b/modules/ocl/src/arithm.cpp @@@ -260,13 -280,13 +257,13 @@@ void cv::ocl::add(const oclMat &src1, c void cv::ocl::subtract(const oclMat &src1, const oclMat &src2, oclMat &dst) { - arithmetic_run(src1, src2, dst, "arithm_sub", &arithm_sub); - arithmetic_run(src1, src2, dst, "arithm_add", &arithm_add, MAT_SUB); ++ arithmetic_run(src1, src2, dst, "arithm_add", &arithm_add); } void cv::ocl::subtract(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask) { - arithmetic_run(src1, src2, dst, mask, "arithm_sub_with_mask", &arithm_sub); - arithmetic_run(src1, src2, dst, mask, "arithm_add_with_mask", &arithm_add, MAT_SUB); ++ arithmetic_run(src1, src2, dst, mask, "arithm_add_with_mask", &arithm_add); } -typedef void (*MulDivFunc)(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, +typedef void (*MulDivFunc)(const oclMat &src1, const oclMat &src2, oclMat &dst, String kernelName, const char **kernelString, void *scalar); void cv::ocl::multiply(const oclMat &src1, const oclMat &src2, oclMat &dst, double scalar) @@@ -451,14 -468,14 +448,16 @@@ void cv::ocl::add(const oclMat &src1, c void cv::ocl::subtract(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask) { - String kernelName = mask.data ? "arithm_s_sub_with_mask" : "arithm_s_sub"; - const char **kernelString = mask.data ? &arithm_sub_scalar_mask : &arithm_sub_scalar; - string kernelName = mask.data ? "arithm_s_add_with_mask" : "arithm_s_add"; ++ String kernelName = mask.data ? "arithm_s_add_with_mask" : "arithm_s_add"; + const char **kernelString = mask.data ? &arithm_add_scalar_mask : &arithm_add_scalar; ++ arithmetic_scalar( src1, src2, dst, mask, kernelName, kernelString, 1); } void cv::ocl::subtract(const Scalar &src2, const oclMat &src1, oclMat &dst, const oclMat &mask) { - String kernelName = mask.data ? "arithm_s_sub_with_mask" : "arithm_s_sub"; - const char **kernelString = mask.data ? &arithm_sub_scalar_mask : &arithm_sub_scalar; - string kernelName = mask.data ? "arithm_s_add_with_mask" : "arithm_s_add"; ++ String kernelName = mask.data ? "arithm_s_add_with_mask" : "arithm_s_add"; + const char **kernelString = mask.data ? &arithm_add_scalar_mask : &arithm_add_scalar; ++ arithmetic_scalar( src1, src2, dst, mask, kernelName, kernelString, -1); } void cv::ocl::divide(double scalar, const oclMat &src, oclMat &dst) diff --cc modules/ocl/src/initialization.cpp index 432dac6,c9ce89f..37a69ce --- a/modules/ocl/src/initialization.cpp +++ b/modules/ocl/src/initialization.cpp @@@ -394,6 -397,15 +394,15 @@@ namespace c } break; + case IS_CPU_DEVICE: + { + cl_device_type devicetype; - openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum], - CL_DEVICE_TYPE, sizeof(cl_device_type), ++ openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum], ++ CL_DEVICE_TYPE, sizeof(cl_device_type), + &devicetype, NULL)); + *(bool*)info = (devicetype == CVCL_DEVICE_TYPE_CPU); + } + break; default: CV_Error(-1, "Invalid device info type"); break; diff --cc modules/ocl/src/opencl/brute_force_match.cl index e76fb1d,7446c77..fc7e515 --- a/modules/ocl/src/opencl/brute_force_match.cl +++ b/modules/ocl/src/opencl/brute_force_match.cl @@@ -13,38 -66,30 +66,30 @@@ int bit1Count(float x return (float)c; } + #ifndef DIST_TYPE + #define DIST_TYPE 0 + #endif + + #if (DIST_TYPE == 0) + #define DIST(x, y) fabs((x) - (y)) + #elif (DIST_TYPE == 1) + #define DIST(x, y) (((x) - (y)) * ((x) - (y))) + #elif (DIST_TYPE == 2) + #define DIST(x, y) bit1Count((uint)(x) ^ (uint)(y)) -#endif ++#endif + + float reduce_block(__local float *s_query, __local float *s_train, - int block_size, int lidx, - int lidy, - int distType + int lidy ) { - /* there are threee types in the reducer. the first is L1Dist, which to sum the abs(v1, v2), the second is L2Dist, which to - sum the (v1 - v2) * (v1 - v2), the third is humming, which to popc(v1 ^ v2), popc is to count the bits are set to 1*/ float result = 0; - switch(distType) + #pragma unroll + for (int j = 0 ; j < BLOCK_SIZE ; j++) { - case 0: - for (int j = 0 ; j < block_size ; j++) - { - result += fabs(s_query[lidy * block_size + j] - s_train[j * block_size + lidx]); - } - break; - case 1: - for (int j = 0 ; j < block_size ; j++) - { - float qr = s_query[lidy * block_size + j] - s_train[j * block_size + lidx]; - result += qr * qr; - } - break; - case 2: - for (int j = 0 ; j < block_size ; j++) - { - result += bit1Count((uint)s_query[lidy * block_size + j] ^ (uint)s_train[(uint)j * block_size + lidx]); - } - break; + result += DIST(s_query[lidy * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + lidx]); } return result; } diff --cc modules/ocl/src/opencl/filter_sep_row.cl index bfe6cd4,5524041..9dc4983 --- a/modules/ocl/src/opencl/filter_sep_row.cl +++ b/modules/ocl/src/opencl/filter_sep_row.cl @@@ -465,4 -465,7 +465,5 @@@ __kernel __attribute__((reqd_work_group start_addr = mad24(y,dst_step_in_pixel,x); dst[start_addr] = sum; } - -} - + +} diff --cc modules/ocl/src/opencl/imgproc_warpPerspective.cl index a37ffa1,edbe42c..ef9e770 --- a/modules/ocl/src/opencl/imgproc_warpPerspective.cl +++ b/modules/ocl/src/opencl/imgproc_warpPerspective.cl @@@ -112,7 -116,7 +116,7 @@@ __kernel void warpPerspectiveNN_C1_D0(_ sval.s1 = scon.s1 ? src[spos.s1] : 0; sval.s2 = scon.s2 ? src[spos.s2] : 0; sval.s3 = scon.s3 ? src[spos.s3] : 0; - dval = convert_uchar4(dcon != 0) ? sval : dval; - dval = convert_uchar4(dcon) != (uchar4)(0,0,0,0) ? sval : dval; ++ dval = convert_uchar4(dcon) != (uchar4)(0,0,0,0) ? sval : dval; *d = dval; } } @@@ -680,5 -684,6 +684,5 @@@ __kernel void warpPerspectiveCubic_C4_D dst[dst_offset+dy*dstStep+dx] = sum; } - } + } } - diff --cc modules/ocl/src/opencl/objdetect_hog.cl index db11ed1,8852fac..64ae3ea --- a/modules/ocl/src/opencl/objdetect_hog.cl +++ b/modules/ocl/src/opencl/objdetect_hog.cl @@@ -53,76 -53,96 +53,96 @@@ //---------------------------------------------------------------------------- // Histogram computation - - __kernel void compute_hists_kernel(const int width, const int cblock_stride_x, const int cblock_stride_y, - const int cnbins, const int cblock_hist_size, const int img_block_width, - const int grad_quadstep, const int qangle_step, - __global const float* grad, __global const uchar* qangle, - const float scale, __global float* block_hists, __local float* smem) + // 12 threads for a cell, 12x4 threads per block + __kernel void compute_hists_kernel( + const int cblock_stride_x, const int cblock_stride_y, - const int cnbins, const int cblock_hist_size, const int img_block_width, ++ const int cnbins, const int cblock_hist_size, const int img_block_width, + const int blocks_in_group, const int blocks_total, + const int grad_quadstep, const int qangle_step, + __global const float* grad, __global const uchar* qangle, + const float scale, __global float* block_hists, __local float* smem) { - const int lidX = get_local_id(0); + const int lx = get_local_id(0); + const int lp = lx / 24; /* local group id */ + const int gid = get_group_id(0) * blocks_in_group + lp;/* global group id */ + const int gidY = gid / img_block_width; + const int gidX = gid - gidY * img_block_width; + + const int lidX = lx - lp * 24; const int lidY = get_local_id(1); - const int gidX = get_group_id(0); - const int gidY = get_group_id(1); - const int cell_x = lidX / 16; + const int cell_x = lidX / 12; const int cell_y = lidY; - const int cell_thread_x = lidX & 0xF; + const int cell_thread_x = lidX - cell_x * 12; - __local float* hists = smem; - __local float* final_hist = smem + cnbins * 48; - __local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X * ++ __local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X * + CELLS_PER_BLOCK_Y * 12 + CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y); - __local float* final_hist = hists + cnbins * ++ __local float* final_hist = hists + cnbins * + (CELLS_PER_BLOCK_X * CELLS_PER_BLOCK_Y * 12); const int offset_x = gidX * cblock_stride_x + (cell_x << 2) + cell_thread_x; const int offset_y = gidY * cblock_stride_y + (cell_y << 2); - __global const float* grad_ptr = grad + offset_y * grad_quadstep + (offset_x << 1); - __global const uchar* qangle_ptr = qangle + offset_y * qangle_step + (offset_x << 1); - - // 12 means that 12 pixels affect on block's cell (in one row) - if (cell_thread_x < 12) - { - __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + cell_thread_x; - for (int bin_id = 0; bin_id < cnbins; ++bin_id) - hist[bin_id * 48] = 0.f; - __global const float* grad_ptr = (gid < blocks_total) ? ++ __global const float* grad_ptr = (gid < blocks_total) ? + grad + offset_y * grad_quadstep + (offset_x << 1) : grad; + __global const uchar* qangle_ptr = (gid < blocks_total) ? + qangle + offset_y * qangle_step + (offset_x << 1) : qangle; - const int dist_x = -4 + cell_thread_x - 4 * cell_x; - __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + ++ __local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) + + cell_thread_x; + for (int bin_id = 0; bin_id < cnbins; ++bin_id) + hist[bin_id * 48] = 0.f; - const int dist_y_begin = -4 - 4 * lidY; - for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y) - { - float2 vote = (float2) (grad_ptr[0], grad_ptr[1]); - uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]); + const int dist_x = -4 + cell_thread_x - 4 * cell_x; + const int dist_center_x = dist_x - 4 * (1 - 2 * cell_x); - grad_ptr += grad_quadstep; - qangle_ptr += qangle_step; + const int dist_y_begin = -4 - 4 * lidY; + for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y) + { + float2 vote = (float2) (grad_ptr[0], grad_ptr[1]); + uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]); - int dist_center_y = dist_y - 4 * (1 - 2 * cell_y); - int dist_center_x = dist_x - 4 * (1 - 2 * cell_x); + grad_ptr += grad_quadstep; + qangle_ptr += qangle_step; - float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * dist_center_x) * scale); - float interp_weight = (8.f - fabs(dist_y + 0.5f)) * (8.f - fabs(dist_x + 0.5f)) / 64.f; + int dist_center_y = dist_y - 4 * (1 - 2 * cell_y); - hist[bin.x * 48] += gaussian * interp_weight * vote.x; - hist[bin.y * 48] += gaussian * interp_weight * vote.y; - } - float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * ++ float gaussian = exp(-(dist_center_y * dist_center_y + dist_center_x * + dist_center_x) * scale); - float interp_weight = (8.f - fabs(dist_y + 0.5f)) * ++ float interp_weight = (8.f - fabs(dist_y + 0.5f)) * + (8.f - fabs(dist_x + 0.5f)) / 64.f; - volatile __local float* hist_ = hist; - for (int bin_id = 0; bin_id < cnbins; ++bin_id, hist_ += 48) - { - if (cell_thread_x < 6) hist_[0] += hist_[6]; - if (cell_thread_x < 3) hist_[0] += hist_[3]; - if (cell_thread_x == 0) - final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = hist_[0] + hist_[1] + hist_[2]; - } + hist[bin.x * 48] += gaussian * interp_weight * vote.x; + hist[bin.y * 48] += gaussian * interp_weight * vote.y; } - barrier(CLK_LOCAL_MEM_FENCE); - __global float* block_hist = block_hists + (gidY * img_block_width + gidX) * cblock_hist_size; + volatile __local float* hist_ = hist; + for (int bin_id = 0; bin_id < cnbins; ++bin_id, hist_ += 48) + { + if (cell_thread_x < 6) + hist_[0] += hist_[6]; + barrier(CLK_LOCAL_MEM_FENCE); + if (cell_thread_x < 3) + hist_[0] += hist_[3]; + #ifdef WAVE_SIZE_1 + barrier(CLK_LOCAL_MEM_FENCE); + #endif + if (cell_thread_x == 0) - final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = ++ final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = + hist_[0] + hist_[1] + hist_[2]; + } + #ifdef WAVE_SIZE_1 + barrier(CLK_LOCAL_MEM_FENCE); + #endif - int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 16 + cell_thread_x; - if (tid < cblock_hist_size) + int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x; + if ((tid < cblock_hist_size) && (gid < blocks_total)) + { - __global float* block_hist = block_hists + ++ __global float* block_hist = block_hists + + (gidY * img_block_width + gidX) * cblock_hist_size; block_hist[tid] = final_hist[tid]; + } } //------------------------------------------------------------- diff --cc modules/ocl/test/test_brute_force_matcher.cpp index 424781f,d658c32..71b299e --- a/modules/ocl/test/test_brute_force_matcher.cpp +++ b/modules/ocl/test/test_brute_force_matcher.cpp @@@ -172,49 -166,33 +166,33 @@@ namespac cv::ocl::BruteForceMatcher_OCL_base matcher(distType); - // assume support atomic. - //if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS)) - //{ - // try - // { - // std::vector< std::vector > matches; - // matcher.radiusMatch(loadMat(query), loadMat(train), matches, radius); - // } - // catch (const cv::Exception& e) - // { - // ASSERT_EQ(CV_StsNotImplemented, e.code); - // } - //} - //else - { - std::vector< std::vector > matches; - matcher.radiusMatch(cv::ocl::oclMat(query), cv::ocl::oclMat(train), matches, radius); + std::vector< std::vector > matches; + matcher.radiusMatch(cv::ocl::oclMat(query), cv::ocl::oclMat(train), matches, radius); - ASSERT_EQ(static_cast(queryDescCount), matches.size()); + ASSERT_EQ(static_cast(queryDescCount), matches.size()); - int badCount = 0; - for (size_t i = 0; i < matches.size(); i++) + int badCount = 0; + for (size_t i = 0; i < matches.size(); i++) + { + if ((int)matches[i].size() != 1) { - if ((int)matches[i].size() != 1) - { + badCount++; + } + else + { + cv::DMatch match = matches[i][0]; + if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i * countFactor) || (match.imgIdx != 0)) badCount++; - } - else - { - cv::DMatch match = matches[i][0]; - if ((match.queryIdx != (int)i) || (match.trainIdx != (int)i * countFactor) || (match.imgIdx != 0)) - badCount++; - } } - - ASSERT_EQ(0, badCount); } + + ASSERT_EQ(0, badCount); } - INSTANTIATE_TEST_CASE_P(GPU_Features2D, BruteForceMatcher, testing::Combine( - //ALL_DEVICES, - testing::Values(DistType(cv::ocl::BruteForceMatcher_OCL_base::L1Dist), DistType(cv::ocl::BruteForceMatcher_OCL_base::L2Dist)), - testing::Values(DescriptorSize(57), DescriptorSize(64), DescriptorSize(83), DescriptorSize(128), DescriptorSize(179), DescriptorSize(256), DescriptorSize(304)))); - INSTANTIATE_TEST_CASE_P(OCL_Features2D, BruteForceMatcher, ++ INSTANTIATE_TEST_CASE_P(OCL_Features2D, BruteForceMatcher, + testing::Combine( + testing::Values(DistType(cv::ocl::BruteForceMatcher_OCL_base::L1Dist), DistType(cv::ocl::BruteForceMatcher_OCL_base::L2Dist)), + testing::Values(DescriptorSize(57), DescriptorSize(64), DescriptorSize(83), DescriptorSize(128), DescriptorSize(179), DescriptorSize(256), DescriptorSize(304)))); } // namespace #endif diff --cc samples/cpp/freak_demo.cpp index b420e45,5112eae..82a435f --- a/samples/cpp/freak_demo.cpp +++ b/samples/cpp/freak_demo.cpp @@@ -72,8 -71,8 +72,8 @@@ int main( int argc, char** argv ) return -1; } - Mat imgB = imread(argv[2], CV_LOAD_IMAGE_GRAYSCALE ); + Mat imgB = imread(argv[2], IMREAD_GRAYSCALE ); - if( !imgA.data ) { + if( !imgB.data ) { std::cout << " --(!) Error reading image " << argv[2] << std::endl; return -1; } diff --cc samples/cpp/lkdemo.cpp index 3913e2c,bbb734c..c665cfd --- a/samples/cpp/lkdemo.cpp +++ b/samples/cpp/lkdemo.cpp @@@ -28,17 -27,19 +27,19 @@@ bool addRemovePt = false static void onMouse( int event, int x, int y, int /*flags*/, void* /*param*/ ) { - if( event == CV_EVENT_LBUTTONDOWN ) + if( event == EVENT_LBUTTONDOWN ) { - point = Point2f((float)x,(float)y); + point = Point2f((float)x, (float)y); addRemovePt = true; } } int main( int argc, char** argv ) { + help(); + VideoCapture cap; - TermCriteria termcrit(CV_TERMCRIT_ITER|CV_TERMCRIT_EPS, 20, 0.03); + TermCriteria termcrit(TermCriteria::COUNT|TermCriteria::EPS,20,0.03); Size subPixWinSize(10,10), winSize(31,31); const int MAX_COUNT = 500; diff --cc samples/python2/common.py index 5996695,f3c4101..3988fe2 --- a/samples/python2/common.py +++ b/samples/python2/common.py @@@ -1,5 -1,7 +1,7 @@@ + #!/usr/bin/env python + ''' -This module contais some common routines used by other samples. +This module contains some common routines used by other samples. ''' import numpy as np