- USE_TEMP_BUFFER fixes the main problem (improper buffer management) and
prevents bad images in the first place
-Added v4l2 support for getting capture property CV_CAP_PROP_POS_MSEC.
+ 11th patch: April 2, 2013, Forrest Reiling forrest.reiling@gmail.com
++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!
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;
}
/* initialize the control structure */
switch (property_id) {
- return 1000 * capture->timestamp.tv_sec + ((double) capture->timestamp.tv_usec) / 1000;
+ case CV_CAP_PROP_POS_MSEC:
+ if (capture->FirstCapture) {
+ return 0;
+ } else {
++ 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;
- #/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
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)
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)
}
break;
- openCLSafeCall(clGetDeviceInfo(impl->devices[impl->devnum],
- CL_DEVICE_TYPE, sizeof(cl_device_type),
+ case IS_CPU_DEVICE:
+ {
+ cl_device_type devicetype;
++ 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;
return (float)c;
}
-#endif
+ #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
+
+
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;
}
start_addr = mad24(y,dst_step_in_pixel,x);
dst[start_addr] = sum;
}
-
-}
-
+
+}
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;
}
}
dst[dst_offset+dy*dstStep+dx] = sum;
}
- }
+ }
}
-
//----------------------------------------------------------------------------
// 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];
+ }
}
//-------------------------------------------------------------
cv::ocl::BruteForceMatcher_OCL_base matcher(distType);
- // assume support atomic.
- //if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS))
- //{
- // try
- // {
- // std::vector< std::vector<cv::DMatch> > 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<cv::DMatch> > matches;
- matcher.radiusMatch(cv::ocl::oclMat(query), cv::ocl::oclMat(train), matches, radius);
+ std::vector< std::vector<cv::DMatch> > matches;
+ matcher.radiusMatch(cv::ocl::oclMat(query), cv::ocl::oclMat(train), matches, radius);
- ASSERT_EQ(static_cast<size_t>(queryDescCount), matches.size());
+ ASSERT_EQ(static_cast<size_t>(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
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;
}
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;
+ #!/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
projects / profile / ivi / opencv.git / commitdiff